Gitiles
Code Review Sign In
gerrit-public.fairphone.software / toolchain / llvm-project / b45d43c397a956ad533829123f476be5c8d3c3ab / . / clang / test / OpenMP / target_parallel_debug_codegen.cpp
blob: 385f7a29cb8fd4f0271ab5c5fcdea4e55fa09acc [file] [log] [blame]
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]1// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
2// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -debug-info-kind=limited | FileCheck %s
3// expected-no-diagnostics
4
5int main() {
6 /* int(*b)[a]; */
7 /* int *(**c)[a]; */
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]8 bool bb;
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]9 int a;
10 int b[10][10];
11 int c[10][10][10];
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]12#pragma omp target parallel firstprivate(a, b) map(tofrom \
13 : c) map(tofrom \
Alexey Bataev7ba57af2017-10-17 16:47:34 +0000[diff] [blame]14 : bb) if (a)
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]15 {
16 int &f = c[1][1][1];
17 int &g = a;
18 int &h = b[1][1];
19 int d = 15;
20 a = 5;
21 b[0][a] = 10;
22 c[0][0][a] = 11;
23 b[0][a] = c[0][0][a];
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]24 bb |= b[0][a];
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]25 }
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]26#pragma omp target parallel firstprivate(a) map(tofrom \
27 : c, b) map(to \
28 : bb)
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]29 {
30 int &f = c[1][1][1];
31 int &g = a;
32 int &h = b[1][1];
33 int d = 15;
34 a = 5;
35 b[0][a] = 10;
36 c[0][0][a] = 11;
37 b[0][a] = c[0][0][a];
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]38 d = bb;
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]39 }
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]40#pragma omp target parallel map(tofrom \
41 : a, c, b) map(from \
42 : bb)
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]43 {
44 int &f = c[1][1][1];
45 int &g = a;
46 int &h = b[1][1];
47 int d = 15;
48 a = 5;
49 b[0][a] = 10;
50 c[0][0][a] = 11;
51 b[0][a] = c[0][0][a];
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]52 bb = b[0][a];
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]53 }
54 return 0;
55}
56
Alexey Bataev7ba57af2017-10-17 16:47:34 +0000[diff] [blame]57// 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{{[^,]+}}, i1 {{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]58// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]59// CHECK: call void [[NONDEBUG_WRAPPER:.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* {{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8* {{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]60
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]61// 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{{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]62// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
63
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]64// 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{{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]65// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]66// 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)* {{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]67
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]68// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]69// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]70// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]71
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]72// 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{{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]73// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
74// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]75// CHECK: call void [[NONDEBUG_WRAPPER:.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* {{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8* {{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]76
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]77// 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)* {{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]78// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
79// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
80
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]81// 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{{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]82// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
83// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]84// 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)* {{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]85
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]86// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]87// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
88// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]89// 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)* {{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]90
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]91// 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{{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]92// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
93// CHECK: addrspacecast i32 addrspace(1)* %{{.+}} to i32*
94// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]95// CHECK: call void [[NONDEBUG_WRAPPER:@.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8* {{[^)]+}})
Alexey Bataeve754b182017-08-09 19:38:53 +0000[diff] [blame]96
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]97// 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{{[^)]+}})
Alexey Bataeve754b182017-08-09 19:38:53 +0000[diff] [blame]98// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
99// CHECK: addrspacecast i32 addrspace(1)* %{{.+}} to i32*
100// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
101
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]102// 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{{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]103// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
104// CHECK: addrspacecast i32* %{{.+}} to i32 addrspace(1)*
105// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]106// 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)* {{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]107
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]108// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i32* dereferenceable{{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
Alexey Bataeve754b182017-08-09 19:38:53 +0000[diff] [blame]109// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
110// CHECK: addrspacecast i32* %{{.+}} to i32 addrspace(1)*
111// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
Alexey Bataev36f2c4d2017-09-13 20:20:59 +0000[diff] [blame]112// 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)* {{[^)]+}})
Alexey Bataev3b8d5582017-08-08 18:04:06 +0000[diff] [blame]113
Alexey Bataevb45d43c2017-11-22 16:02:03 +0000[diff] [blame^]114// CHECK: !DILocalVariable(name: ".global_tid.",
115// CHECK-SAME: DIFlagArtificial
116// CHECK: !DILocalVariable(name: ".bound_tid.",
117// CHECK-SAME: DIFlagArtificial
118// CHECK: !DILocalVariable(name: "c",
119// CHECK-SAMEi-NOT: DIFlagArtificial
120// CHECK: !DILocalVariable(name: "a",
121// CHECK-SAMEi-NOT: DIFlagArtificial
122// CHECK: !DILocalVariable(name: "b",
123// CHECK-SAMEi-NOT: DIFlagArtificial