[clang] e0326b6 - [OpenMP] Map `omp_default_mem_alloc` to global memory (#104790)

[via cfe-commits]

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);
   }