This is an archive of the discontinued LLVM Phabricator instance.

Paths

Table of Contentst

-
clang/
-
lib/CodeGen/
-
CodeGen/
3/6
CGOpenMPRuntime.cpp
-
test/OpenMP/
-
OpenMP/
-
target_depend_codegen.cpp

Differential D84470

[OpenMP 5.0] Fix user-defined mapper privatization in tasks
ClosedPublic

Authored by ABataev on Jul 23 2020, 3:51 PM.

Download Raw Diff

Details

Reviewers

protze.joachim
jdoerfert
lildmh
ronlieb

Commits

rGd5ce8233bfcf: [OpenMP 5.0] Fix user-defined mapper privatization in tasks

Summary

This patch fixes the problem that user-defined mapper array is not correctly privatized inside a task. This problem causes openmp/libomptarget/test/offloading/target_depend_nowait.cpp fails.

Diff Detail

Repository: rG LLVM Github Monorepo

Event Timeline

lildmh created this revision.Jul 23 2020, 3:51 PM

Herald added a reviewer: jdoerfert. · View Herald TranscriptJul 23 2020, 3:51 PM

Herald added subscribers: sstefan1, guansong, yaxunl. · View Herald Transcript

I tried the patch and indeed it fixes the problem with target_depend_nowait.cpp. I'll let someone else do the review though.

I'm not enough invested in the compiler implementation to judge this patch.
I can confirm that this patch fixes the segmentation fault I have seen.

Are these tests supposed to work for offloading to nvidia GPU as well?
I tried and added:

// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda

which gives me:

ptxas /tmp/target_depend_nowait-37c5c4.s, line 355; error   : Call has wrong number of parameters
ptxas /tmp/target_depend_nowait-37c5c4.s, line 628; error   : Call has wrong number of parameters
ptxas /tmp/target_depend_nowait-37c5c4.s, line 2331; warning : Instruction 'vote' without '.sync' may produce unpredictable results on sm_70 and later architectures
ptxas /tmp/target_depend_nowait-37c5c4.s, line 2331; warning : Instruction 'vote' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas fatal   : Ptx assembly aborted due to errors

I guess this error is not directly related to your patch?

clang/lib/CodeGen/CGOpenMPRuntime.cpp
10524	Does this mean, that a task is only created if depend clauses are provided? Does waiting with taskwait work without this wrapper task?

The problem you had is probably not related to this patch. A lot of libomptarget tests are not working for nvptx now. I suppose someone is going to fix them.

clang/lib/CodeGen/CGOpenMPRuntime.cpp
10524	The case you describe, which is to have a task wrapped around target, is solved with modification on line 9592 This part is to deal with other cases, which is consistent with line 10455 if you look a bit further

lildmh added a reviewer: ABataev.Jul 24 2020, 3:26 PM

ABataev added inline comments.Jul 27 2020, 8:47 AM

clang/lib/CodeGen/CGOpenMPRuntime.cpp
8915	Do you really need this `IsTask` flag? The array is created always when `Info.NumberOfPtrs > 0`. Maybe, just do not check for HasMapper here?
10524	Better to add a flag `RequiresOuterTask = D.hasClausesOfKind<OMPDependClause>()` and use it here and where the task is generated (line 10455)

lildmh marked 2 inline comments as done.Jul 27 2020, 10:51 AM

lildmh added inline comments.

clang/lib/CodeGen/CGOpenMPRuntime.cpp
8915	This is an optimization. The mapper array can be optimized out when there is no actual mappers, so supposed to have less memory accesses in this common case. Not checking `HasMapper` will disable this optimization.
10524	Sure, please accept this patch after that.

protze.joachim mentioned this in D84557: [OpenMP][Tests] Enable nvptx64 testing for most libomptarget tests.Jul 28 2020, 2:15 AM

Add a test

What is the current status of this patch?
@lildmh could you update this patch? I'd like to test it against
https://bugs.llvm.org/show_bug.cgi?id=47122

In D84470#2191752, @ABataev wrote:

Add a test

@ABataev did you ask for a codegen test, or is the runtime test openmp/libomptarget/test/offloading/target_depend_nowait.cpp enough?
The test fails without this patch and succeeds with this patch.

In D84470#2223678, @protze.joachim wrote:

In D84470#2191752, @ABataev wrote:

Add a test

@ABataev did you ask for a codegen test, or is the runtime test openmp/libomptarget/test/offloading/target_depend_nowait.cpp enough?
The test fails without this patch and succeeds with this patch.

Yes, I know that it fails, but each functional change in the frontend must have the corresponding test in clang itself.

I'll add a test to it. Please give me some time as I am busy with other stuffs.

It looks like @lildmh is not able to provide a test any time soon. @grokos Could you take over this mapper patch as well? I think the current segfault state is rather unfortunate.

i applied this patch to our downstream build for amdgcn. it resolves the problem i was seeing with traps on tests with task wait depend tests from sollve.
perhaps rebase it to latest llvm.

I can update this patch in a week when I'm back from vacation.

[AMD Official Use Only - Internal Distribution Only]

Hi Alexey
Any ETA on this one ?

Get Outlook for iOShttps://aka.ms/o0ukef

From: Alexey Bataev via Phabricator <reviews@reviews.llvm.org>
Sent: Tuesday, September 8, 2020 1:39:06 PM
To: lildmh@gmail.com <lildmh@gmail.com>; protze.joachim@gmail.com <protze.joachim@gmail.com>; jdoerfert@anl.gov <jdoerfert@anl.gov>; a.bataev@hotmail.com <a.bataev@hotmail.com>
Cc: Lieberman, Ron <Ron.Lieberman@amd.com>; xw111luoye@gmail.com <xw111luoye@gmail.com>; georgios.rokos@intel.com <georgios.rokos@intel.com>; Liu, Yaxun (Sam) <Yaxun.Liu@amd.com>; zhang.guansong@gmail.com <zhang.guansong@gmail.com>; stefomeister@gmail.com <stefomeister@gmail.com>
Subject: [PATCH] D84470: [OpenMP 5.0] Fix user-defined mapper privatization in tasks

[CAUTION: External Email]

ABataev added a comment.

I can update this patch in a week when I'm back from vacation.

CHANGES SINCE LAST ACTION

