jyu2 created this revision.
jyu2 added reviewers: ABataev, jdoerfert, mikerice.
jyu2 added a project: OpenMP.
Herald added subscribers: guansong, yaxunl.
Herald added a project: All.
jyu2 requested review of this revision.
Herald added subscribers: openmp-commits, cfe-commits, sstefan1.
Herald added a project: clang.
It is due to data mapping ordering.
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.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D134556
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_data_codegen.cpp
openmp/libomptarget/test/mapping/target_use_device_addr.c
Index: openmp/libomptarget/test/mapping/target_use_device_addr.c
===================================================================
--- /dev/null
+++ openmp/libomptarget/test/mapping/target_use_device_addr.c
@@ -0,0 +1,21 @@
+// 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]);
+}
Index: clang/test/OpenMP/target_data_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_data_codegen.cpp
+++ clang/test/OpenMP/target_data_codegen.cpp
@@ -596,15 +596,18 @@
}
#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 @@
// 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
///==========================================================================///
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8556,7 +8556,7 @@
llvm::MapVector<CanonicalDeclPtr<const Decl>,
SmallVector<DeferredDevicePtrEntryTy, 4>>
DeferredInfo;
- MapCombinedInfoTy UseDevicePtrCombinedInfo;
+ MapCombinedInfoTy UseDeviceDataCombinedInfo;
for (const auto *Cl : Clauses) {
const auto *C = dyn_cast<OMPUseDevicePtrClause>(Cl);
@@ -8626,13 +8626,13 @@
} 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(
+ 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));
- UseDevicePtrCombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM);
- UseDevicePtrCombinedInfo.Mappers.push_back(nullptr);
+ UseDeviceDataCombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM);
+ UseDeviceDataCombinedInfo.Mappers.push_back(nullptr);
}
}
}
@@ -8702,13 +8702,13 @@
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(
+ 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));
- CombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM);
- CombinedInfo.Mappers.push_back(nullptr);
+ UseDeviceDataCombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM);
+ UseDeviceDataCombinedInfo.Mappers.push_back(nullptr);
}
}
}
@@ -8798,7 +8798,7 @@
CombinedInfo.append(CurInfo);
}
// Append data for use_device_ptr clauses.
- CombinedInfo.append(UseDevicePtrCombinedInfo);
+ CombinedInfo.append(UseDeviceDataCombinedInfo);
}
public:
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
