[clang] aaa33b6 - Fix assert "DeclRefExpr for Decl not entered in LocalDeclMap?"

[Jennifer Yu via cfe-commits]

Author: Jennifer Yu
Date: 2023-05-30T16:40:06-07:00
New Revision: aaa33b6a98de2be7cdc827b13e60c103206d6461

URL: https://github.com/llvm/llvm-project/commit/aaa33b6a98de2be7cdc827b13e60c103206d6461
DIFF: https://github.com/llvm/llvm-project/commit/aaa33b6a98de2be7cdc827b13e60c103206d6461.diff

LOG: Fix assert "DeclRefExpr for Decl not entered in LocalDeclMap?"

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).

Differential Revision: https://reviews.llvm.org/D151743

Added: 
    

Modified: 
    clang/lib/CodeGen/CGOpenMPRuntime.cpp
    clang/test/OpenMP/target_uses_allocators.c

Removed: 
    


################################################################################
diff  --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 1f1db83378233..5957e59097709 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -6041,7 +6041,7 @@ void CGOpenMPRuntime::emitUsesAllocatorsInit(CodeGenFunction &CGF,
                               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 =

diff  --git a/clang/test/OpenMP/target_uses_allocators.c b/clang/test/OpenMP/target_uses_allocators.c
index eab202671e793..0352a5874bf12 100644
--- a/clang/test/OpenMP/target_uses_allocators.c
+++ b/clang/test/OpenMP/target_uses_allocators.c
@@ -64,6 +64,35 @@ void fie(void) {
   {}
 }
 
+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)
@@ -140,3 +169,15 @@ void fie(void) {
 // 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]])