https://nam11.safelinks.protection.outlook.com/?url=https%3A%2F%2Freviews.llvm.org%2FD84470%2Fnew%2F&amp;data=02%7C01%7Cron.lieberman%40amd.com%7C6582a2f2d31d4bbf9b4a08d8541e16cb%7C3dd8961fe4884e608e11a82d994e183d%7C0%7C0%7C637351835496784231&amp;sdata=8%2BJbFULmzVN%2FdafC0QYX%2B9LwN7d1HFZ9Qxt6k4y85%2FA%3D&amp;reserved=0

https://nam11.safelinks.protection.outlook.com/?url=https%3A%2F%2Freviews.llvm.org%2FD84470&data=02%7C01%7Cron.lieberman%40amd.com%7C6582a2f2d31d4bbf9b4a08d8541e16cb%7C3dd8961fe4884e608e11a82d994e183d%7C0%7C0%7C637351835496794187&sdata=gyj3%2B4K3oPlvedNjOzbsnosciYUIxuTOoKq%2BsfnsAOQ%3D&reserved=0

Hi Ron, hope to prepare it by the end of this week, maybe earlier, if everything goes well. I have some technical issues, but anyway will update the patch and the tests.

ABataev commandeered this revision.Sep 16 2020, 1:31 PM

ABataev edited reviewers, added: lildmh; removed: ABataev.

Updated + added clang test.

Herald added a project: Restricted Project. · View Herald TranscriptSep 16 2020, 1:32 PM

Herald added a subscriber: cfe-commits. · View Herald Transcript

Harbormaster completed remote builds in B71924: Diff 292330.Sep 16 2020, 2:28 PM

The latest patch applied cleanly to our downstream port.
builds fine, tests very nicely as well. All the failing SOLLVE task wait depend tests now pass.

In D84470#2277991, @ronlieb wrote:

The latest patch applied cleanly to our downstream port.
builds fine, tests very nicely as well. All the failing SOLLVE task wait depend tests now pass.

Could you accept it then?

accepted. not sure if there were any outstanding comments from Johannes. your call

This revision is now accepted and ready to land.Sep 17 2020, 6:21 AM

This revision was landed with ongoing or failed builds.Sep 17 2020, 8:22 AM

Closed by commit rGd5ce8233bfcf: [OpenMP 5.0] Fix user-defined mapper privatization in tasks (authored by ABataev). · Explain Why

This revision was automatically updated to reflect the committed changes.

ABataev added a commit: rGd5ce8233bfcf: [OpenMP 5.0] Fix user-defined mapper privatization in tasks.

grokos mentioned this in D87946: [OpenMP] Add Location Fields to Libomptarget Runtime for Debugging.Oct 12 2020, 1:21 PM

Revision Contents

Path

Size

clang/

lib/

CodeGen/

CGOpenMPRuntime.cpp

54 lines

test/

OpenMP/

target_depend_codegen.cpp

89 lines

Diff 292515

clang/lib/CodeGen/CGOpenMPRuntime.cpp

This file is larger than 256 KB, so syntax highlighting is disabled by default.

