[OMP] Fix compiler assert "DeclRefExpr for Decl not entered in LocalDeclMap?"

[jyu2]

Summary

Currently compiler assert when passing variable "memspace" in
omp_init_allocator.

omp_allocator_handle_t alloc=omp_init_allocator(memspace,1,traits)

The problem is memspace is not mapping to the target region. During
the call to emitAllocatorInit, calls to EmitVarDecl for "alloc", then
emit initialization of "alloc" that cause to assert.

If I understant correct, it is not necessary to emit variable
initialization, since "allocator" is private to target region.

To fix this call CGF.EmitAutoVarAlloca(allocator) instead
CGF.EmitVarDecl(allocator).

[ABataev]

LG

clang/lib/CodeGen/CGOpenMPRuntime.cpp

AllocatorTraitsLVal = CGF.MakeAddrLValue(Addr, CGF.getContext().VoidPtrTy,
                                            AllocatorTraitsLVal.getTBAAInfo());
   llvm::Value *Traits = Addr.getPointer();
 
   llvm::Value *AllocatorVal =
       CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                               CGM.getModule(), OMPRTL___kmpc_init_allocator),
                           {ThreadId, MemSpaceHandle, NumTraits, Traits});
   // Store to allocator.
-  CGF.EmitVarDecl(*cast<VarDecl>(
+  CGF.EmitAutoVarAlloca(*cast<VarDecl>(
       cast<DeclRefExpr>(Allocator->IgnoreParenImpCasts())->getDecl()));
   LValue AllocatorLVal = CGF.EmitLValue(Allocator->IgnoreParenImpCasts());
   AllocatorVal =
       CGF.EmitScalarConversion(AllocatorVal, CGF.getContext().VoidPtrTy,
                                Allocator->getType(), Allocator->getExprLoc());
   CGF.EmitStoreOfScalar(AllocatorVal, AllocatorLVal);
 }
 

clang/test/OpenMP/target_uses_allocators.c

void fie(void) {
   #pragma omp target uses_allocators(omp_pteam_mem_alloc) allocate(omp_pteam_mem_alloc: x) firstprivate(x)
   {}
   #pragma omp target uses_allocators(omp_thread_mem_alloc) allocate(omp_thread_mem_alloc: x) firstprivate(x) // expected-warning {{allocator with the 'thread' trait access has unspecified behavior on 'target' directive}}
   {}
 #pragma omp target uses_allocators(omp_null_allocator, omp_thread_mem_alloc, my_allocator(traits))
   {}
 }
 
+typedef enum omp_memspace_handle_t {
+  omp_default_mem_space = 0,
+  omp_large_cap_mem_space = 1,
+  omp_const_mem_space = 2,
+  omp_high_bw_mem_space = 3,
+  omp_low_lat_mem_space = 4,
+  llvm_omp_target_host_mem_space = 100,
+  llvm_omp_target_shared_mem_space = 101,
+  llvm_omp_target_device_mem_space = 102,
+  KMP_MEMSPACE_MAX_HANDLE = __UINTPTR_MAX__
+} omp_memspace_handle_t;
+
+extern omp_allocator_handle_t
+omp_init_allocator(omp_memspace_handle_t memspace, int ntraits,
+                   const omp_alloctrait_t traits[]);
+
+void *omp_aligned_alloc(unsigned long alignment, unsigned long size,
+                        omp_allocator_handle_t allocator);
+extern void * omp_alloc(int size, omp_allocator_handle_t a);
+#define N 1024
+
+void foo() {
+  int errors = 0;
+  omp_memspace_handle_t  memspace = omp_default_mem_space;
+  omp_alloctrait_t       traits[1] = {{omp_atk_alignment, 64}};
+  omp_allocator_handle_t alloc = omp_init_allocator(memspace,1,traits);
+  #pragma omp target map(tofrom: errors) uses_allocators(alloc(traits))
+  { }
+}
 #endif
 
 // CHECK: %[[#R0:]] = call i32 @__kmpc_global_thread_num(ptr @1)
 // CHECK-NEXT: store i64 %x, ptr %x.addr, align 8
 // CHECK-NEXT: %.x..void.addr = call ptr @__kmpc_alloc(i32 %[[#R0]], i64 4, ptr null)
 // CHECK-NEXT: %[[#R1:]] = load i32, ptr %x.addr, align 4
 // CHECK-NEXT: store i32 %[[#R1]], ptr %.x..void.addr, align 4
 // CHECK-NEXT: call void @__kmpc_free(i32 %[[#R0]], ptr %.x..void.addr, ptr null)

 // CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
 // Destroy allocator upon exit from the region.
 // CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
 // CHECK: [[CONV:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
 // CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV]])
+
+// CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
+// CHECK: [[MY_ALLOCATOR_ADDR:%alloc]] = alloca i64,
+// CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 1, ptr [[TRAITS_ADDR]])
+// CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
+// CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
+
+// Destroy allocator upon exit from the region.
+// CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
+// CHECK: [[CONV1:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
+// CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV1]])