Author: Joseph Huber Date: 2024-08-20T12:00:41-05:00 New Revision: e0326b668efda1d4ed7969c1a6cca44a487d24b6
URL: https://github.com/llvm/llvm-project/commit/e0326b668efda1d4ed7969c1a6cca44a487d24b6 DIFF: https://github.com/llvm/llvm-project/commit/e0326b668efda1d4ed7969c1a6cca44a487d24b6.diff LOG: [OpenMP] Map `omp_default_mem_alloc` to global memory (#104790) Summary: Currently, we assign this to private memory. This causes failures on some SOLLVE tests. The standard isn't clear on the semantics of this allocation type, but there seems to be a consensus that it's supposed to be shared memory. Added: Modified: clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp clang/test/OpenMP/nvptx_allocate_codegen.cpp offload/test/api/omp_device_alloc.c Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp index 8965a14d88a6fb..9e095a37552196 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp @@ -2048,14 +2048,12 @@ Address CGOpenMPRuntimeGPU::getAddressOfLocalVariable(CodeGenFunction &CGF, const auto *A = VD->getAttr<OMPAllocateDeclAttr>(); auto AS = LangAS::Default; switch (A->getAllocatorType()) { - // Use the default allocator here as by default local vars are - // threadlocal. case OMPAllocateDeclAttr::OMPNullMemAlloc: case OMPAllocateDeclAttr::OMPDefaultMemAlloc: - case OMPAllocateDeclAttr::OMPThreadMemAlloc: case OMPAllocateDeclAttr::OMPHighBWMemAlloc: case OMPAllocateDeclAttr::OMPLowLatMemAlloc: - // Follow the user decision - use default allocation. + break; + case OMPAllocateDeclAttr::OMPThreadMemAlloc: return Address::invalid(); case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc: // TODO: implement aupport for user-defined allocators. @@ -2208,11 +2206,11 @@ bool CGOpenMPRuntimeGPU::hasAllocateAttributeForGlobalVar(const VarDecl *VD, case OMPAllocateDeclAttr::OMPNullMemAlloc: case OMPAllocateDeclAttr::OMPDefaultMemAlloc: // Not supported, fallback to the default mem space. - case OMPAllocateDeclAttr::OMPThreadMemAlloc: case OMPAllocateDeclAttr::OMPLargeCapMemAlloc: case OMPAllocateDeclAttr::OMPCGroupMemAlloc: case OMPAllocateDeclAttr::OMPHighBWMemAlloc: case OMPAllocateDeclAttr::OMPLowLatMemAlloc: + case OMPAllocateDeclAttr::OMPThreadMemAlloc: AS = LangAS::Default; return true; case OMPAllocateDeclAttr::OMPConstMemAlloc: diff --git a/clang/test/OpenMP/nvptx_allocate_codegen.cpp b/clang/test/OpenMP/nvptx_allocate_codegen.cpp index 3f3457dab33c2d..4f38e2c50efe3c 100644 --- a/clang/test/OpenMP/nvptx_allocate_codegen.cpp +++ b/clang/test/OpenMP/nvptx_allocate_codegen.cpp @@ -87,10 +87,9 @@ void bar() { // CHECK1-SAME: () #[[ATTR0:[0-9]+]] { // CHECK1-NEXT: entry: // CHECK1-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 -// CHECK1-NEXT: [[B:%.*]] = alloca double, align 8 // CHECK1-NEXT: store i32 0, ptr [[RETVAL]], align 4 // CHECK1-NEXT: store i32 2, ptr @_ZZ4mainE1a, align 4 -// CHECK1-NEXT: store double 3.000000e+00, ptr [[B]], align 8 +// CHECK1-NEXT: store double 3.000000e+00, ptr @b1, align 8 // CHECK1-NEXT: [[CALL:%.*]] = call noundef i32 @_Z3fooIiET_v() #[[ATTR7:[0-9]+]] // CHECK1-NEXT: ret i32 [[CALL]] // diff --git a/offload/test/api/omp_device_alloc.c b/offload/test/api/omp_device_alloc.c index 368c6cfe42949b..b4cfe442d9ded5 100644 --- a/offload/test/api/omp_device_alloc.c +++ b/offload/test/api/omp_device_alloc.c @@ -5,13 +5,19 @@ #include <stdio.h> int main() { -#pragma omp target teams num_teams(4) -#pragma omp parallel +#pragma omp target { - int *ptr = (int *)omp_alloc(sizeof(int), omp_default_mem_alloc); + int *ptr; +#pragma omp allocate(ptr) allocator(omp_default_mem_alloc) + ptr = omp_alloc(sizeof(int), omp_default_mem_alloc); assert(ptr && "Ptr is (null)!"); - *ptr = 1; - assert(*ptr == 1 && "Ptr is not 1"); + *ptr = 0; +#pragma omp parallel num_threads(32) + { +#pragma omp atomic + *ptr += 1; + } + assert(*ptr == 32 && "Ptr is not 32"); omp_free(ptr, omp_default_mem_alloc); } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