Show First 20 Lines • Show All 3,778 Lines • ▼ Show 20 Lines	OpenMPDirectiveKind Kind = isOpenMPTaskLoopDirective(D.getDirectiveKind())
? OMPD_taskloop		? OMPD_taskloop
: OMPD_task;		: OMPD_task;
const CapturedStmt &CS = *D.getCapturedStmt(Kind);		const CapturedStmt &CS = *D.getCapturedStmt(Kind);
CodeGenFunction::CGCapturedStmtInfo CapturesInfo(CS);		CodeGenFunction::CGCapturedStmtInfo CapturesInfo(CS);
LValue SrcBase;		LValue SrcBase;
bool IsTargetTask =		bool IsTargetTask =
isOpenMPTargetDataManagementDirective(D.getDirectiveKind()) \|\|		isOpenMPTargetDataManagementDirective(D.getDirectiveKind()) \|\|
isOpenMPTargetExecutionDirective(D.getDirectiveKind());		isOpenMPTargetExecutionDirective(D.getDirectiveKind());
// For target-based directives skip 3 firstprivate arrays BasePointersArray,		// For target-based directives skip 4 firstprivate arrays BasePointersArray,
// PointersArray and SizesArray. The original variables for these arrays are		// PointersArray, SizesArray, and MappersArray. The original variables for
// not captured and we get their addresses explicitly.		// these arrays are not captured and we get their addresses explicitly.
if ((!IsTargetTask && !Data.FirstprivateVars.empty() && ForDup) \|\|		if ((!IsTargetTask && !Data.FirstprivateVars.empty() && ForDup) \|\|
(IsTargetTask && KmpTaskSharedsPtr.isValid())) {		(IsTargetTask && KmpTaskSharedsPtr.isValid())) {
SrcBase = CGF.MakeAddrLValue(		SrcBase = CGF.MakeAddrLValue(
CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(		CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
KmpTaskSharedsPtr, CGF.ConvertTypeForMem(SharedsPtrTy)),		KmpTaskSharedsPtr, CGF.ConvertTypeForMem(SharedsPtrTy)),
SharedsTy);		SharedsTy);
}		}
FI = cast<RecordDecl>(FI->getType()->getAsTagDecl())->field_begin();		FI = cast<RecordDecl>(FI->getType()->getAsTagDecl())->field_begin();
for (const PrivateDataTy &Pair : Privates) {		for (const PrivateDataTy &Pair : Privates) {
// Do not initialize private locals.		// Do not initialize private locals.
if (Pair.second.isLocalPrivate()) {		if (Pair.second.isLocalPrivate()) {
++FI;		++FI;
continue;		continue;
}		}
const VarDecl *VD = Pair.second.PrivateCopy;		const VarDecl *VD = Pair.second.PrivateCopy;
const Expr *Init = VD->getAnyInitializer();		const Expr *Init = VD->getAnyInitializer();
if (Init && (!ForDup \|\| (isa<CXXConstructExpr>(Init) &&		if (Init && (!ForDup \|\| (isa<CXXConstructExpr>(Init) &&
!CGF.isTrivialInitializer(Init)))) {		!CGF.isTrivialInitializer(Init)))) {
LValue PrivateLValue = CGF.EmitLValueForField(PrivatesBase, *FI);		LValue PrivateLValue = CGF.EmitLValueForField(PrivatesBase, *FI);
if (const VarDecl *Elem = Pair.second.PrivateElemInit) {		if (const VarDecl *Elem = Pair.second.PrivateElemInit) {
const VarDecl *OriginalVD = Pair.second.Original;		const VarDecl *OriginalVD = Pair.second.Original;
// Check if the variable is the target-based BasePointersArray,		// Check if the variable is the target-based BasePointersArray,
// PointersArray or SizesArray.		// PointersArray, SizesArray, or MappersArray.
LValue SharedRefLValue;		LValue SharedRefLValue;
QualType Type = PrivateLValue.getType();		QualType Type = PrivateLValue.getType();
const FieldDecl *SharedField = CapturesInfo.lookup(OriginalVD);		const FieldDecl *SharedField = CapturesInfo.lookup(OriginalVD);
if (IsTargetTask && !SharedField) {		if (IsTargetTask && !SharedField) {
assert(isa<ImplicitParamDecl>(OriginalVD) &&		assert(isa<ImplicitParamDecl>(OriginalVD) &&
isa<CapturedDecl>(OriginalVD->getDeclContext()) &&		isa<CapturedDecl>(OriginalVD->getDeclContext()) &&
cast<CapturedDecl>(OriginalVD->getDeclContext())		cast<CapturedDecl>(OriginalVD->getDeclContext())
->getNumParams() == 0 &&		->getNumParams() == 0 &&
▲ Show 20 Lines • Show All 5,040 Lines • ▼ Show 20 Lines	for (unsigned I = 0; I < Info.NumberOfPtrs; ++I) {
Info.HasMapper = true;		Info.HasMapper = true;
}		}
Address MAddr = CGF.Builder.CreateConstArrayGEP(MappersArray, I);		Address MAddr = CGF.Builder.CreateConstArrayGEP(MappersArray, I);
CGF.Builder.CreateStore(MFunc, MAddr);		CGF.Builder.CreateStore(MFunc, MAddr);
}		}
}		}
}		}

		namespace {
		/// Additional arguments for emitOffloadingArraysArgument function.
		struct ArgumentsOptions {
		bool ForEndCall = false;
		bool IsTask = false;
		ArgumentsOptions() = default;
		ArgumentsOptions(bool ForEndCall, bool IsTask)
		: ForEndCall(ForEndCall), IsTask(IsTask) {}
		};
		} // namespace

/// Emit the arguments to be passed to the runtime library based on the		/// Emit the arguments to be passed to the runtime library based on the
/// arrays of base pointers, pointers, sizes, map types, and mappers. If		/// arrays of base pointers, pointers, sizes, map types, and mappers. If
/// ForEndCall, emit map types to be passed for the end of the region instead of		/// ForEndCall, emit map types to be passed for the end of the region instead of
/// the beginning.		/// the beginning.
static void emitOffloadingArraysArgument(		static void emitOffloadingArraysArgument(
CodeGenFunction &CGF, llvm::Value *&BasePointersArrayArg,		CodeGenFunction &CGF, llvm::Value *&BasePointersArrayArg,
llvm::Value &PointersArrayArg, llvm::Value &SizesArrayArg,		llvm::Value &PointersArrayArg, llvm::Value &SizesArrayArg,
llvm::Value &MapTypesArrayArg, llvm::Value &MappersArrayArg,		llvm::Value &MapTypesArrayArg, llvm::Value &MappersArrayArg,
CGOpenMPRuntime::TargetDataInfo &Info, bool ForEndCall = false) {		CGOpenMPRuntime::TargetDataInfo &Info,
assert((!ForEndCall \|\| Info.separateBeginEndCalls()) &&		const ArgumentsOptions &Options = ArgumentsOptions()) {
		assert((!Options.ForEndCall \|\| Info.separateBeginEndCalls()) &&
"expected region end call to runtime only when end call is separate");		"expected region end call to runtime only when end call is separate");
CodeGenModule &CGM = CGF.CGM;		CodeGenModule &CGM = CGF.CGM;
if (Info.NumberOfPtrs) {		if (Info.NumberOfPtrs) {
BasePointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(		BasePointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),		llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
Info.BasePointersArray,		Info.BasePointersArray,
/Idx0=/0, /Idx1=/0);		/Idx0=/0, /Idx1=/0);
PointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(		PointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),		llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
Info.PointersArray,		Info.PointersArray,
/Idx0=/0,		/Idx0=/0,
/Idx1=/0);		/Idx1=/0);
SizesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(		SizesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
llvm::ArrayType::get(CGM.Int64Ty, Info.NumberOfPtrs), Info.SizesArray,		llvm::ArrayType::get(CGM.Int64Ty, Info.NumberOfPtrs), Info.SizesArray,
/Idx0=/0, /Idx1=/0);		/Idx0=/0, /Idx1=/0);
MapTypesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(		MapTypesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
llvm::ArrayType::get(CGM.Int64Ty, Info.NumberOfPtrs),		llvm::ArrayType::get(CGM.Int64Ty, Info.NumberOfPtrs),
ForEndCall && Info.MapTypesArrayEnd ? Info.MapTypesArrayEnd		Options.ForEndCall && Info.MapTypesArrayEnd ? Info.MapTypesArrayEnd
: Info.MapTypesArray,		: Info.MapTypesArray,
/Idx0=/0,		/Idx0=/0,
/Idx1=/0);		/Idx1=/0);
		// Always emit the mapper array address in case of a target task for
		// privatization.
		if (!Options.IsTask && !Info.HasMapper)
		MappersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
		ABataevAuthorUnsubmitted Not Done Reply Inline Actions Do you really need this `IsTask` flag? The array is created always when `Info.NumberOfPtrs > 0`. Maybe, just do not check for HasMapper here? ABataev: Do you really need this `IsTask` flag? The array is created always when `Info.NumberOfPtrs > 0`.
		lildmhUnsubmitted Done Reply Inline Actions This is an optimization. The mapper array can be optimized out when there is no actual mappers, so supposed to have less memory accesses in this common case. Not checking `HasMapper` will disable this optimization. lildmh: This is an optimization. The mapper array can be optimized out when there is no actual mappers…
		else
MappersArrayArg =		MappersArrayArg =
Info.HasMapper		CGF.Builder.CreatePointerCast(Info.MappersArray, CGM.VoidPtrPtrTy);
? CGF.Builder.CreatePointerCast(Info.MappersArray, CGM.VoidPtrPtrTy)
: llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
} else {		} else {
BasePointersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);		BasePointersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
PointersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);		PointersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
SizesArrayArg = llvm::ConstantPointerNull::get(CGM.Int64Ty->getPointerTo());		SizesArrayArg = llvm::ConstantPointerNull::get(CGM.Int64Ty->getPointerTo());
MapTypesArrayArg =		MapTypesArrayArg =
llvm::ConstantPointerNull::get(CGM.Int64Ty->getPointerTo());		llvm::ConstantPointerNull::get(CGM.Int64Ty->getPointerTo());
MappersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);		MappersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
}		}
▲ Show 20 Lines • Show All 731 Lines • ▼ Show 20 Lines	auto &&TargetThenGen = [this, &ThenGen, &D, &InputInfo, &MapTypesArray,
// Map any list items in a map clause that were not captures because they		// Map any list items in a map clause that were not captures because they
// weren't referenced within the construct.		// weren't referenced within the construct.
MEHandler.generateAllInfo(CombinedInfo, /NotTargetParams=/true,		MEHandler.generateAllInfo(CombinedInfo, /NotTargetParams=/true,
MappedVarSet);		MappedVarSet);

TargetDataInfo Info;		TargetDataInfo Info;
// Fill up the arrays and create the arguments.		// Fill up the arrays and create the arguments.
emitOffloadingArrays(CGF, CombinedInfo, Info);		emitOffloadingArrays(CGF, CombinedInfo, Info);
		bool HasDependClauses = D.hasClausesOfKind<OMPDependClause>();
emitOffloadingArraysArgument(CGF, Info.BasePointersArray,		emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
Info.PointersArray, Info.SizesArray,		Info.PointersArray, Info.SizesArray,
Info.MapTypesArray, Info.MappersArray, Info);		Info.MapTypesArray, Info.MappersArray, Info,
		{/ForEndTask=/false, HasDependClauses});
InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;		InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;
InputInfo.BasePointersArray =		InputInfo.BasePointersArray =
Address(Info.BasePointersArray, CGM.getPointerAlign());		Address(Info.BasePointersArray, CGM.getPointerAlign());
InputInfo.PointersArray =		InputInfo.PointersArray =
Address(Info.PointersArray, CGM.getPointerAlign());		Address(Info.PointersArray, CGM.getPointerAlign());
InputInfo.SizesArray = Address(Info.SizesArray, CGM.getPointerAlign());		InputInfo.SizesArray = Address(Info.SizesArray, CGM.getPointerAlign());
InputInfo.MappersArray = Address(Info.MappersArray, CGM.getPointerAlign());		InputInfo.MappersArray = Address(Info.MappersArray, CGM.getPointerAlign());
MapTypesArray = Info.MapTypesArray;		MapTypesArray = Info.MapTypesArray;
▲ Show 20 Lines • Show All 594 Lines • ▼ Show 20 Lines	auto &&BeginThenGen = [this, &D, Device, &Info,

llvm::Value *BasePointersArrayArg = nullptr;		llvm::Value *BasePointersArrayArg = nullptr;
llvm::Value *PointersArrayArg = nullptr;		llvm::Value *PointersArrayArg = nullptr;
llvm::Value *SizesArrayArg = nullptr;		llvm::Value *SizesArrayArg = nullptr;
llvm::Value *MapTypesArrayArg = nullptr;		llvm::Value *MapTypesArrayArg = nullptr;
llvm::Value *MappersArrayArg = nullptr;		llvm::Value *MappersArrayArg = nullptr;
emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,		emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
SizesArrayArg, MapTypesArrayArg,		SizesArrayArg, MapTypesArrayArg,
MappersArrayArg, Info, /ForEndCall=/false);		MappersArrayArg, Info);

// Emit device ID if any.		// Emit device ID if any.
llvm::Value *DeviceID = nullptr;		llvm::Value *DeviceID = nullptr;
if (Device) {		if (Device) {
DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),		DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
CGF.Int64Ty, /isSigned=/true);		CGF.Int64Ty, /isSigned=/true);
} else {		} else {
DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);		DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);
Show All 23 Lines	auto &&EndThenGen = [this, Device, &Info](CodeGenFunction &CGF,

llvm::Value *BasePointersArrayArg = nullptr;		llvm::Value *BasePointersArrayArg = nullptr;
llvm::Value *PointersArrayArg = nullptr;		llvm::Value *PointersArrayArg = nullptr;
llvm::Value *SizesArrayArg = nullptr;		llvm::Value *SizesArrayArg = nullptr;
llvm::Value *MapTypesArrayArg = nullptr;		llvm::Value *MapTypesArrayArg = nullptr;
llvm::Value *MappersArrayArg = nullptr;		llvm::Value *MappersArrayArg = nullptr;
emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,		emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
SizesArrayArg, MapTypesArrayArg,		SizesArrayArg, MapTypesArrayArg,
MappersArrayArg, Info, /ForEndCall=/true);		MappersArrayArg, Info,
		{/ForEndCall=/true, /IsTask=/false});

