jyu2 created this revision.
jyu2 added reviewers: ABataev, mikerice, jdoerfert.
Herald added subscribers: sunshaoce, guansong, yaxunl.
Herald added a project: All.
jyu2 requested review of this revision.
Herald added subscribers: openmp-commits, cfe-commits, sstefan1.
Herald added projects: clang, OpenMP.
My change of D14093 <https://reviews.llvm.org/D14093> is only fixed problem for
"pragma target data".
The problem still here for "pragma target"
what I am missing is:
When processing "pragma target data", the VD is passed when call to
emitCombinedEntry, so check VD is null as map for this pointer.
But when processing "pragma target" the VD is passed as nullptr, so
check VD is null is not working.
To fix this I add a new parameter IsMapThis. During the call to
emitCombinedEntry passes true if it is capturing this pointer and use
that instead check of "!VD".
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D146000
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_map_member_expr_codegen.cpp
openmp/libomptarget/test/mapping/target_map_for_member_data.cpp
Index: openmp/libomptarget/test/mapping/target_map_for_member_data.cpp
===================================================================
--- openmp/libomptarget/test/mapping/target_map_for_member_data.cpp
+++ openmp/libomptarget/test/mapping/target_map_for_member_data.cpp
@@ -68,11 +68,13 @@
auto Asize = 4;
auto Csize = 4;
-#pragma omp target data map(to : d.A) map(from : d.C)
+#pragma omp target data map(from : d.C)
{
#pragma omp target teams firstprivate(Csize)
d.C = 1;
}
+#pragma omp target map(from : d.A)
+ d.A = 3;
}
};
@@ -91,4 +93,6 @@
z.bar(d);
// CHECK 1
printf("%d\n", d.C);
+ // CHECK 3
+ printf("%d\n", d.A);
}
Index: clang/test/OpenMP/target_map_member_expr_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_map_member_expr_codegen.cpp
+++ clang/test/OpenMP/target_map_member_expr_codegen.cpp
@@ -56,6 +56,9 @@
for (int i = 0; i < Csize; ++i)
d.C[i] = 1;
}
+ #pragma omp target map(from:d.C[0:Csize])
+ for (int i = 0; i < Csize; ++i)
+ d.C[i] = 1;
}
};
@@ -66,9 +69,11 @@
descriptor<float> d;
c.bar(d);
}
+
// CHECK: @.offload_sizes = private unnamed_addr constant [4 x i64] [i64 12, i64 4, i64 4, i64 4]
// CHECK-NOT: @.offload_sizes = private unnamed_addr constant [4 x i64] [i64 0, i64 4, i64 4, i64 4]
-
+// CHECK: @.offload_sizes.4 = private unnamed_addr constant [3 x i64] [i64 4, i64 0, i64 0]
+// CHECK-NOT: @.offload_sizes.4 = private unnamed_addr constant [3 x i64] [i64 4, i64 1, i64 0]
// CHECK-LABEL: define {{[^@]+}}@_Z3foov
// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
// CHECK-NEXT: entry:
@@ -189,6 +194,12 @@
// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS9:%.*]] = alloca [2 x ptr], align 8
// CHECK-NEXT: [[DOTOFFLOAD_PTRS10:%.*]] = alloca [2 x ptr], align 8
// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS11:%.*]] = alloca [2 x ptr], align 8
+// CHECK-NEXT: [[_TMP12:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[CSIZE_CASTED13:%.*]] = alloca i64, align 8
+// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS18:%.*]] = alloca [3 x ptr], align 8
+// CHECK-NEXT: [[DOTOFFLOAD_PTRS19:%.*]] = alloca [3 x ptr], align 8
+// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS20:%.*]] = alloca [3 x ptr], align 8
+// CHECK-NEXT: [[DOTOFFLOAD_SIZES21:%.*]] = alloca [3 x i64], align 8
// CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 8
// CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
@@ -323,6 +334,87 @@
// CHECK-NEXT: [[TMP71:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NEXT: [[TMP72:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 3, ptr [[TMP70]], ptr [[TMP71]], ptr [[TMP72]], ptr @.offload_maptypes.1, ptr null, ptr null)
+// CHECK-NEXT: [[TMP73:%.*]] = load ptr, ptr [[D_ADDR]], align 8
+// CHECK-NEXT: store ptr [[TMP73]], ptr [[_TMP12]], align 8
+// CHECK-NEXT: [[TMP74:%.*]] = load i32, ptr [[CSIZE]], align 4
+// CHECK-NEXT: store i32 [[TMP74]], ptr [[CSIZE_CASTED13]], align 4
+// CHECK-NEXT: [[TMP75:%.*]] = load i64, ptr [[CSIZE_CASTED13]], align 8
+// CHECK-NEXT: [[TMP76:%.*]] = load ptr, ptr [[_TMP12]], align 8
+// CHECK-NEXT: [[TMP77:%.*]] = load ptr, ptr [[_TMP12]], align 8
+// CHECK-NEXT: [[TMP78:%.*]] = load ptr, ptr [[_TMP12]], align 8
+// CHECK-NEXT: [[C14:%.*]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], ptr [[TMP78]], i32 0, i32 1
+// CHECK-NEXT: [[TMP79:%.*]] = load ptr, ptr [[_TMP12]], align 8
+// CHECK-NEXT: [[C15:%.*]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], ptr [[TMP79]], i32 0, i32 1
+// CHECK-NEXT: [[TMP80:%.*]] = load ptr, ptr [[C15]], align 8
+// CHECK-NEXT: [[ARRAYIDX16:%.*]] = getelementptr inbounds float, ptr [[TMP80]], i64 0
+// CHECK-NEXT: [[TMP81:%.*]] = load i32, ptr [[CSIZE]], align 4
+// CHECK-NEXT: [[CONV17:%.*]] = zext i32 [[TMP81]] to i64
+// CHECK-NEXT: [[TMP82:%.*]] = mul nuw i64 [[CONV17]], 4
+// CHECK-NEXT: [[TMP83:%.*]] = getelementptr ptr, ptr [[C14]], i32 1
+// CHECK-NEXT: [[TMP84:%.*]] = ptrtoint ptr [[TMP83]] to i64
+// CHECK-NEXT: [[TMP85:%.*]] = ptrtoint ptr [[C14]] to i64
+// CHECK-NEXT: [[TMP86:%.*]] = sub i64 [[TMP84]], [[TMP85]]
+// CHECK-NEXT: [[TMP87:%.*]] = sdiv exact i64 [[TMP86]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
+// CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES21]], ptr align 8 @.offload_sizes.4, i64 24, i1 false)
+// CHECK-NEXT: [[TMP88:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS18]], i32 0, i32 0
+// CHECK-NEXT: store i64 [[TMP75]], ptr [[TMP88]], align 8
+// CHECK-NEXT: [[TMP89:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS19]], i32 0, i32 0
+// CHECK-NEXT: store i64 [[TMP75]], ptr [[TMP89]], align 8
+// CHECK-NEXT: [[TMP90:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS20]], i64 0, i64 0
+// CHECK-NEXT: store ptr null, ptr [[TMP90]], align 8
+// CHECK-NEXT: [[TMP91:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS18]], i32 0, i32 1
+// CHECK-NEXT: store ptr [[TMP77]], ptr [[TMP91]], align 8
+// CHECK-NEXT: [[TMP92:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS19]], i32 0, i32 1
+// CHECK-NEXT: store ptr [[C14]], ptr [[TMP92]], align 8
+// CHECK-NEXT: [[TMP93:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES21]], i32 0, i32 1
+// CHECK-NEXT: store i64 [[TMP87]], ptr [[TMP93]], align 8
+// CHECK-NEXT: [[TMP94:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS20]], i64 0, i64 1
+// CHECK-NEXT: store ptr null, ptr [[TMP94]], align 8
+// CHECK-NEXT: [[TMP95:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS18]], i32 0, i32 2
+// CHECK-NEXT: store ptr [[C14]], ptr [[TMP95]], align 8
+// CHECK-NEXT: [[TMP96:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS19]], i32 0, i32 2
+// CHECK-NEXT: store ptr [[ARRAYIDX16]], ptr [[TMP96]], align 8
+// CHECK-NEXT: [[TMP97:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES21]], i32 0, i32 2
+// CHECK-NEXT: store i64 [[TMP82]], ptr [[TMP97]], align 8
+// CHECK-NEXT: [[TMP98:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS20]], i64 0, i64 2
+// CHECK-NEXT: store ptr null, ptr [[TMP98]], align 8
+// CHECK-NEXT: [[TMP99:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS18]], i32 0, i32 0
+// CHECK-NEXT: [[TMP100:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS19]], i32 0, i32 0
+// CHECK-NEXT: [[TMP101:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES21]], i32 0, i32 0
+// CHECK-NEXT: [[KERNEL_ARGS22:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8
+// CHECK-NEXT: [[TMP102:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS22]], i32 0, i32 0
+// CHECK-NEXT: store i32 2, ptr [[TMP102]], align 4
+// CHECK-NEXT: [[TMP103:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS22]], i32 0, i32 1
+// CHECK-NEXT: store i32 3, ptr [[TMP103]], align 4
+// CHECK-NEXT: [[TMP104:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS22]], i32 0, i32 2
+// CHECK-NEXT: store ptr [[TMP99]], ptr [[TMP104]], align 8
+// CHECK-NEXT: [[TMP105:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS22]], i32 0, i32 3
+// CHECK-NEXT: store ptr [[TMP100]], ptr [[TMP105]], align 8
+// CHECK-NEXT: [[TMP106:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS22]], i32 0, i32 4
+// CHECK-NEXT: store ptr [[TMP101]], ptr [[TMP106]], align 8
+// CHECK-NEXT: [[TMP107:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS22]], i32 0, i32 5
+// CHECK-NEXT: store ptr @.offload_maptypes.5, ptr [[TMP107]], align 8
+// CHECK-NEXT: [[TMP108:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS22]], i32 0, i32 6
+// CHECK-NEXT: store ptr null, ptr [[TMP108]], align 8
+// CHECK-NEXT: [[TMP109:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS22]], i32 0, i32 7
+// CHECK-NEXT: store ptr null, ptr [[TMP109]], align 8
+// CHECK-NEXT: [[TMP110:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS22]], i32 0, i32 8
+// CHECK-NEXT: store i64 0, ptr [[TMP110]], align 8
+// CHECK-NEXT: [[TMP111:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS22]], i32 0, i32 9
+// CHECK-NEXT: store i64 0, ptr [[TMP111]], align 8
+// CHECK-NEXT: [[TMP112:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS22]], i32 0, i32 10
+// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP112]], align 4
+// CHECK-NEXT: [[TMP113:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS22]], i32 0, i32 11
+// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP113]], align 4
+// CHECK-NEXT: [[TMP114:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS22]], i32 0, i32 12
+// CHECK-NEXT: store i32 0, ptr [[TMP114]], align 4
+// CHECK-NEXT: [[TMP115:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1C3barER10descriptorIfE_l59.region_id, ptr [[KERNEL_ARGS22]])
+// CHECK-NEXT: [[TMP116:%.*]] = icmp ne i32 [[TMP115]], 0
+// CHECK-NEXT: br i1 [[TMP116]], label [[OMP_OFFLOAD_FAILED23:%.*]], label [[OMP_OFFLOAD_CONT24:%.*]]
+// CHECK: omp_offload.failed23:
+// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1C3barER10descriptorIfE_l59(i64 [[TMP75]], ptr [[TMP76]]) #[[ATTR3]]
+// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT24]]
+// CHECK: omp_offload.cont24:
// CHECK-NEXT: ret void
//
//
@@ -433,13 +525,49 @@
// CHECK-NEXT: [[TMP6:%.*]] = load i32, ptr [[I]], align 4
// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP6]], 1
// CHECK-NEXT: store i32 [[INC]], ptr [[I]], align 4
-// CHECK-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP5:![0-9]+]]
+// CHECK-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP6:![0-9]+]]
+// CHECK: for.end:
+// CHECK-NEXT: ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN1C3barER10descriptorIfE_l59
+// CHECK-SAME: (i64 noundef [[CSIZE:%.*]], ptr noundef nonnull align 8 dereferenceable(40) [[D:%.*]]) #[[ATTR2]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[CSIZE_ADDR:%.*]] = alloca i64, align 8
+// CHECK-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[I:%.*]] = alloca i32, align 4
+// CHECK-NEXT: store i64 [[CSIZE]], ptr [[CSIZE_ADDR]], align 8
+// CHECK-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 8
+// CHECK-NEXT: store ptr [[TMP0]], ptr [[TMP]], align 8
+// CHECK-NEXT: store i32 0, ptr [[I]], align 4
+// CHECK-NEXT: br label [[FOR_COND:%.*]]
+// CHECK: for.cond:
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[I]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[CSIZE_ADDR]], align 4
+// CHECK-NEXT: [[CMP:%.*]] = icmp ult i32 [[TMP1]], [[TMP2]]
+// CHECK-NEXT: br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]]
+// CHECK: for.body:
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr, ptr [[TMP]], align 8
+// CHECK-NEXT: [[C:%.*]] = getelementptr inbounds [[STRUCT_DESCRIPTOR:%.*]], ptr [[TMP3]], i32 0, i32 1
+// CHECK-NEXT: [[TMP4:%.*]] = load ptr, ptr [[C]], align 8
+// CHECK-NEXT: [[TMP5:%.*]] = load i32, ptr [[I]], align 4
+// CHECK-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP5]] to i64
+// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds float, ptr [[TMP4]], i64 [[IDXPROM]]
+// CHECK-NEXT: store float 1.000000e+00, ptr [[ARRAYIDX]], align 4
+// CHECK-NEXT: br label [[FOR_INC:%.*]]
+// CHECK: for.inc:
+// CHECK-NEXT: [[TMP6:%.*]] = load i32, ptr [[I]], align 4
+// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP6]], 1
+// CHECK-NEXT: store i32 [[INC]], ptr [[I]], align 4
+// CHECK-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP10:![0-9]+]]
// CHECK: for.end:
// CHECK-NEXT: ret void
//
//
// CHECK-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg
-// CHECK-SAME: () #[[ATTR4:[0-9]+]] section ".text.startup" {
+// CHECK-SAME: () #[[ATTR5:[0-9]+]] section ".text.startup" {
// CHECK-NEXT: entry:
// CHECK-NEXT: call void @__tgt_register_requires(i64 1)
// CHECK-NEXT: ret void
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8378,7 +8378,8 @@
// individual members mapped. Emit an extra combined entry.
if (PartialStruct.Base.isValid()) {
CurInfo.NonContigInfo.Dims.push_back(0);
- emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct, VD);
+ emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct,
+ /*IsMapThis*/ !VD, VD);
}
// We need to append the results of this capture to what we already
@@ -8445,7 +8446,7 @@
void emitCombinedEntry(MapCombinedInfoTy &CombinedInfo,
MapFlagsArrayTy &CurTypes,
const StructRangeInfoTy &PartialStruct,
- const ValueDecl *VD = nullptr,
+ bool IsMapThis, const ValueDecl *VD = nullptr,
bool NotTargetParams = true) const {
if (CurTypes.size() == 1 &&
((CurTypes.back() & OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF) !=
@@ -8466,8 +8467,7 @@
const CXXMethodDecl *MD =
CGF.CurFuncDecl ? dyn_cast<CXXMethodDecl>(CGF.CurFuncDecl) : nullptr;
const CXXRecordDecl *RD = MD ? MD->getParent() : nullptr;
- // When VD is not null, it is not field of class, skip generating this[:1].
- bool HasBaseClass = RD && !VD ? RD->getNumBases() > 0 : false;
+ bool HasBaseClass = RD && IsMapThis ? RD->getNumBases() > 0 : false;
// There should not be a mapper for a combined entry.
if (HasBaseClass) {
// OpenMP 5.2 148:21:
@@ -10052,8 +10052,8 @@
if (PartialStruct.Base.isValid()) {
CombinedInfo.append(PartialStruct.PreliminaryMapData);
MEHandler.emitCombinedEntry(
- CombinedInfo, CurInfo.Types, PartialStruct, nullptr,
- !PartialStruct.PreliminaryMapData.BasePointers.empty());
+ CombinedInfo, CurInfo.Types, PartialStruct, CI->capturesThis(),
+ nullptr, !PartialStruct.PreliminaryMapData.BasePointers.empty());
}
// We need to append the results of this capture to what we already have.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
