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 @@ -8691,7 +8691,7 @@ DeferredInfo[nullptr].emplace_back(IE, VD, /*ForDeviceAddr=*/true); } else { llvm::Value *Ptr; - if (IE->isGLValue()) + if (IE->isGLValue() && !IE->getType()->isPointerType()) Ptr = CGF.EmitLValue(IE).getPointer(CGF); else Ptr = CGF.EmitScalarExpr(IE); diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -7241,6 +7241,7 @@ // declaration used by the mapping logic. In some cases we may get // OMPCapturedExprDecl that refers to the original declaration. const ValueDecl *MatchingVD = OrigVD; + bool isPartOfAStruct = false; if (const auto *OED = dyn_cast(MatchingVD)) { // OMPCapturedExprDecl are used to privative fields of the current // structure. @@ -7248,6 +7249,7 @@ assert(isa(ME->getBase()) && "Base should be the current struct!"); MatchingVD = ME->getMemberDecl(); + isPartOfAStruct = true; } // If we don't have information about the current list item, move on to @@ -7259,8 +7261,11 @@ Address PrivAddr = InitAddrIt->getSecond(); // For declrefs and variable length array need to load the pointer for // correct mapping, since the pointer to the data was passed to the runtime. - if (isa(Ref->IgnoreParenImpCasts()) || - MatchingVD->getType()->isArrayType()) { + // Pointer types are already mapped correctly so no need to do a load unless + // the pointer type is part of a struct. + if ((isa(Ref->IgnoreParenImpCasts()) || + MatchingVD->getType()->isArrayType()) && + (isPartOfAStruct || !MatchingVD->getType()->isPointerType())) { QualType PtrTy = getContext().getPointerType( OrigVD->getType().getNonReferenceType()); PrivAddr = EmitLoadOfPointer( diff --git a/clang/test/OpenMP/target_data_use_device_addr_codegen_ptr.cpp b/clang/test/OpenMP/target_data_use_device_addr_codegen_ptr.cpp new file mode 100755 --- /dev/null +++ b/clang/test/OpenMP/target_data_use_device_addr_codegen_ptr.cpp @@ -0,0 +1,41 @@ +// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s +// expected-no-diagnostics + +#ifndef HEADER +#define HEADER + +int main() { + float x_array[256]; + float *x = &x_array[0]; + + // make x available on the GPU + #pragma omp target data map(tofrom:x[0:256]) + { + #pragma omp target data use_device_addr(x) + { + x[0] = 2; + } + } + return x[0] == 2; +} + +// CHECK-LABEL: @main() +// CHECK: [[X:%.+]] = alloca ptr, align 8 +// CHECK: call void @__tgt_target_data_begin_mapper( +// CHECK: [[LOADED_X:%.+]] = load ptr, ptr [[X]], align 8 +// CHECK: [[BASE_PTR_GEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_BASE_PTR:%.+]], i32 0, i32 0 +// CHECK: store ptr [[LOADED_X]], ptr [[BASE_PTR_GEP]], align 8 +// CHECK: [[OFFLOAD_PTR_GEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_PTR:%.+]], i32 0, i32 0 +// CHECK: store ptr [[LOADED_X]], ptr [[OFFLOAD_PTR_GEP]], align 8 +// CHECK: getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_BASE_PTR]], i32 0, i32 0 +// CHECK: getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_PTR]], i32 0, i32 0 +// CHECK: call void @__tgt_target_data_begin_mapper( +// CHECK: [[LOADED_DEVICE_X:%.+]] = load ptr, ptr [[BASE_PTR_GEP]], align 8 +// CHECK: %arrayidx5 = getelementptr inbounds float, ptr %13, i64 0 +// CHECK: store float 2.000000e+00, ptr %arrayidx5, align 4 +// CHECK: getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_BASE_PTR]], i32 0, i32 0 +// CHECK: getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_PTR]], i32 0, i32 0 +// CHECK: call void @__tgt_target_data_end_mapper( +// CHECK: call void @__tgt_target_data_end_mapper( + +#endif