[OpenMP] Map omp_default_mem_alloc to global memory

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.

Author: jhuber6

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

@@ -2048,15 +2048,15 @@ 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.
-      return Address::invalid();
+      AS = LangAS::opencl_global;
+      break;
+    case OMPAllocateDeclAttr::OMPThreadMemAlloc:
+      AS = LangAS::opencl_private;
+      break;
     case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:
       // TODO: implement aupport for user-defined allocators.
       return Address::invalid();
@@ -2208,12 +2208,14 @@ 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:
-    AS = LangAS::Default;
+    AS = LangAS::opencl_global;
+    return true;
+  case OMPAllocateDeclAttr::OMPThreadMemAlloc:
+    AS = LangAS::opencl_private;
     return true;
   case OMPAllocateDeclAttr::OMPConstMemAlloc:
     AS = LangAS::cuda_constant;

clang/test/OpenMP/nvptx_allocate_codegen.cpp

@@ -87,18 +87,17 @@ 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 addrspacecast (ptr addrspace(1) @b1 to ptr), align 8
 // CHECK1-NEXT:    [[CALL:%.*]] = call noundef i32 @_Z3fooIiET_v() #[[ATTR7:[0-9]+]]
 // CHECK1-NEXT:    ret i32 [[CALL]]
 //
 //
 // CHECK1-LABEL: define {{[^@]+}}@_Z3fooIiET_v
 // CHECK1-SAME: () #[[ATTR1:[0-9]+]] comdat {
 // CHECK1-NEXT:  entry:
-// CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr @_ZN2STIiE1mE, align 4
+// CHECK1-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @_ZN2STIiE1mE to ptr), align 4
 // CHECK1-NEXT:    store i32 [[TMP0]], ptr @v, align 4
 // CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr @v, align 4
 // CHECK1-NEXT:    ret i32 [[TMP1]]
@@ -120,13 +119,12 @@ void bar() {
 // CHECK1-NEXT:  entry:
 // CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
 // CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
-// CHECK1-NEXT:    [[BAR_A:%.*]] = alloca float, align 4
 // CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
 // CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
-// CHECK1-NEXT:    [[TMP0:%.*]] = load float, ptr [[BAR_A]], align 4
+// CHECK1-NEXT:    [[TMP0:%.*]] = load float, ptr @bar_a, align 4
 // CHECK1-NEXT:    [[CONV:%.*]] = fpext float [[TMP0]] to double
 // CHECK1-NEXT:    store double [[CONV]], ptr addrspacecast (ptr addrspace(3) @bar_b to ptr), align 8
-// CHECK1-NEXT:    call void @_Z3bazRf(ptr noundef nonnull align 4 dereferenceable(4) [[BAR_A]]) #[[ATTR7]]
+// CHECK1-NEXT:    call void @_Z3bazRf(ptr noundef nonnull align 4 dereferenceable(4) @bar_a) #[[ATTR7]]
 // CHECK1-NEXT:    ret void
 //
 //

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