diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -8556,7 +8556,7 @@ llvm::MapVector, SmallVector> DeferredInfo; - MapCombinedInfoTy UseDevicePtrCombinedInfo; + MapCombinedInfoTy UseDeviceDataCombinedInfo; for (const auto *Cl : Clauses) { const auto *C = dyn_cast(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: diff --git a/clang/test/OpenMP/target_data_codegen.cpp b/clang/test/OpenMP/target_data_codegen.cpp --- a/clang/test/OpenMP/target_data_codegen.cpp +++ b/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 ///==========================================================================/// 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 --- /dev/null +++ b/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 +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]); +}