Author: abataev
Date: Mon Apr 23 10:33:41 2018
New Revision: 330620
URL: http://llvm.org/viewvc/llvm-project?rev=330620&view=rev
Log:
[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.
Modified:
cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
cfe/trunk/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp?rev=330620&r1=330619&r2=330620&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp Mon Apr 23 10:33:41 2018
@@ -3037,7 +3037,8 @@ llvm::Function *CGOpenMPRuntimeNVPTX::cr
/*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());
}
Modified:
cfe/trunk/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp?rev=330620&r1=330619&r2=330620&view=diff
==============================================================================
---
cfe/trunk/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp
(original)
+++
cfe/trunk/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp
Mon Apr 23 10:33:41 2018
@@ -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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