// Emit device ID if any.		// Emit device ID if any.
llvm::Value *DeviceID = nullptr;		llvm::Value *DeviceID = nullptr;
if (Device) {		if (Device) {
DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),		DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
CGF.Int64Ty, /isSigned=/true);		CGF.Int64Ty, /isSigned=/true);
} else {		} else {
DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);		DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);
▲ Show 20 Lines • Show All 181 Lines • ▼ Show 20 Lines	auto &&TargetThenGen = [this, &ThenGen, &D, &InputInfo, &MapTypesArray](

// Get map clause information.		// Get map clause information.
MappableExprsHandler MEHandler(D, CGF);		MappableExprsHandler MEHandler(D, CGF);
MEHandler.generateAllInfo(CombinedInfo);		MEHandler.generateAllInfo(CombinedInfo);

TargetDataInfo Info;		TargetDataInfo Info;
// Fill up the arrays and create the arguments.		// Fill up the arrays and create the arguments.
emitOffloadingArrays(CGF, CombinedInfo, Info);		emitOffloadingArrays(CGF, CombinedInfo, Info);
		bool HasDependClauses = D.hasClausesOfKind<OMPDependClause>();
emitOffloadingArraysArgument(CGF, Info.BasePointersArray,		emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
Info.PointersArray, Info.SizesArray,		Info.PointersArray, Info.SizesArray,
Info.MapTypesArray, Info.MappersArray, Info);		Info.MapTypesArray, Info.MappersArray, Info,
		{/ForEndTask=/false, HasDependClauses});
		protze.joachimUnsubmitted Not Done Reply Inline Actions Does this mean, that a task is only created if depend clauses are provided? Does waiting with taskwait work without this wrapper task? protze.joachim: Does this mean, that a task is only created if depend clauses are provided? Does waiting with…
		lildmhUnsubmitted Done Reply Inline Actions The case you describe, which is to have a task wrapped around target, is solved with modification on line 9592 This part is to deal with other cases, which is consistent with line 10455 if you look a bit further lildmh: The case you describe, which is to have a task wrapped around target, is solved with…
		ABataevAuthorUnsubmitted Not Done Reply Inline Actions Better to add a flag `RequiresOuterTask = D.hasClausesOfKind<OMPDependClause>()` and use it here and where the task is generated (line 10455) ABataev: Better to add a flag `RequiresOuterTask = D.hasClausesOfKind<OMPDependClause>()` and use it…
		lildmhUnsubmitted Done Reply Inline Actions Sure, please accept this patch after that. lildmh: Sure, please accept this patch after that.
InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;		InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;
InputInfo.BasePointersArray =		InputInfo.BasePointersArray =
Address(Info.BasePointersArray, CGM.getPointerAlign());		Address(Info.BasePointersArray, CGM.getPointerAlign());
InputInfo.PointersArray =		InputInfo.PointersArray =
Address(Info.PointersArray, CGM.getPointerAlign());		Address(Info.PointersArray, CGM.getPointerAlign());
InputInfo.SizesArray =		InputInfo.SizesArray =
Address(Info.SizesArray, CGM.getPointerAlign());		Address(Info.SizesArray, CGM.getPointerAlign());
InputInfo.MappersArray = Address(Info.MappersArray, CGM.getPointerAlign());		InputInfo.MappersArray = Address(Info.MappersArray, CGM.getPointerAlign());
MapTypesArray = Info.MapTypesArray;		MapTypesArray = Info.MapTypesArray;
if (D.hasClausesOfKind<OMPDependClause>())		if (HasDependClauses)
CGF.EmitOMPTargetTaskBasedDirective(D, ThenGen, InputInfo);		CGF.EmitOMPTargetTaskBasedDirective(D, ThenGen, InputInfo);
else		else
emitInlinedDirective(CGF, D.getDirectiveKind(), ThenGen);		emitInlinedDirective(CGF, D.getDirectiveKind(), ThenGen);
};		};

if (IfCond) {		if (IfCond) {
emitIfClause(CGF, IfCond, TargetThenGen,		emitIfClause(CGF, IfCond, TargetThenGen,
[](CodeGenFunction &CGF, PrePostActionTy &) {});		[](CodeGenFunction &CGF, PrePostActionTy &) {});
▲ Show 20 Lines • Show All 1,655 Lines • Show Last 20 Lines

clang/test/OpenMP/target_depend_codegen.cpp

Show All 37 Lines
#ifndef HEADER		#ifndef HEADER
#define HEADER		#define HEADER

// CHECK-DAG: [[TT:%.+]] = type { i64, i8 }		// CHECK-DAG: [[TT:%.+]] = type { i64, i8 }
// CHECK-DAG: [[ENTTY:%.+]] = type { i8, i8, i[[SZ:32\|64]], i32, i32 }		// CHECK-DAG: [[ENTTY:%.+]] = type { i8, i8, i[[SZ:32\|64]], i32, i32 }

// TCHECK: [[ENTTY:%.+]] = type { i8, i8, i{{32\|64}}, i32, i32 }		// TCHECK: [[ENTTY:%.+]] = type { i8, i8, i{{32\|64}}, i32, i32 }

// CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 4]		// CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [3 x i64] [i64 0, i64 4, i64 {{16\|12}}]
// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, i64 800]		// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 3]
// CHECK-DAG: @{{.*}} = weak constant i8 0		// CHECK-DAG: @{{.*}} = weak constant i8 0

// TCHECK: @{{.+}} = weak constant [[ENTTY]]		// TCHECK: @{{.+}} = weak constant [[ENTTY]]
// TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]		// TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]		// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]

// Check target registration is registered as a Ctor.		// Check target registration is registered as a Ctor.
// CHECK: appending global [1 x { i32, void (), i8 }] [{ i32, void (), i8 } { i32 0, void ()* @.omp_offloading.requires_reg, i8* null }]		// CHECK: appending global [1 x { i32, void (), i8 }] [{ i32, void (), i8 } { i32 0, void ()* @.omp_offloading.requires_reg, i8* null }]


template<typename tx, typename ty>		template<typename tx, typename ty>
struct TT{		struct TT{
tx X;		tx X;
ty Y;		ty Y;
};		};

		#pragma omp declare mapper(id \
		: TT <long long, char> \
		s) map(s.X, s.Y)
int global;		int global;
extern int global;		extern int global;

// CHECK: define {{.*}}[[FOO:@.+]](		// CHECK: define {{.*}}[[FOO:@.+]](
int foo(int n) {		int foo(int n) {
int a = 0;		int a = 0;
short aa = 0;		short aa = 0;
float b[10];		float b[10];
Show All 25 Lines	#pragma omp target device(global + a) depend(in \
}		}

// CHECK: [[ADD:%.+]] = add nsw i32		// CHECK: [[ADD:%.+]] = add nsw i32
// CHECK: store i32 [[ADD]], i32* [[DEVICE_CAP:%.+]],		// CHECK: store i32 [[ADD]], i32* [[DEVICE_CAP:%.+]],

// CHECK: [[BOOL:%.+]] = icmp ne i32 %{{.+}}, 0		// CHECK: [[BOOL:%.+]] = icmp ne i32 %{{.+}}, 0
// CHECK: br i1 [[BOOL]], label %[[THEN:.+]], label %[[ELSE:.+]]		// CHECK: br i1 [[BOOL]], label %[[THEN:.+]], label %[[ELSE:.+]]
// CHECK: [[THEN]]:		// CHECK: [[THEN]]:
// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [2 x i8], [2 x i8]* [[BP:%.+]], i32 0, i32 0		// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x i8], [3 x i8]* [[BP:%.+]], i32 0, i32 0
// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [2 x i8], [2 x i8]* [[P:%.+]], i32 0, i32 0		// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [3 x i8], [3 x i8]* [[P:%.+]], i32 0, i32 0
		// CHECK-DAG: [[MADDR0:%.+]] = getelementptr inbounds [3 x i8], [3 x i8]* [[M:%.+]], i[[SZ]] 0, i[[SZ]] 0
// CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8 [[BPADDR0]] to i[[SZ]]		// CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8 [[BPADDR0]] to i[[SZ]]
// CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8 [[PADDR0]] to i[[SZ]]		// CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8 [[PADDR0]] to i[[SZ]]
// CHECK-DAG: store i[[SZ]]* [[BP0:%[^,]+]], i[[SZ]]** [[CBPADDR0]]		// CHECK-DAG: store i[[SZ]]* [[BP0:%[^,]+]], i[[SZ]]** [[CBPADDR0]]
// CHECK-DAG: store i[[SZ]]* [[BP0]], i[[SZ]]** [[CPADDR0]]		// CHECK-DAG: store i[[SZ]]* [[BP0]], i[[SZ]]** [[CPADDR0]]
		// CHECK-DAG: store i8* null, i8** [[MADDR0]],

// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [2 x i8], [2 x i8]* [[BP]], i32 0, i32 1		// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [3 x i8], [3 x i8]* [[BP]], i32 0, i32 1
// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [2 x i8], [2 x i8]* [[P]], i32 0, i32 1		// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [3 x i8], [3 x i8]* [[P]], i32 0, i32 1
		// CHECK-DAG: [[MADDR1:%.+]] = getelementptr inbounds [3 x i8], [3 x i8]* [[M]], i[[SZ]] 0, i[[SZ]] 1
// CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]*		// CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]*
// CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]*		// CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]*
// CHECK-DAG: store i[[SZ]] [[BP1:%[^,]+]], i[[SZ]]* [[CBPADDR1]]		// CHECK-DAG: store i[[SZ]] [[BP1:%[^,]+]], i[[SZ]]* [[CBPADDR1]]
// CHECK-DAG: store i[[SZ]] [[BP1]], i[[SZ]]* [[CPADDR1]]		// CHECK-DAG: store i[[SZ]] [[BP1]], i[[SZ]]* [[CPADDR1]]
// CHECK-DAG: getelementptr inbounds [2 x i8], [2 x i8]* [[BP]], i32 0, i32 0		// CHECK-DAG: store i8* null, i8** [[MADDR1]],
// CHECK-DAG: getelementptr inbounds [2 x i8], [2 x i8]* [[P]], i32 0, i32 0
		// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [3 x i8], [3 x i8]* [[BP]], i32 0, i32 2
		// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [3 x i8], [3 x i8]* [[P]], i32 0, i32 2
		// CHECK-DAG: [[MADDR2:%.+]] = getelementptr inbounds [3 x i8], [3 x i8]* [[M]], i[[SZ]] 0, i[[SZ]] 2
		// CHECK-DAG: [[CBPADDR2:%.+]] = bitcast i8 [[BPADDR2]] to [[STRUCT_TT:%.+]]
		// CHECK-DAG: [[CPADDR2:%.+]] = bitcast i8 [[PADDR2]] to [[STRUCT_TT]]
		// CHECK-DAG: store [[STRUCT_TT]]* [[D_ADDR:%.+]], [[STRUCT_TT]]** [[CBPADDR2]]
		// CHECK-DAG: store [[STRUCT_TT]]* [[D_ADDR]], [[STRUCT_TT]]** [[CPADDR2]]
		// CHECK-DAG: store i8* bitcast (void (i8, i8, i8, i64, i64) [[MAPPER_ID:@.+]] to i8), i8* [[MADDR2]],

		// CHECK-DAG: [[BP_START:%.+]] = getelementptr inbounds [3 x i8], [3 x i8]* [[BP]], i32 0, i32 0
		// CHECK-DAG: [[P_START:%.+]] = getelementptr inbounds [3 x i8], [3 x i8]* [[P]], i32 0, i32 0
		// CHECK-DAG: [[M_START:%.+]] = bitcast [3 x i8] [[M]] to i8**
// CHECK: [[GEP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2		// CHECK: [[GEP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
// CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],		// CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],
// CHECK: store i32 [[DEV]], i32* [[GEP]],		// CHECK: store i32 [[DEV]], i32* [[GEP]],
// CHECK: [[DEV1:%.+]] = load i32, i32* [[DEVICE_CAP]],		// CHECK: [[DEV1:%.+]] = load i32, i32* [[DEVICE_CAP]],
// CHECK: [[DEV2:%.+]] = sext i32 [[DEV1]] to i64		// CHECK: [[DEV2:%.+]] = sext i32 [[DEV1]] to i64

// CHECK: [[TASK:%.+]] = call i8* @__kmpc_omp_target_task_alloc(%struct.ident_t* @{{.}}, i32 [[GTID]], i32 1, i[[SZ]] {{120\|68}}, i[[SZ]] {{16\|12}}, i32 (i32, i8)* bitcast (i32 (i32, %{{.+}}) [[TASK_ENTRY1_:@.+]] to i32 (i32, i8)), i64 [[DEV2]])		// CHECK: [[TASK:%.+]] = call i8* @__kmpc_omp_target_task_alloc(%struct.ident_t* @{{.}}, i32 [[GTID]], i32 1, i[[SZ]] {{152\|88}}, i[[SZ]] {{16\|12}}, i32 (i32, i8)* bitcast (i32 (i32, %{{.+}}) [[TASK_ENTRY1_:@.+]] to i32 (i32, i8)), i64 [[DEV2]])
// CHECK: [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY1_:%.+]]*		// CHECK: [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY1_:%.+]]*
		// CHECK: [[BASE:%.+]] = getelementptr inbounds [[TASK_TY1_]], [[TASK_TY1_]]* [[BC_TASK]], i32 0, i32 1
		// CHECK-64: [[BP_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY:%.+]], [[PRIVS_TY:%.+]]* [[BASE]], i32 0, i32 1
		// CHECK-64: [[BP_CAST:%.+]] = bitcast [3 x i8] [[BP_BASE]] to i8*
		// CHECK-64: [[BP_SRC:%.+]] = bitcast i8** [[BP_START]] to i8*
		// CHECK-64: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[BP_CAST]], i8* align 8 [[BP_SRC]], i64 24, i1 false)
		// CHECK-64: [[P_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 2
		// CHECK-64: [[P_CAST:%.+]] = bitcast [3 x i8] [[P_BASE]] to i8*
		// CHECK-64: [[P_SRC:%.+]] = bitcast i8** [[P_START]] to i8*
		// CHECK-64: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[P_CAST]], i8* align 8 [[P_SRC]], i64 24, i1 false)
		// CHECK-64: [[SZ_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 3
		// CHECK-64: [[SZ_CAST:%.+]] = bitcast [3 x i64]* [[SZ_BASE]] to i8*
		// CHECK-64: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[SZ_CAST]], i8* align 8 bitcast ([3 x i64]* [[SIZET]] to i8*), i64 24, i1 false)
		// CHECK-64: [[M_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 4
		// CHECK-64: [[M_CAST:%.+]] = bitcast [3 x i8] [[M_BASE]] to i8*
		// CHECK-64: [[M_SRC:%.+]] = bitcast i8** [[M_START]] to i8*
		// CHECK-64: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[M_CAST]], i8* align 8 [[M_SRC]], i64 24, i1 false)
		// CHECK-32: [[SZ_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY:%.+]], [[PRIVS_TY:%.+]]* [[BASE]], i32 0, i32 0
		// CHECK-32: [[SZ_CAST:%.+]] = bitcast [3 x i64]* [[SZ_BASE]] to i8*
		// CHECK-32: call void @llvm.memcpy.p0i8.p0i8.i32(i8* align 4 [[SZ_CAST]], i8* align 4 bitcast ([3 x i64]* [[SIZET]] to i8*), i32 24, i1 false)
		// CHECK-32: [[BP_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 3
		// CHECK-32: [[BP_CAST:%.+]] = bitcast [3 x i8] [[BP_BASE]] to i8*
		// CHECK-32: [[BP_SRC:%.+]] = bitcast i8** [[BP_START]] to i8*
		// CHECK-32: call void @llvm.memcpy.p0i8.p0i8.i32(i8* align 4 [[BP_CAST]], i8* align 4 [[BP_SRC]], i32 12, i1 false)
		// CHECK-32: [[P_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 4
		// CHECK-32: [[P_CAST:%.+]] = bitcast [3 x i8] [[P_BASE]] to i8*
		// CHECK-32: [[P_SRC:%.+]] = bitcast i8** [[P_START]] to i8*
		// CHECK-32: call void @llvm.memcpy.p0i8.p0i8.i32(i8* align 4 [[P_CAST]], i8* align 4 [[P_SRC]], i32 12, i1 false)
		// CHECK-32: [[M_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 5
		// CHECK-32: [[M_CAST:%.+]] = bitcast [3 x i8] [[M_BASE]] to i8*
		// CHECK-32: [[M_SRC:%.+]] = bitcast i8** [[M_START]] to i8*
		// CHECK-32: call void @llvm.memcpy.p0i8.p0i8.i32(i8* align 4 [[M_CAST]], i8* align 4 [[M_SRC]], i32 12, i1 false)
// CHECK: getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START:%.+]], i[[SZ]] 1		// CHECK: getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START:%.+]], i[[SZ]] 1
// CHECK: getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START]], i[[SZ]] 2		// CHECK: getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START]], i[[SZ]] 2
// CHECK: [[DEP:%.+]] = bitcast %struct.kmp_depend_info* [[DEP_START]] to i8*		// CHECK: [[DEP:%.+]] = bitcast %struct.kmp_depend_info* [[DEP_START]] to i8*
// CHECK: call i32 @__kmpc_omp_task_with_deps(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]], i32 3, i8* [[DEP]], i32 0, i8* null)		// CHECK: call i32 @__kmpc_omp_task_with_deps(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]], i32 3, i8* [[DEP]], i32 0, i8* null)
// CHECK: br label %[[EXIT:.+]]		// CHECK: br label %[[EXIT:.+]]

// CHECK: [[ELSE]]:		// CHECK: [[ELSE]]:
// CHECK-NOT: getelementptr inbounds [2 x i8], [2 x i8]*		// CHECK-NOT: getelementptr inbounds [2 x i8], [2 x i8]*
// CHECK: [[GEP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2		// CHECK: [[GEP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
// CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],		// CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],
// CHECK: store i32 [[DEV]], i32* [[GEP]],		// CHECK: store i32 [[DEV]], i32* [[GEP]],
// CHECK: [[DEV1:%.+]] = load i32, i32* [[DEVICE_CAP]],		// CHECK: [[DEV1:%.+]] = load i32, i32* [[DEVICE_CAP]],
// CHECK: [[DEV2:%.+]] = sext i32 [[DEV1]] to i64		// CHECK: [[DEV2:%.+]] = sext i32 [[DEV1]] to i64

// CHECK: [[TASK:%.+]] = call i8* @__kmpc_omp_target_task_alloc(%struct.ident_t* @1, i32 [[GTID]], i32 1, i[[SZ]] {{56\|28}}, i[[SZ]] {{16\|12}}, i32 (i32, i8) bitcast (i32 (i32, %{{.+}}) [[TASK_ENTRY1__:@.+]] to i32 (i32, i8)), i64 [[DEV2]])		// CHECK: [[TASK:%.+]] = call i8* @__kmpc_omp_target_task_alloc(%struct.ident_t* @1, i32 [[GTID]], i32 1, i[[SZ]] {{56\|28}}, i[[SZ]] {{16\|12}}, i32 (i32, i8) bitcast (i32 (i32, %{{.+}}) [[TASK_ENTRY1__:@.+]] to i32 (i32, i8)), i64 [[DEV2]])
// CHECK: [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY1__:%.+]]*		// CHECK: [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY1__:%.+]]*
// CHECK: getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START:%.+]], i[[SZ]] 1		// CHECK: getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START:%.+]], i[[SZ]] 1
// CHECK: getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START]], i[[SZ]] 2		// CHECK: getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START]], i[[SZ]] 2
// CHECK: [[DEP:%.+]] = bitcast %struct.kmp_depend_info* [[DEP_START]] to i8*		// CHECK: [[DEP:%.+]] = bitcast %struct.kmp_depend_info* [[DEP_START]] to i8*
// CHECK: call i32 @__kmpc_omp_task_with_deps(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]], i32 3, i8* [[DEP]], i32 0, i8* null)		// CHECK: call i32 @__kmpc_omp_task_with_deps(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]], i32 3, i8* [[DEP]], i32 0, i8* null)
// CHECK: br label %[[EXIT:.+]]		// CHECK: br label %[[EXIT:.+]]
// CHECK: [[EXIT]]:		// CHECK: [[EXIT]]:

#pragma omp target device(global + a) nowait depend(inout \		#pragma omp target device(global + a) nowait depend(inout \
: global, a, bn) if (a)		: global, a, bn) if (a) map(mapper(id), tofrom \
		: d)
{		{
static int local1;		static int local1;
*plocal = global;		*plocal = global;
local1 = global;		local1 = global;
}		}

