Skip to content

Commit d5ce823

Browse files
alexey-bataevalexey-bataev
committed
[OpenMP 5.0] Fix user-defined mapper privatization in tasks
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. Differential Revision: https://reviews.llvm.org/D84470
1 parent 5b533d6 commit d5ce823

File tree

2 files changed

+111
-32
lines changed

2 files changed

+111
-32
lines changed

‎clang/lib/CodeGen/CGOpenMPRuntime.cpp

Lines changed: 37 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -3784,9 +3784,9 @@ static void emitPrivatesInit(CodeGenFunction &CGF,
37843784
bool IsTargetTask =
37853785
isOpenMPTargetDataManagementDirective(D.getDirectiveKind()) ||
37863786
isOpenMPTargetExecutionDirective(D.getDirectiveKind());
3787-
// For target-based directives skip 3 firstprivate arrays BasePointersArray,
3788-
// PointersArray and SizesArray. The original variables for these arrays are
3789-
// not captured and we get their addresses explicitly.
3787+
// For target-based directives skip 4 firstprivate arrays BasePointersArray,
3788+
// PointersArray, SizesArray, and MappersArray. The original variables for
3789+
// these arrays are not captured and we get their addresses explicitly.
37903790
if ((!IsTargetTask && !Data.FirstprivateVars.empty() && ForDup) ||
37913791
(IsTargetTask && KmpTaskSharedsPtr.isValid())) {
37923792
SrcBase = CGF.MakeAddrLValue(
@@ -3809,7 +3809,7 @@ static void emitPrivatesInit(CodeGenFunction &CGF,
38093809
if (const VarDecl *Elem = Pair.second.PrivateElemInit) {
38103810
const VarDecl *OriginalVD = Pair.second.Original;
38113811
// Check if the variable is the target-based BasePointersArray,
3812-
// PointersArray or SizesArray.
3812+
// PointersArray, SizesArray, or MappersArray.
38133813
LValue SharedRefLValue;
38143814
QualType Type = PrivateLValue.getType();
38153815
const FieldDecl *SharedField = CapturesInfo.lookup(OriginalVD);
@@ -8866,6 +8866,17 @@ emitOffloadingArrays(CodeGenFunction &CGF,
88668866
}
88678867
}
88688868

8869+
namespace {
8870+
/// Additional arguments for emitOffloadingArraysArgument function.
8871+
struct ArgumentsOptions {
8872+
bool ForEndCall = false;
8873+
bool IsTask = false;
8874+
ArgumentsOptions() = default;
8875+
ArgumentsOptions(bool ForEndCall, bool IsTask)
8876+
: ForEndCall(ForEndCall), IsTask(IsTask) {}
8877+
};
8878+
} // namespace
8879+
88698880
/// Emit the arguments to be passed to the runtime library based on the
88708881
/// arrays of base pointers, pointers, sizes, map types, and mappers. If
88718882
/// ForEndCall, emit map types to be passed for the end of the region instead of
@@ -8874,8 +8885,9 @@ static void emitOffloadingArraysArgument(
88748885
CodeGenFunction &CGF, llvm::Value *&BasePointersArrayArg,
88758886
llvm::Value *&PointersArrayArg, llvm::Value *&SizesArrayArg,
88768887
llvm::Value *&MapTypesArrayArg, llvm::Value *&MappersArrayArg,
8877-
CGOpenMPRuntime::TargetDataInfo &Info, bool ForEndCall = false) {
8878-
assert((!ForEndCall || Info.separateBeginEndCalls()) &&
8888+
CGOpenMPRuntime::TargetDataInfo &Info,
8889+
const ArgumentsOptions &Options = ArgumentsOptions()) {
8890+
assert((!Options.ForEndCall || Info.separateBeginEndCalls()) &&
88798891
"expected region end call to runtime only when end call is separate");
88808892
CodeGenModule &CGM = CGF.CGM;
88818893
if (Info.NumberOfPtrs) {
@@ -8893,14 +8905,17 @@ static void emitOffloadingArraysArgument(
88938905
/*Idx0=*/0, /*Idx1=*/0);
88948906
MapTypesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
88958907
llvm::ArrayType::get(CGM.Int64Ty, Info.NumberOfPtrs),
8896-
ForEndCall && Info.MapTypesArrayEnd ? Info.MapTypesArrayEnd
8897-
: Info.MapTypesArray,
8908+
Options.ForEndCall && Info.MapTypesArrayEnd ? Info.MapTypesArrayEnd
8909+
: Info.MapTypesArray,
88988910
/*Idx0=*/0,
88998911
/*Idx1=*/0);
8900-
MappersArrayArg =
8901-
Info.HasMapper
8902-
? CGF.Builder.CreatePointerCast(Info.MappersArray, CGM.VoidPtrPtrTy)
8903-
: llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
8912+
// Always emit the mapper array address in case of a target task for
8913+
// privatization.
8914+
if (!Options.IsTask && !Info.HasMapper)
8915+
MappersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
8916+
else
8917+
MappersArrayArg =
8918+
CGF.Builder.CreatePointerCast(Info.MappersArray, CGM.VoidPtrPtrTy);
89048919
} else {
89058920
BasePointersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
89068921
PointersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
@@ -9648,9 +9663,11 @@ void CGOpenMPRuntime::emitTargetCall(
96489663
TargetDataInfo Info;
96499664
// Fill up the arrays and create the arguments.
96509665
emitOffloadingArrays(CGF, CombinedInfo, Info);
9666+
bool HasDependClauses = D.hasClausesOfKind<OMPDependClause>();
96519667
emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
96529668
Info.PointersArray, Info.SizesArray,
9653-
Info.MapTypesArray, Info.MappersArray, Info);
9669+
Info.MapTypesArray, Info.MappersArray, Info,
9670+
{/*ForEndTask=*/false, HasDependClauses});
96549671
InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;
96559672
InputInfo.BasePointersArray =
96569673
Address(Info.BasePointersArray, CGM.getPointerAlign());
@@ -10261,7 +10278,7 @@ void CGOpenMPRuntime::emitTargetDataCalls(
1026110278
llvm::Value *MappersArrayArg = nullptr;
1026210279
emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
1026310280
SizesArrayArg, MapTypesArrayArg,
10264-
MappersArrayArg, Info, /*ForEndCall=*/false);
10281+
MappersArrayArg, Info);
1026510282

1026610283
// Emit device ID if any.
1026710284
llvm::Value *DeviceID = nullptr;
@@ -10301,7 +10318,8 @@ void CGOpenMPRuntime::emitTargetDataCalls(
1030110318
llvm::Value *MappersArrayArg = nullptr;
1030210319
emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
1030310320
SizesArrayArg, MapTypesArrayArg,
10304-
MappersArrayArg, Info, /*ForEndCall=*/true);
10321+
MappersArrayArg, Info,
10322+
{/*ForEndCall=*/true, /*IsTask=*/false});
1030510323

1030610324
// Emit device ID if any.
1030710325
llvm::Value *DeviceID = nullptr;
@@ -10499,9 +10517,11 @@ void CGOpenMPRuntime::emitTargetDataStandAloneCall(
1049910517
TargetDataInfo Info;
1050010518
// Fill up the arrays and create the arguments.
1050110519
emitOffloadingArrays(CGF, CombinedInfo, Info);
10520+
bool HasDependClauses = D.hasClausesOfKind<OMPDependClause>();
1050210521
emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
1050310522
Info.PointersArray, Info.SizesArray,
10504-
Info.MapTypesArray, Info.MappersArray, Info);
10523+
Info.MapTypesArray, Info.MappersArray, Info,
10524+
{/*ForEndTask=*/false, HasDependClauses});
1050510525
InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;
1050610526
InputInfo.BasePointersArray =
1050710527
Address(Info.BasePointersArray, CGM.getPointerAlign());
@@ -10511,7 +10531,7 @@ void CGOpenMPRuntime::emitTargetDataStandAloneCall(
1051110531
Address(Info.SizesArray, CGM.getPointerAlign());
1051210532
InputInfo.MappersArray = Address(Info.MappersArray, CGM.getPointerAlign());
1051310533
MapTypesArray = Info.MapTypesArray;
10514-
if (D.hasClausesOfKind<OMPDependClause>())
10534+
if (HasDependClauses)
1051510535
CGF.EmitOMPTargetTaskBasedDirective(D, ThenGen, InputInfo);
1051610536
else
1051710537
emitInlinedDirective(CGF, D.getDirectiveKind(), ThenGen);

‎clang/test/OpenMP/target_depend_codegen.cpp

Lines changed: 74 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -43,8 +43,8 @@
4343

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

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

5050
// TCHECK: @{{.+}} = weak constant [[ENTTY]]
@@ -61,6 +61,9 @@ struct TT{
6161
ty Y;
6262
};
6363

64+
#pragma omp declare mapper(id \
65+
: TT <long long, char> \
66+
s) map(s.X, s.Y)
6467
int global;
6568
extern int global;
6669

@@ -102,29 +105,75 @@ int foo(int n) {
102105
// CHECK: [[BOOL:%.+]] = icmp ne i32 %{{.+}}, 0
103106
// CHECK: br i1 [[BOOL]], label %[[THEN:.+]], label %[[ELSE:.+]]
104107
// CHECK: [[THEN]]:
105-
// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP:%.+]], i32 0, i32 0
106-
// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P:%.+]], i32 0, i32 0
108+
// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP:%.+]], i32 0, i32 0
109+
// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P:%.+]], i32 0, i32 0
110+
// CHECK-DAG: [[MADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[M:%.+]], i[[SZ]] 0, i[[SZ]] 0
107111
// CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to i[[SZ]]**
108112
// CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to i[[SZ]]**
109113
// CHECK-DAG: store i[[SZ]]* [[BP0:%[^,]+]], i[[SZ]]** [[CBPADDR0]]
110114
// CHECK-DAG: store i[[SZ]]* [[BP0]], i[[SZ]]** [[CPADDR0]]
115+
// CHECK-DAG: store i8* null, i8** [[MADDR0]],
111116

112-
// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP]], i32 0, i32 1
113-
// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P]], i32 0, i32 1
117+
// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 1
118+
// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 1
119+
// CHECK-DAG: [[MADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[M]], i[[SZ]] 0, i[[SZ]] 1
114120
// CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]*
115121
// CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]*
116122
// CHECK-DAG: store i[[SZ]] [[BP1:%[^,]+]], i[[SZ]]* [[CBPADDR1]]
117123
// CHECK-DAG: store i[[SZ]] [[BP1]], i[[SZ]]* [[CPADDR1]]
118-
// CHECK-DAG: getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP]], i32 0, i32 0
119-
// CHECK-DAG: getelementptr inbounds [2 x i8*], [2 x i8*]* [[P]], i32 0, i32 0
124+
// CHECK-DAG: store i8* null, i8** [[MADDR1]],
125+
126+
// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 2
127+
// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 2
128+
// CHECK-DAG: [[MADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[M]], i[[SZ]] 0, i[[SZ]] 2
129+
// CHECK-DAG: [[CBPADDR2:%.+]] = bitcast i8** [[BPADDR2]] to [[STRUCT_TT:%.+]]**
130+
// CHECK-DAG: [[CPADDR2:%.+]] = bitcast i8** [[PADDR2]] to [[STRUCT_TT]]**
131+
// CHECK-DAG: store [[STRUCT_TT]]* [[D_ADDR:%.+]], [[STRUCT_TT]]** [[CBPADDR2]]
132+
// CHECK-DAG: store [[STRUCT_TT]]* [[D_ADDR]], [[STRUCT_TT]]** [[CPADDR2]]
133+
// CHECK-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MAPPER_ID:@.+]] to i8*), i8** [[MADDR2]],
134+
135+
// CHECK-DAG: [[BP_START:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 0
136+
// CHECK-DAG: [[P_START:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 0
137+
// CHECK-DAG: [[M_START:%.+]] = bitcast [3 x i8*]* [[M]] to i8**
120138
// CHECK: [[GEP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
121139
// CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],
122140
// CHECK: store i32 [[DEV]], i32* [[GEP]],
123141
// CHECK: [[DEV1:%.+]] = load i32, i32* [[DEVICE_CAP]],
124142
// CHECK: [[DEV2:%.+]] = sext i32 [[DEV1]] to i64
125143

126-
// 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]])
144+
// 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]])
127145
// CHECK: [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY1_:%.+]]*
146+
// CHECK: [[BASE:%.+]] = getelementptr inbounds [[TASK_TY1_]], [[TASK_TY1_]]* [[BC_TASK]], i32 0, i32 1
147+
// CHECK-64: [[BP_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY:%.+]], [[PRIVS_TY:%.+]]* [[BASE]], i32 0, i32 1
148+
// CHECK-64: [[BP_CAST:%.+]] = bitcast [3 x i8*]* [[BP_BASE]] to i8*
149+
// CHECK-64: [[BP_SRC:%.+]] = bitcast i8** [[BP_START]] to i8*
150+
// CHECK-64: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[BP_CAST]], i8* align 8 [[BP_SRC]], i64 24, i1 false)
151+
// CHECK-64: [[P_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 2
152+
// CHECK-64: [[P_CAST:%.+]] = bitcast [3 x i8*]* [[P_BASE]] to i8*
153+
// CHECK-64: [[P_SRC:%.+]] = bitcast i8** [[P_START]] to i8*
154+
// CHECK-64: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[P_CAST]], i8* align 8 [[P_SRC]], i64 24, i1 false)
155+
// CHECK-64: [[SZ_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 3
156+
// CHECK-64: [[SZ_CAST:%.+]] = bitcast [3 x i64]* [[SZ_BASE]] to i8*
157+
// 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)
158+
// CHECK-64: [[M_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 4
159+
// CHECK-64: [[M_CAST:%.+]] = bitcast [3 x i8*]* [[M_BASE]] to i8*
160+
// CHECK-64: [[M_SRC:%.+]] = bitcast i8** [[M_START]] to i8*
161+
// CHECK-64: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[M_CAST]], i8* align 8 [[M_SRC]], i64 24, i1 false)
162+
// CHECK-32: [[SZ_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY:%.+]], [[PRIVS_TY:%.+]]* [[BASE]], i32 0, i32 0
163+
// CHECK-32: [[SZ_CAST:%.+]] = bitcast [3 x i64]* [[SZ_BASE]] to i8*
164+
// 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)
165+
// CHECK-32: [[BP_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 3
166+
// CHECK-32: [[BP_CAST:%.+]] = bitcast [3 x i8*]* [[BP_BASE]] to i8*
167+
// CHECK-32: [[BP_SRC:%.+]] = bitcast i8** [[BP_START]] to i8*
168+
// CHECK-32: call void @llvm.memcpy.p0i8.p0i8.i32(i8* align 4 [[BP_CAST]], i8* align 4 [[BP_SRC]], i32 12, i1 false)
169+
// CHECK-32: [[P_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 4
170+
// CHECK-32: [[P_CAST:%.+]] = bitcast [3 x i8*]* [[P_BASE]] to i8*
171+
// CHECK-32: [[P_SRC:%.+]] = bitcast i8** [[P_START]] to i8*
172+
// CHECK-32: call void @llvm.memcpy.p0i8.p0i8.i32(i8* align 4 [[P_CAST]], i8* align 4 [[P_SRC]], i32 12, i1 false)
173+
// CHECK-32: [[M_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 5
174+
// CHECK-32: [[M_CAST:%.+]] = bitcast [3 x i8*]* [[M_BASE]] to i8*
175+
// CHECK-32: [[M_SRC:%.+]] = bitcast i8** [[M_START]] to i8*
176+
// CHECK-32: call void @llvm.memcpy.p0i8.p0i8.i32(i8* align 4 [[M_CAST]], i8* align 4 [[M_SRC]], i32 12, i1 false)
128177
// CHECK: getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START:%.+]], i[[SZ]] 1
129178
// CHECK: getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START]], i[[SZ]] 2
130179
// CHECK: [[DEP:%.+]] = bitcast %struct.kmp_depend_info* [[DEP_START]] to i8*
@@ -148,8 +197,9 @@ int foo(int n) {
148197
// CHECK: br label %[[EXIT:.+]]
149198
// CHECK: [[EXIT]]:
150199

151-
#pragma omp target device(global + a) nowait depend(inout \
152-
: global, a, bn) if (a)
200+
#pragma omp target device(global + a) nowait depend(inout \
201+
: global, a, bn) if (a) map(mapper(id), tofrom \
202+
: d)
153203
{
154204
static int local1;
155205
*plocal = global;
@@ -193,13 +243,22 @@ int foo(int n) {
193243

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

196-
// CHECK: define internal{{.*}} i32 [[TASK_ENTRY1_]](i32{{.*}}, [[TASK_TY1_]]* noalias %1)
197-
// CHECK: call void (i8*, ...) %
198-
// CHECK: [[SZT:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 0
246+
// CHECK: define internal void [[MAPPER_ID]](i8* %{{.+}}, i8* %{{.+}}, i8* %{{.+}}, i64 %{{.+}}, i64 %{{.+}})
247+
248+
// CHECK: define internal{{.*}} i32 [[TASK_ENTRY1_]](i32{{.*}}, [[TASK_TY1_]]* noalias %{{.+}})
249+
// 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:%.+]])
250+
// CHECK: [[BPTR_REF:%.+]] = load [3 x i8*]*, [3 x i8*]** [[BPTR_ADDR]],
251+
// CHECK: [[PTR_REF:%.+]] = load [3 x i8*]*, [3 x i8*]** [[PTR_ADDR]],
252+
// CHECK: [[SZ_REF:%.+]] = load [3 x i64]*, [3 x i64]** [[SZ_ADDR]],
253+
// CHECK: [[M_REF:%.+]] = load [3 x i8*]*, [3 x i8*]** [[M_ADDR]],
254+
// CHECK: [[BPR:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPTR_REF]], i[[SZ]] 0, i[[SZ]] 0
255+
// CHECK: [[PR:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTR_REF]], i[[SZ]] 0, i[[SZ]] 0
256+
// CHECK: [[SZT:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[SZ_REF]], i[[SZ]] 0, i[[SZ]] 0
257+
// CHECK: [[M:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[M_REF]], i[[SZ]] 0, i[[SZ]] 0
199258
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
200259
// CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],
201260
// CHECK: [[DEVICE:%.+]] = sext i32 [[DEV]] to i64
202-
// 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:%[^,]+]])
261+
// 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]])
203262

204263
// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
205264
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]

0 commit comments

Comments
 (0)