[Openmp-commits] [openmp] 30cc712 - [Clang][OpenMP] Fix run time crash when use_device_addr is used.
Jennifer Yu via Openmp-commits
openmp-commits at lists.llvm.org
Tue Sep 27 11:57:52 PDT 2022
Author: Jennifer Yu
Date: 2022-09-27T11:53:57-07:00
New Revision: 30cc712eb6f23a5c7beaae669bf2ab6beede7f20
URL: https://github.com/llvm/llvm-project/commit/30cc712eb6f23a5c7beaae669bf2ab6beede7f20
DIFF: https://github.com/llvm/llvm-project/commit/30cc712eb6f23a5c7beaae669bf2ab6beede7f20.diff
LOG: [Clang][OpenMP] Fix run time crash when use_device_addr is used.
It is data mapping ordering problem.
According omp spec
If one or more map clauses are present, the list item conversions that
are performed for any use_device_ptr or use_device_addr clause occur
after all variables are mapped on entry to the region according to those
map clauses.
The change is to put mapping data for use_device_addr at end of data
mapping array.
Differential Revision: https://reviews.llvm.org/D134556
Added:
openmp/libomptarget/test/mapping/target_use_device_addr.c
Modified:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_data_codegen.cpp
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 6ff36c72e0319..5bf0e0815111e 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8547,49 +8547,92 @@ class MappableExprsHandler {
}
}
- // Look at the use_device_ptr clause information and mark the existing map
- // entries as such. If there is no map information for an entry in the
- // use_device_ptr list, we create one with map type 'alloc' and zero size
- // section. It is the user fault if that was not mapped before. If there is
- // no map information and the pointer is a struct member, then we defer the
- // emission of that entry until the whole struct has been processed.
+ // Look at the use_device_ptr and use_device_addr clauses information and
+ // mark the existing map entries as such. If there is no map information for
+ // an entry in the use_device_ptr and use_device_addr list, we create one
+ // with map type 'alloc' and zero size section. It is the user fault if that
+ // was not mapped before. If there is no map information and the pointer is
+ // a struct member, then we defer the emission of that entry until the whole
+ // struct has been processed.
llvm::MapVector<CanonicalDeclPtr<const Decl>,
SmallVector<DeferredDevicePtrEntryTy, 4>>
DeferredInfo;
- MapCombinedInfoTy UseDevicePtrCombinedInfo;
+ MapCombinedInfoTy UseDeviceDataCombinedInfo;
+
+ auto &&UseDeviceDataCombinedInfoGen =
+ [&UseDeviceDataCombinedInfo](const ValueDecl *VD, llvm::Value *Ptr,
+ CodeGenFunction &CGF) {
+ UseDeviceDataCombinedInfo.Exprs.push_back(VD);
+ UseDeviceDataCombinedInfo.BasePointers.emplace_back(Ptr, VD);
+ UseDeviceDataCombinedInfo.Pointers.push_back(Ptr);
+ UseDeviceDataCombinedInfo.Sizes.push_back(
+ llvm::Constant::getNullValue(CGF.Int64Ty));
+ UseDeviceDataCombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM);
+ UseDeviceDataCombinedInfo.Mappers.push_back(nullptr);
+ };
- for (const auto *Cl : Clauses) {
- const auto *C = dyn_cast<OMPUseDevicePtrClause>(Cl);
- if (!C)
- continue;
- for (const auto L : C->component_lists()) {
- OMPClauseMappableExprCommon::MappableExprComponentListRef Components =
- std::get<1>(L);
- assert(!Components.empty() &&
- "Not expecting empty list of components!");
- const ValueDecl *VD = Components.back().getAssociatedDeclaration();
- VD = cast<ValueDecl>(VD->getCanonicalDecl());
- const Expr *IE = Components.back().getAssociatedExpression();
- // If the first component is a member expression, we have to look into
- // 'this', which maps to null in the map of map information. Otherwise
- // look directly for the information.
- auto It = Info.find(isa<MemberExpr>(IE) ? nullptr : VD);
-
- // We potentially have map information for this declaration already.
- // Look for the first set of components that refer to it.
- if (It != Info.end()) {
- bool Found = false;
- for (auto &Data : It->second) {
- auto *CI = llvm::find_if(Data, [VD](const MapInfo &MI) {
- return MI.Components.back().getAssociatedDeclaration() == VD;
- });
- // If we found a map entry, signal that the pointer has to be
- // returned and move on to the next declaration. Exclude cases where
- // the base pointer is mapped as array subscript, array section or
- // array shaping. The base address is passed as a pointer to base in
- // this case and cannot be used as a base for use_device_ptr list
- // item.
- if (CI != Data.end()) {
+ auto &&MapInfoGen =
+ [&DeferredInfo, &UseDeviceDataCombinedInfoGen,
+ &InfoGen](CodeGenFunction &CGF, const Expr *IE, const ValueDecl *VD,
+ OMPClauseMappableExprCommon::MappableExprComponentListRef
+ Components,
+ bool IsImplicit, bool IsDevAddr) {
+ // We didn't find any match in our map information - generate a zero
+ // size array section - if the pointer is a struct member we defer
+ // this action until the whole struct has been processed.
+ if (isa<MemberExpr>(IE)) {
+ // Insert the pointer into Info to be processed by
+ // generateInfoForComponentList. Because it is a member pointer
+ // without a pointee, no entry will be generated for it, therefore
+ // we need to generate one after the whole struct has been
+ // processed. Nonetheless, generateInfoForComponentList must be
+ // called to take the pointer into account for the calculation of
+ // the range of the partial struct.
+ InfoGen(nullptr, Other, Components, OMPC_MAP_unknown, llvm::None,
+ llvm::None, /*ReturnDevicePointer=*/false, IsImplicit,
+ nullptr, nullptr, IsDevAddr);
+ DeferredInfo[nullptr].emplace_back(IE, VD, IsDevAddr);
+ } else {
+ llvm::Value *Ptr;
+ if (IsDevAddr) {
+ if (IE->isGLValue())
+ Ptr = CGF.EmitLValue(IE).getPointer(CGF);
+ else
+ Ptr = CGF.EmitScalarExpr(IE);
+ } else {
+ Ptr = CGF.EmitLoadOfScalar(CGF.EmitLValue(IE), IE->getExprLoc());
+ }
+ UseDeviceDataCombinedInfoGen(VD, Ptr, CGF);
+ }
+ };
+
+ auto &&IsMapInfoExist = [&Info](CodeGenFunction &CGF, const ValueDecl *VD,
+ const Expr *IE, bool IsDevAddr) -> bool {
+ // We potentially have map information for this declaration already.
+ // Look for the first set of components that refer to it. If found,
+ // return true.
+ // If the first component is a member expression, we have to look into
+ // 'this', which maps to null in the map of map information. Otherwise
+ // look directly for the information.
+ auto It = Info.find(isa<MemberExpr>(IE) ? nullptr : VD);
+ if (It != Info.end()) {
+ bool Found = false;
+ for (auto &Data : It->second) {
+ auto *CI = llvm::find_if(Data, [VD](const MapInfo &MI) {
+ return MI.Components.back().getAssociatedDeclaration() == VD;
+ });
+ // If we found a map entry, signal that the pointer has to be
+ // returned and move on to the next declaration. Exclude cases where
+ // the base pointer is mapped as array subscript, array section or
+ // array shaping. The base address is passed as a pointer to base in
+ // this case and cannot be used as a base for use_device_ptr list
+ // item.
+ if (CI != Data.end()) {
+ if (IsDevAddr) {
+ CI->ReturnDevicePointer = true;
+ Found = true;
+ break;
+ } else {
auto PrevCI = std::next(CI->Components.rbegin());
const auto *VarD = dyn_cast<VarDecl>(VD);
if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
@@ -8604,51 +8647,45 @@ class MappableExprsHandler {
}
}
}
- if (Found)
- continue;
- }
-
- // We didn't find any match in our map information - generate a zero
- // size array section - if the pointer is a struct member we defer this
- // action until the whole struct has been processed.
- if (isa<MemberExpr>(IE)) {
- // Insert the pointer into Info to be processed by
- // generateInfoForComponentList. Because it is a member pointer
- // without a pointee, no entry will be generated for it, therefore
- // we need to generate one after the whole struct has been processed.
- // Nonetheless, generateInfoForComponentList must be called to take
- // the pointer into account for the calculation of the range of the
- // partial struct.
- InfoGen(nullptr, Other, Components, OMPC_MAP_unknown, llvm::None,
- llvm::None, /*ReturnDevicePointer=*/false, C->isImplicit(),
- nullptr);
- DeferredInfo[nullptr].emplace_back(IE, VD, /*ForDeviceAddr=*/false);
- } else {
- llvm::Value *Ptr =
- CGF.EmitLoadOfScalar(CGF.EmitLValue(IE), IE->getExprLoc());
- UseDevicePtrCombinedInfo.Exprs.push_back(VD);
- UseDevicePtrCombinedInfo.BasePointers.emplace_back(Ptr, VD);
- UseDevicePtrCombinedInfo.Pointers.push_back(Ptr);
- UseDevicePtrCombinedInfo.Sizes.push_back(
- llvm::Constant::getNullValue(CGF.Int64Ty));
- UseDevicePtrCombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM);
- UseDevicePtrCombinedInfo.Mappers.push_back(nullptr);
}
+ return Found;
}
- }
+ return false;
+ };
- // Look at the use_device_addr clause information and mark the existing map
+ // Look at the use_device_ptr clause information and mark the existing map
// entries as such. If there is no map information for an entry in the
- // use_device_addr list, we create one with map type 'alloc' and zero size
+ // use_device_ptr list, we create one with map type 'alloc' and zero size
// section. It is the user fault if that was not mapped before. If there is
// no map information and the pointer is a struct member, then we defer the
// emission of that entry until the whole struct has been processed.
+ for (const auto *Cl : Clauses) {
+ const auto *C = dyn_cast<OMPUseDevicePtrClause>(Cl);
+ if (!C)
+ continue;
+ for (const auto L : C->component_lists()) {
+ OMPClauseMappableExprCommon::MappableExprComponentListRef Components =
+ std::get<1>(L);
+ assert(!Components.empty() &&
+ "Not expecting empty list of components!");
+ const ValueDecl *VD = Components.back().getAssociatedDeclaration();
+ VD = cast<ValueDecl>(VD->getCanonicalDecl());
+ const Expr *IE = Components.back().getAssociatedExpression();
+ if (IsMapInfoExist(CGF, VD, IE, /*IsDevAddr=*/false))
+ continue;
+ MapInfoGen(CGF, IE, VD, Components, C->isImplicit(),
+ /*IsDevAddr=*/false);
+ }
+ }
+
llvm::SmallDenseSet<CanonicalDeclPtr<const Decl>, 4> Processed;
for (const auto *Cl : Clauses) {
const auto *C = dyn_cast<OMPUseDeviceAddrClause>(Cl);
if (!C)
continue;
for (const auto L : C->component_lists()) {
+ OMPClauseMappableExprCommon::MappableExprComponentListRef Components =
+ std::get<1>(L);
assert(!std::get<1>(L).empty() &&
"Not expecting empty list of components!");
const ValueDecl *VD = std::get<1>(L).back().getAssociatedDeclaration();
@@ -8656,60 +8693,10 @@ class MappableExprsHandler {
continue;
VD = cast<ValueDecl>(VD->getCanonicalDecl());
const Expr *IE = std::get<1>(L).back().getAssociatedExpression();
- // If the first component is a member expression, we have to look into
- // 'this', which maps to null in the map of map information. Otherwise
- // look directly for the information.
- auto It = Info.find(isa<MemberExpr>(IE) ? nullptr : VD);
-
- // We potentially have map information for this declaration already.
- // Look for the first set of components that refer to it.
- if (It != Info.end()) {
- bool Found = false;
- for (auto &Data : It->second) {
- auto *CI = llvm::find_if(Data, [VD](const MapInfo &MI) {
- return MI.Components.back().getAssociatedDeclaration() == VD;
- });
- // If we found a map entry, signal that the pointer has to be
- // returned and move on to the next declaration.
- if (CI != Data.end()) {
- CI->ReturnDevicePointer = true;
- Found = true;
- break;
- }
- }
- if (Found)
- continue;
- }
-
- // We didn't find any match in our map information - generate a zero
- // size array section - if the pointer is a struct member we defer this
- // action until the whole struct has been processed.
- if (isa<MemberExpr>(IE)) {
- // Insert the pointer into Info to be processed by
- // generateInfoForComponentList. Because it is a member pointer
- // without a pointee, no entry will be generated for it, therefore
- // we need to generate one after the whole struct has been processed.
- // Nonetheless, generateInfoForComponentList must be called to take
- // the pointer into account for the calculation of the range of the
- // partial struct.
- InfoGen(nullptr, Other, std::get<1>(L), OMPC_MAP_unknown, llvm::None,
- llvm::None, /*ReturnDevicePointer=*/false, C->isImplicit(),
- nullptr, nullptr, /*ForDeviceAddr=*/true);
- DeferredInfo[nullptr].emplace_back(IE, VD, /*ForDeviceAddr=*/true);
- } else {
- llvm::Value *Ptr;
- if (IE->isGLValue())
- Ptr = CGF.EmitLValue(IE).getPointer(CGF);
- else
- Ptr = CGF.EmitScalarExpr(IE);
- CombinedInfo.Exprs.push_back(VD);
- CombinedInfo.BasePointers.emplace_back(Ptr, VD);
- CombinedInfo.Pointers.push_back(Ptr);
- CombinedInfo.Sizes.push_back(
- llvm::Constant::getNullValue(CGF.Int64Ty));
- CombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM);
- CombinedInfo.Mappers.push_back(nullptr);
- }
+ if (IsMapInfoExist(CGF, VD, IE, /*IsDevAddr=*/true))
+ continue;
+ MapInfoGen(CGF, IE, VD, Components, C->isImplicit(),
+ /*IsDevAddr=*/true);
}
}
@@ -8798,7 +8785,7 @@ class MappableExprsHandler {
CombinedInfo.append(CurInfo);
}
// Append data for use_device_ptr clauses.
- CombinedInfo.append(UseDevicePtrCombinedInfo);
+ CombinedInfo.append(UseDeviceDataCombinedInfo);
}
public:
diff --git a/clang/test/OpenMP/target_data_codegen.cpp b/clang/test/OpenMP/target_data_codegen.cpp
index 73a3bac1fc60e..3ea267258fed4 100644
--- a/clang/test/OpenMP/target_data_codegen.cpp
+++ b/clang/test/OpenMP/target_data_codegen.cpp
@@ -596,15 +596,18 @@ void test_close_modifier(int arg) {
}
#endif
///==========================================================================///
-// RUN: %clang_cc1 -no-opaque-pointers -DCK7 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK7 --check-prefix CK7-64
-// RUN: %clang_cc1 -no-opaque-pointers -DCK7 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
-// RUN: %clang_cc1 -no-opaque-pointers -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK7 --check-prefix CK7-64
+// RUN: %clang_cc1 -no-opaque-pointers -DCK7 -verify -fopenmp -fopenmp-targets=x86_64 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK7 --check-prefix CK7-64
+// RUN: %clang_cc1 -no-opaque-pointers -DCK7 -fopenmp -fopenmp-targets=x86_64 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -no-opaque-pointers -fopenmp -fopenmp-targets=x86_64 -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK7 --check-prefix CK7-64
-// RUN: %clang_cc1 -no-opaque-pointers -DCK7 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY7 %s
-// RUN: %clang_cc1 -no-opaque-pointers -DCK7 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
-// RUN: %clang_cc1 -no-opaque-pointers -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY7 %s
+// RUN: %clang_cc1 -no-opaque-pointers -DCK7 -verify -fopenmp-simd -fopenmp-targets=x86_64 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY7 %s
+// RUN: %clang_cc1 -no-opaque-pointers -DCK7 -fopenmp-simd -fopenmp-targets=x86_64 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -no-opaque-pointers -fopenmp-simd -fopenmp-targets=x86_64 -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY7 %s
// SIMD-ONLY7-NOT: {{__kmpc|__tgt}}
#ifdef CK7
+// CK7: private unnamed_addr constant [2 x i64] [i64 64, i64 64]
+// CK7: private unnamed_addr constant [2 x i64] [i64 3, i64 64]
+// CK7-NOT: private unnamed_addr constant [2 x i64] [i64 64, i64 3]
// CK7: test_device_ptr_addr
void test_device_ptr_addr(int arg) {
int *p;
@@ -612,6 +615,16 @@ void test_device_ptr_addr(int arg) {
// CK7: add nsw i32
#pragma omp target data use_device_ptr(p) use_device_addr(arg)
{ ++arg, ++(*p); }
+
+ short x[10];
+ short *xp = &x[0];
+
+ x[1] = 111;
+
+ #pragma omp target data map(tofrom: x) use_device_addr(xp[1:3])
+ {
+ xp[1] = 222;
+ }
}
#endif
///==========================================================================///
diff --git a/openmp/libomptarget/test/mapping/target_use_device_addr.c b/openmp/libomptarget/test/mapping/target_use_device_addr.c
new file mode 100644
index 0000000000000..05a5aea88a892
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/target_use_device_addr.c
@@ -0,0 +1,18 @@
+// RUN: %libomptarget-compile-generic -fopenmp-version=51
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+#include <stdio.h>
+int main() {
+ short x[10];
+ short *xp = &x[0];
+
+ x[1] = 111;
+
+ printf("%d, %p\n", xp[1], &xp[1]);
+#pragma omp target data use_device_addr(xp [1:3]) map(tofrom : x)
+#pragma omp target is_device_ptr(xp)
+ { xp[1] = 222; }
+ // CHECK: 222
+ printf("%d, %p\n", xp[1], &xp[1]);
+}
More information about the Openmp-commits
mailing list