jyu2 created this revision.
jyu2 added reviewers: ABataev, jdoerfert.
jyu2 added projects: OpenMP, clang.
Herald added a project: All.
jyu2 requested review of this revision.
Herald added subscribers: cfe-commits, jplehr, sstefan1.
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).
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D151743
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_uses_allocators.c
Index: clang/test/OpenMP/target_uses_allocators.c
===================================================================
--- clang/test/OpenMP/target_uses_allocators.c
+++ 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]])
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ 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<VarDecl>(
+ CGF.EmitAutoVarAlloca(*cast<VarDecl>(
cast<DeclRefExpr>(Allocator->IgnoreParenImpCasts())->getDecl()));
LValue AllocatorLVal = CGF.EmitLValue(Allocator->IgnoreParenImpCasts());
AllocatorVal =
Index: clang/test/OpenMP/target_uses_allocators.c
===================================================================
--- clang/test/OpenMP/target_uses_allocators.c
+++ 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]])
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ 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<VarDecl>(
+ CGF.EmitAutoVarAlloca(*cast<VarDecl>(
cast<DeclRefExpr>(Allocator->IgnoreParenImpCasts())->getDecl()));
LValue AllocatorLVal = CGF.EmitLValue(Allocator->IgnoreParenImpCasts());
AllocatorVal =
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
