Author: abataev
Date: Wed Sep 13 13:20:59 2017
New Revision: 313186
URL: http://llvm.org/viewvc/llvm-project?rev=313186&view=rev
Log:
[OPENMP] Fix types for the target specific parameters in debug mode.
Used incorrect types for target specific parameters in debug mode,
should use original pointers rather than the pointee types.
Modified:
cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
cfe/trunk/test/OpenMP/target_parallel_debug_codegen.cpp
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp?rev=313186&r1=313185&r2=313186&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp Wed Sep 13 13:20:59 2017
@@ -2274,9 +2274,9 @@ CGOpenMPRuntimeNVPTX::getParameterAddres
QualType NativePointeeTy = cast<ReferenceType>(NonQualTy)->getPointeeType();
unsigned NativePointeeAddrSpace =
NativePointeeTy.getQualifiers().getAddressSpace();
- QualType TargetPointeeTy = TargetParam->getType()->getPointeeType();
+ QualType TargetTy = TargetParam->getType();
llvm::Value *TargetAddr = CGF.EmitLoadOfScalar(
- LocalAddr, /*Volatile=*/false, TargetPointeeTy, SourceLocation());
+ LocalAddr, /*Volatile=*/false, TargetTy, SourceLocation());
// First cast to generic.
TargetAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
TargetAddr, TargetAddr->getType()->getPointerElementType()->getPointerTo(
@@ -2287,7 +2287,7 @@ CGOpenMPRuntimeNVPTX::getParameterAddres
NativePointeeAddrSpace));
Address NativeParamAddr = CGF.CreateMemTemp(NativeParamType);
CGF.EmitStoreOfScalar(TargetAddr, NativeParamAddr, /*Volatile=*/false,
- NativeParam->getType());
+ NativeParamType);
return NativeParamAddr;
}
Modified: cfe/trunk/test/OpenMP/target_parallel_debug_codegen.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_parallel_debug_codegen.cpp?rev=313186&r1=313185&r2=313186&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_parallel_debug_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/target_parallel_debug_codegen.cpp Wed Sep 13 13:20:59
2017
@@ -5,11 +5,13 @@
int main() {
/* int(*b)[a]; */
/* int *(**c)[a]; */
+ bool bb;
int a;
int b[10][10];
int c[10][10][10];
-#pragma omp target parallel firstprivate(a, b) map(tofrom \
- : c)
+#pragma omp target parallel firstprivate(a, b) map(tofrom \
+ : c) map(tofrom \
+ : bb)
{
int &f = c[1][1][1];
int &g = a;
@@ -19,9 +21,11 @@ int main() {
b[0][a] = 10;
c[0][0][a] = 11;
b[0][a] = c[0][0][a];
+ bb |= b[0][a];
}
-#pragma omp target parallel firstprivate(a) map(tofrom \
- : c, b)
+#pragma omp target parallel firstprivate(a) map(tofrom \
+ : c, b) map(to \
+ : bb)
{
int &f = c[1][1][1];
int &g = a;
@@ -31,9 +35,11 @@ int main() {
b[0][a] = 10;
c[0][0][a] = 11;
b[0][a] = c[0][0][a];
+ d = bb;
}
-#pragma omp target parallel map(tofrom \
- : a, c, b)
+#pragma omp target parallel map(tofrom \
+ : a, c, b) map(from \
+ : bb)
{
int &f = c[1][1][1];
int &g = a;
@@ -43,64 +49,65 @@ int main() {
b[0][a] = 10;
c[0][0][a] = 11;
b[0][a] = c[0][0][a];
+ bb = b[0][a];
}
return 0;
}
-// CHECK: define internal void @__omp_offloading{{[^(]+}}([10 x [10 x [10 x
i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^)]+}})
+// CHECK: define internal void @__omp_offloading{{[^(]+}}([10 x [10 x [10 x
i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8
addrspace(1)* noalias{{[^)]+}})
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10
x [10 x [10 x i32]]]*
-// CHECK: call void [[NONDEBUG_WRAPPER:.+]](i32* {{[^,]+}}, i32* {{[^,]+}},
[10 x [10 x [10 x i32]]]* {{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]*
{{[^)]+}})
+// CHECK: call void [[NONDEBUG_WRAPPER:.+]](i32* {{[^,]+}}, i32* {{[^,]+}},
[10 x [10 x [10 x i32]]]* {{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]*
{{[^,]+}}, i8* {{[^)]+}})
-// CHECK: define internal void [[DEBUG_PARALLEL:@.+]](i32* {{[^,]+}}, i32*
{{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32
{{[^,]+}}, [10 x [10 x i32]]* noalias{{[^)]+}})
+// CHECK: define internal void [[DEBUG_PARALLEL:@.+]](i32* {{[^,]+}}, i32*
{{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32
{{[^,]+}}, [10 x [10 x i32]]* noalias{{[^,]+}}, i8 addrspace(1)*
noalias{{[^)]+}})
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10
x [10 x [10 x i32]]]*
-// CHECK: define internal void [[NONDEBUG_WRAPPER]](i32* {{[^,]+}}, i32*
{{[^,]+}}, [10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}},
[10 x [10 x i32]]* dereferenceable{{[^)]+}})
+// CHECK: define internal void [[NONDEBUG_WRAPPER]](i32* {{[^,]+}}, i32*
{{[^,]+}}, [10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}},
[10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x
i32]]] addrspace(1)*
-// CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x
[10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]*
{{[^)]+}})
+// CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x
[10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]*
{{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
-// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]*
dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]*
dereferenceable{{[^)]+}})
+// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]*
dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]*
dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x
i32]]] addrspace(1)*
-// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]
addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^)]+}})
+// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]
addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8
addrspace(1)* {{[^)]+}})
-// CHECK: define internal void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x
i32]]] addrspace(1)* noalias {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]
addrspace(1)* noalias {{[^)]+}})
+// CHECK: define internal void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x
i32]]] addrspace(1)* noalias{{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]
addrspace(1)* noalias{{[^,]+}}, i8 addrspace(1)* noalias{{[^)]+}})
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10
x [10 x [10 x i32]]]*
// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x
i32]]*
-// CHECK: call void [[NONDEBUG_WRAPPER:.+]](i32* {{[^,]+}}, i32* {{[^,]+}},
[10 x [10 x [10 x i32]]]* {{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]*
{{[^)]+}})
+// CHECK: call void [[NONDEBUG_WRAPPER:.+]](i32* {{[^,]+}}, i32* {{[^,]+}},
[10 x [10 x [10 x i32]]]* {{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]*
{{[^,]+}}, i8* {{[^)]+}})
-// CHECK: define internal void [[DEBUG_PARALLEL:@.+]](i32* {{[^,]+}}, i32*
{{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32
{{[^,]+}}, [10 x [10 x i32]] addrspace(1)* noalias{{[^)]+}})
+// CHECK: define internal void [[DEBUG_PARALLEL:@.+]](i32* {{[^,]+}}, i32*
{{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32
{{[^,]+}}, [10 x [10 x i32]] addrspace(1)* noalias{{[^,]+}}, i8 addrspace(1)*
{{[^)]+}})
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10
x [10 x [10 x i32]]]*
// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x
i32]]*
-// CHECK: define internal void [[NONDEBUG_WRAPPER]](i32* {{[^,]+}}, i32*
{{[^,]+}}, [10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}},
[10 x [10 x i32]]* dereferenceable{{[^)]+}})
+// CHECK: define internal void [[NONDEBUG_WRAPPER]](i32* {{[^,]+}}, i32*
{{[^,]+}}, [10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}},
[10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x
i32]]] addrspace(1)*
// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]]
addrspace(1)*
-// CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x
[10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]
addrspace(1)* {{[^)]+}})
+// CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x
[10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]
addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
-// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]*
dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]*
dereferenceable{{[^)]+}})
+// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]*
dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]*
dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x
i32]]] addrspace(1)*
// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]]
addrspace(1)*
-// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]
addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)*
{{[^)]+}})
+// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]
addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)*
{{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
-// CHECK: define internal void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x
i32]]] addrspace(1)* noalias {{[^,]+}}, i32 addrspace(1)* noalias {{[^,]+}},
[10 x [10 x i32]] addrspace(1)* noalias {{[^)]+}})
+// CHECK: define internal void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x
i32]]] addrspace(1)* noalias{{[^,]+}}, i32 addrspace(1)* noalias{{[^,]+}}, [10
x [10 x i32]] addrspace(1)* noalias{{[^,]+}}, i8 addrspace(1)* noalias{{[^)]+}})
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10
x [10 x [10 x i32]]]*
// CHECK: addrspacecast i32 addrspace(1)* %{{.+}} to i32*
// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x
i32]]*
-// CHECK: call void [[NONDEBUG_WRAPPER:@.+]](i32* {{[^,]+}}, i32* {{[^,]+}},
[10 x [10 x [10 x i32]]]* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x i32]]*
{{[^)]+}})
+// CHECK: call void [[NONDEBUG_WRAPPER:@.+]](i32* {{[^,]+}}, i32* {{[^,]+}},
[10 x [10 x [10 x i32]]]* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x i32]]*
{{[^,]+}}, i8* {{[^)]+}})
-// CHECK: define internal void [[DEBUG_PARALLEL:@.+]](i32* {{[^,]+}}, i32*
{{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32
addrspace(1)* noalias{{[^,]+}}, [10 x [10 x i32]] addrspace(1)*
noalias{{[^)]+}})
+// CHECK: define internal void [[DEBUG_PARALLEL:@.+]](i32* {{[^,]+}}, i32*
{{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32
addrspace(1)* noalias{{[^,]+}}, [10 x [10 x i32]] addrspace(1)*
noalias{{[^,]+}}, i8 addrspace(1)* noalias{{[^)]+}})
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10
x [10 x [10 x i32]]]*
// CHECK: addrspacecast i32 addrspace(1)* %{{.+}} to i32*
// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x
i32]]*
-// CHECK: define internal void [[NONDEBUG_WRAPPER]](i32* {{[^,]+}}, i32*
{{[^,]+}}, [10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i32*
dereferenceable{{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^)]+}})
+// CHECK: define internal void [[NONDEBUG_WRAPPER]](i32* {{[^,]+}}, i32*
{{[^,]+}}, [10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i32*
dereferenceable{{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8*
dereferenceable{{[^)]+}})
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x
i32]]] addrspace(1)*
// CHECK: addrspacecast i32* %{{.+}} to i32 addrspace(1)*
// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]]
addrspace(1)*
-// CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x
[10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 addrspace(1)* {{[^,]+}}, [10 x
[10 x i32]] addrspace(1)* {{[^)]+}})
+// CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x
[10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 addrspace(1)* {{[^,]+}}, [10 x
[10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
-// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]*
dereferenceable{{[^,]+}}, i32* dereferenceable{{[^,]+}}, [10 x [10 x i32]]*
dereferenceable{{[^)]+}})
+// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]*
dereferenceable{{[^,]+}}, i32* dereferenceable{{[^,]+}}, [10 x [10 x i32]]*
dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x
i32]]] addrspace(1)*
// CHECK: addrspacecast i32* %{{.+}} to i32 addrspace(1)*
// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]]
addrspace(1)*
-// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]
addrspace(1)* {{[^,]+}}, i32 addrspace(1)* {{[^,]+}}, [10 x [10 x i32]]
addrspace(1)* {{[^)]+}})
+// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]
addrspace(1)* {{[^,]+}}, i32 addrspace(1)* {{[^,]+}}, [10 x [10 x i32]]
addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
