maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[OPENMP] Do not cast captured by value variables with pointer types in 

NVPTX target.

When generating the wrapper function for the offloading region, we need
to call the outlined function and cast the arguments correctly to follow
the ABI. Usually, variables captured by value are casted to `uintptr_t`
type. But this should not performed for the variables with pointer type.

llvm-svn: 330620
This commit is contained in:
Alexey Bataev 2018-04-23 17:33:41 +00:00
parent 061fcd586d
commit 2091ca6c97
2 changed files with 15 additions and 5 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
View File

@ -3037,7 +3037,8 @@ llvm::Function *CGOpenMPRuntimeNVPTX::createParallelDataSharingWrapper(
/*Volatile=*/false,
CGFContext.getPointerType(ElemTy),
CI->getLocation());
if (CI->capturesVariableByCopy()) {
if (CI->capturesVariableByCopy() &&
!CI->getCapturedVar()->getType()->isAnyPointerType()) {
Arg = castValueToType(CGF, Arg, ElemTy, CGFContext.getUIntPtrType(),
CI->getLocation());
}

17
clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp
View File

@ -9,10 +9,11 @@
#define HEADER
// Check that the execution mode of all 2 target regions on the gpu is set to SPMD Mode.
// CHECK-DAG: {{@__omp_offloading_.+l30}}_exec_mode = weak constant i8 0
// CHECK-DAG: {{@__omp_offloading_.+l36}}_exec_mode = weak constant i8 0
// CHECK-DAG: {{@__omp_offloading_.+l41}}_exec_mode = weak constant i8 0
// CHECK-DAG: {{@__omp_offloading_.+l46}}_exec_mode = weak constant i8 0
// CHECK-DAG: {{@__omp_offloading_.+l32}}_exec_mode = weak constant i8 0
// CHECK-DAG: {{@__omp_offloading_.+l38}}_exec_mode = weak constant i8 0
// CHECK-DAG: {{@__omp_offloading_.+l43}}_exec_mode = weak constant i8 0
// CHECK-DAG: {{@__omp_offloading_.+l48}}_exec_mode = weak constant i8 0
// CHECK-DAG: {{@__omp_offloading_.+l56}}_exec_mode = weak constant i8 0
#define N 1000
#define M 10
@ -26,6 +27,7 @@ tx ftemplate(int n) {
tx f = n;
tx l;
int k;
tx *v;
#pragma omp target teams distribute parallel for lastprivate(l) dist_schedule(static,128) schedule(static,32)
for(int i = 0; i < n; i++) {
@ -51,6 +53,9 @@ tx ftemplate(int n) {
}
}
#pragma omp target teams distribute parallel for map(a, v[:N])
for(int i = 0; i < n; i++)
a[i] = v[i];
return a[0];
}
@ -120,4 +125,8 @@ int bar(int n){
// CHECK: call void @__kmpc_for_static_fini(
// CHECK: ret void
// CHECK: define void @__omp_offloading_{{.*}}_l56(i[[SZ:64|32]] %{{[^,]+}}, [1000 x i32]* dereferenceable{{.*}}, i32* %{{[^)]+}})
// CHECK: call void [[OUTLINED:@__omp_outlined.*]](i32* %{{.+}}, i32* %{{.+}}, i[[SZ]] %{{.*}}, i[[SZ]] %{{.*}}, i[[SZ]] %{{.*}}, [1000 x i32]* %{{.*}}, i32* %{{.*}})
// CHECK: define internal void [[OUTLINED]](i32* noalias %{{.*}}, i32* noalias %{{.*}} i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, [1000 x i32]* dereferenceable{{.*}}, i32* %{{.*}})
#endif