// CHECK: [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @1, i32 [[GTID]], i32 1, i[[SZ]] {{48\|24}}, i[[SZ]] 4, i32 (i32, i8) bitcast (i32 (i32, %{{.+}}) [[TASK_ENTRY2:@.+]] to i32 (i32, i8)))		// CHECK: [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @1, i32 [[GTID]], i32 1, i[[SZ]] {{48\|24}}, i[[SZ]] 4, i32 (i32, i8) bitcast (i32 (i32, %{{.+}}) [[TASK_ENTRY2:@.+]] to i32 (i32, i8)))
// CHECK: [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY2:%.+]]*		// CHECK: [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY2:%.+]]*
Show All 27 Lines
// CHECK: [[FAIL]]		// CHECK: [[FAIL]]
// CHECK: call void [[HVT0]]()		// CHECK: call void [[HVT0]]()
// CHECK-NEXT: br label %[[END]]		// CHECK-NEXT: br label %[[END]]
// CHECK: [[END]]		// CHECK: [[END]]
// CHECK: ret i32 0		// CHECK: ret i32 0

// CHECK: define internal void [[HVT1:@.+]](i[[SZ]]* %{{.+}}, i[[SZ]] %{{.+}})		// CHECK: define internal void [[HVT1:@.+]](i[[SZ]]* %{{.+}}, i[[SZ]] %{{.+}})

