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]]) _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits