diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -6041,7 +6041,7 @@ CGM.getModule(), OMPRTL___kmpc_init_allocator), {ThreadId, MemSpaceHandle, NumTraits, Traits}); // Store to allocator. - CGF.EmitVarDecl(*cast( + CGF.EmitAutoVarAlloca(*cast( cast(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 --- a/clang/test/OpenMP/target_uses_allocators.c +++ b/clang/test/OpenMP/target_uses_allocators.c @@ -64,6 +64,35 @@ {} } +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 @@ // 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]])