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 @@ -8547,49 +8547,92 @@ } } - // 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, SmallVector> 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(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(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(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(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(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(VD); if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() || @@ -8604,51 +8647,45 @@ } } } - 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(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(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(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, 4> Processed; for (const auto *Cl : Clauses) { const auto *C = dyn_cast(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 @@ continue; VD = cast(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(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(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 @@ 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,18 @@ +// 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]); +}