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: openmp-commits, cfe-commits, jplehr, sstefan1.
1> Fix wrong error message when compiling C souce code:
Currently emit error as following for uses_allocators(alloc(traits)):
called object type 'omp_allocator_handle_t' (aka
'enum omp_allocator_handle_t') is not a function or function pointer
To fix this, since "alloc" is Id expresison(spce 5.2), during the parser
(in ParseOpenMP.cpp), using tryParseCXXIdExpression instead of
ParseExpression for C.
2> Fix runtime problem when call to __kmpc_init_allocator
It seem load of traits.addr should be passed in rumtime call. Curently
the load of load traits.addr gets passed cause runtime to fail.
3> change to use CGF.EmitAutoVarAlloca(allocator) instead
CGF.EmitVarDecl(allocator).
If I understant correct, it is not necessary to emit variable init,
since "allocator" is private to target region.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D151517
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/Parse/ParseOpenMP.cpp
clang/test/OpenMP/target_parallel_for_simd_uses_allocators_codegen.cpp
clang/test/OpenMP/target_parallel_for_uses_allocators_codegen.cpp
clang/test/OpenMP/target_parallel_uses_allocators_codegen.cpp
clang/test/OpenMP/target_simd_uses_allocators_codegen.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_simd_uses_allocators_codegen.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_uses_allocators_codegen.cpp
clang/test/OpenMP/target_teams_distribute_simd_uses_allocators_codegen.cpp
clang/test/OpenMP/target_teams_distribute_uses_allocators_codegen.cpp
clang/test/OpenMP/target_teams_uses_allocators_codegen.cpp
clang/test/OpenMP/target_uses_allocators.c
clang/test/OpenMP/target_uses_allocators_codegen.cpp
openmp/libomptarget/test/mapping/target_uses_allocator.c
Index: openmp/libomptarget/test/mapping/target_uses_allocator.c
===================================================================
--- /dev/null
+++ openmp/libomptarget/test/mapping/target_uses_allocator.c
@@ -0,0 +1,56 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+#define N 1024
+
+int test_omp_aligned_alloc_on_device() {
+ int errors = 0;
+
+ omp_memspace_handle_t memspace = omp_default_mem_space;
+ omp_alloctrait_t traits[2] = {{omp_atk_alignment, 64}, {omp_atk_access, 64}};
+ omp_allocator_handle_t alloc =
+ omp_init_allocator(omp_default_mem_space, 1, traits);
+
+#pragma omp target map(tofrom : errors) uses_allocators(alloc(traits))
+ {
+ int *x;
+ int not_correct_array_values = 0;
+
+ x = (int *)omp_aligned_alloc(64, N * sizeof(int), alloc);
+ if (x == NULL) {
+ errors++;
+ } else {
+#pragma omp parallel for simd simdlen(16) aligned(x : 64)
+ for (int i = 0; i < N; i++) {
+ x[i] = i;
+ }
+
+#pragma omp parallel for simd simdlen(16) aligned(x : 64)
+ for (int i = 0; i < N; i++) {
+ if (x[i] != i) {
+#pragma omp atomic write
+ not_correct_array_values = 1;
+ }
+ }
+ if (not_correct_array_values) {
+ errors++;
+ }
+ omp_free(x, alloc);
+ }
+ }
+
+ omp_destroy_allocator(alloc);
+
+ return errors;
+}
+
+int main() {
+ int errors = 0;
+ if (test_omp_aligned_alloc_on_device())
+ printf("FAILE\n");
+ else
+ // CHECK: PASSED
+ printf("PASSED\n");
+}
Index: clang/test/OpenMP/target_uses_allocators_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_uses_allocators_codegen.cpp
+++ clang/test/OpenMP/target_uses_allocators_codegen.cpp
@@ -78,8 +78,7 @@
// CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
// CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
// CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// 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]],
Index: clang/test/OpenMP/target_uses_allocators.c
===================================================================
--- clang/test/OpenMP/target_uses_allocators.c
+++ clang/test/OpenMP/target_uses_allocators.c
@@ -2,6 +2,8 @@
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -verify -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -verify -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t %s -emit-llvm -o - | FileCheck %s
#ifndef HEADER
#define HEADER
@@ -19,9 +21,27 @@
KMP_ALLOCATOR_MAX_HANDLE = __UINTPTR_MAX__
} omp_allocator_handle_t;
+typedef enum omp_alloctrait_key_t { omp_atk_sync_hint = 1,
+ omp_atk_alignment = 2,
+ omp_atk_access = 3,
+ omp_atk_pool_size = 4,
+ omp_atk_fallback = 5,
+ omp_atk_fb_data = 6,
+ omp_atk_pinned = 7,
+ omp_atk_partition = 8
+} omp_alloctrait_key_t;
+
+typedef struct omp_alloctrait_t {
+ omp_alloctrait_key_t key;
+ __UINTPTR_TYPE__ value;
+} omp_alloctrait_t;
+
+
// CHECK: define {{.*}}[[FIE:@.+]]()
void fie(void) {
int x;
+ omp_allocator_handle_t my_allocator;
+ omp_alloctrait_t traits[10];
#pragma omp target uses_allocators(omp_null_allocator) allocate(omp_null_allocator: x) firstprivate(x)
{}
#pragma omp target uses_allocators(omp_default_mem_alloc) allocate(omp_default_mem_alloc: x) firstprivate(x)
@@ -40,6 +60,8 @@
{}
#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))
+ {}
}
#endif
@@ -106,3 +128,15 @@
// 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 inttoptr (i64 8 to ptr))
+
+// CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
+// CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
+// CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
+// 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]])
Index: clang/test/OpenMP/target_teams_uses_allocators_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_teams_uses_allocators_codegen.cpp
+++ clang/test/OpenMP/target_teams_uses_allocators_codegen.cpp
@@ -78,8 +78,7 @@
// CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
// CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
// CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// 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]],
Index: clang/test/OpenMP/target_teams_distribute_uses_allocators_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_uses_allocators_codegen.cpp
+++ clang/test/OpenMP/target_teams_distribute_uses_allocators_codegen.cpp
@@ -79,8 +79,7 @@
// CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
// CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
// CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// 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]],
Index: clang/test/OpenMP/target_teams_distribute_simd_uses_allocators_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_simd_uses_allocators_codegen.cpp
+++ clang/test/OpenMP/target_teams_distribute_simd_uses_allocators_codegen.cpp
@@ -79,8 +79,7 @@
// CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
// CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
// CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// 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]],
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_uses_allocators_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_uses_allocators_codegen.cpp
+++ clang/test/OpenMP/target_teams_distribute_parallel_for_uses_allocators_codegen.cpp
@@ -79,8 +79,7 @@
// CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
// CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
// CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// 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]],
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_simd_uses_allocators_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_simd_uses_allocators_codegen.cpp
+++ clang/test/OpenMP/target_teams_distribute_parallel_for_simd_uses_allocators_codegen.cpp
@@ -79,8 +79,7 @@
// CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
// CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
// CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// 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]],
Index: clang/test/OpenMP/target_simd_uses_allocators_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_simd_uses_allocators_codegen.cpp
+++ clang/test/OpenMP/target_simd_uses_allocators_codegen.cpp
@@ -79,8 +79,7 @@
// CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
// CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
// CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// 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]],
Index: clang/test/OpenMP/target_parallel_uses_allocators_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_parallel_uses_allocators_codegen.cpp
+++ clang/test/OpenMP/target_parallel_uses_allocators_codegen.cpp
@@ -78,8 +78,7 @@
// CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
// CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
// CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// 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]],
Index: clang/test/OpenMP/target_parallel_for_uses_allocators_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_parallel_for_uses_allocators_codegen.cpp
+++ clang/test/OpenMP/target_parallel_for_uses_allocators_codegen.cpp
@@ -79,8 +79,7 @@
// CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
// CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
// CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// 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]],
Index: clang/test/OpenMP/target_parallel_for_simd_uses_allocators_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_parallel_for_simd_uses_allocators_codegen.cpp
+++ clang/test/OpenMP/target_parallel_for_simd_uses_allocators_codegen.cpp
@@ -79,8 +79,7 @@
// CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
// CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
// CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// 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]],
Index: clang/lib/Parse/ParseOpenMP.cpp
===================================================================
--- clang/lib/Parse/ParseOpenMP.cpp
+++ clang/lib/Parse/ParseOpenMP.cpp
@@ -3103,8 +3103,13 @@
return nullptr;
SmallVector<Sema::UsesAllocatorsData, 4> Data;
do {
+ CXXScopeSpec SS;
+ Token Replacement;
ExprResult Allocator =
- getLangOpts().CPlusPlus ? ParseCXXIdExpression() : ParseExpression();
+ getLangOpts().CPlusPlus
+ ? ParseCXXIdExpression()
+ : tryParseCXXIdExpression(SS, /*isAddressOfOperand=*/false,
+ Replacement);
if (Allocator.isInvalid()) {
SkipUntil(tok::comma, tok::r_paren, tok::annot_pragma_openmp_end,
StopBeforeMatch);
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -6034,15 +6034,14 @@
AllocatorTraitsLVal = CGF.MakeAddrLValue(Addr, CGF.getContext().VoidPtrTy,
AllocatorTraitsLVal.getBaseInfo(),
AllocatorTraitsLVal.getTBAAInfo());
- llvm::Value *Traits =
- CGF.EmitLoadOfScalar(AllocatorTraitsLVal, AllocatorTraits->getExprLoc());
+ 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 =
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