// CHECK: define internal{{.}} i32 [[TASK_ENTRY1_]](i32{{.}}, [[TASK_TY1_]]* noalias %1)		// CHECK: define internal void [[MAPPER_ID]](i8* %{{.+}}, i8* %{{.+}}, i8* %{{.+}}, i64 %{{.+}}, i64 %{{.+}})
// CHECK: call void (i8*, ...) %
// CHECK: [[SZT:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 0		// CHECK: define internal{{.}} i32 [[TASK_ENTRY1_]](i32{{.}}, [[TASK_TY1_]]* noalias %{{.+}})
		// CHECK: call void (i8, ...) %{{.+}}(i8 %{{.+}}, i[[SZ]]* %{{.+}}, i32 %{{.+}}, [3 x i8]* [[BPTR_ADDR:%.+]], [3 x i8]* [[PTR_ADDR:%.+]], [3 x i64]** [[SZ_ADDR:%.+]], [3 x i8]* [[M_ADDR:%.+]])
		// CHECK: [[BPTR_REF:%.+]] = load [3 x i8], [3 x i8]* [[BPTR_ADDR]],
		// CHECK: [[PTR_REF:%.+]] = load [3 x i8], [3 x i8]* [[PTR_ADDR]],
		// CHECK: [[SZ_REF:%.+]] = load [3 x i64], [3 x i64]* [[SZ_ADDR]],
		// CHECK: [[M_REF:%.+]] = load [3 x i8], [3 x i8]* [[M_ADDR]],
		// CHECK: [[BPR:%.+]] = getelementptr inbounds [3 x i8], [3 x i8]* [[BPTR_REF]], i[[SZ]] 0, i[[SZ]] 0
		// CHECK: [[PR:%.+]] = getelementptr inbounds [3 x i8], [3 x i8]* [[PTR_REF]], i[[SZ]] 0, i[[SZ]] 0
		// CHECK: [[SZT:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[SZ_REF]], i[[SZ]] 0, i[[SZ]] 0
		// CHECK: [[M:%.+]] = getelementptr inbounds [3 x i8], [3 x i8]* [[M_REF]], i[[SZ]] 0, i[[SZ]] 0
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2		// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
// CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],		// CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],
// CHECK: [[DEVICE:%.+]] = sext i32 [[DEV]] to i64		// CHECK: [[DEVICE:%.+]] = sext i32 [[DEV]] to i64
// CHECK: [[RET:%.+]] = call i32 @__tgt_target_nowait_mapper(i64 [[DEVICE]], i8* @{{[^,]+}}, i32 2, i8 [[BPR:%[^,]+]], i8 [[PR:%[^,]+]], i64* [[SZT]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPT]], i32 0, i32 0), i8** [[M:%[^,]+]])		// CHECK: [[RET:%.+]] = call i32 @__tgt_target_nowait_mapper(i64 [[DEVICE]], i8* @{{[^,]+}}, i32 3, i8 [[BPR]], i8 [[PR]], i64* [[SZT]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT]], i32 0, i32 0), i8** [[M]])

// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0		// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]		// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
// CHECK: [[FAIL]]		// CHECK: [[FAIL]]
// CHECK: [[BP0:%.+]] = load i[[SZ]], i[[SZ]]* %		// CHECK: [[BP0:%.+]] = load i[[SZ]], i[[SZ]]* %
// CHECK: [[BP1_I32:%.+]] = load i32, i32* @		// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*		// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],		// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
▲ Show 20 Lines • Show All 42 Lines • Show Last 20 Lines