| 1 | // Test target codegen - host bc file has to be created first. |
| 2 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc |
| 3 | // RUN: %clang_cc1 -debug-info-kind=limited -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 - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64 |
| 4 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc |
| 5 | // RUN: %clang_cc1 -debug-info-kind=limited -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32 |
| 6 | // expected-no-diagnostics |
| 7 | #ifndef HEADER |
| 8 | #define HEADER |
| 9 | |
| 10 | template <typename tx, typename ty> |
| 11 | struct TT { |
| 12 | tx X; |
| 13 | ty Y; |
| 14 | }; |
| 15 | |
| 16 | // TCHECK-DAG: [[TT:%.+]] = type { i64, i8 } |
| 17 | // TCHECK-DAG: [[TTII:%.+]] = type { i32, i32 } |
| 18 | // TCHECK-DAG: [[S1:%.+]] = type { double } |
| 19 | |
| 20 | // TCHECK: @__omp_offloading_firstprivate__{{.+}}_e_l27 = internal addrspace(4) global [[TTII]] zeroinitializer |
| 21 | int foo(int n, double *ptr) { |
| 22 | int a = 0; |
| 23 | short aa = 0; |
| 24 | float b[10]; |
| 25 | double c[5][10]; |
| 26 | TT<long long, char> d; |
| 27 | const TT<int, int> e = {n, n}; |
| 28 | |
| 29 | #pragma omp target firstprivate(a, e) map(tofrom \ |
| 30 | : b) |
| 31 | { |
| 32 | b[a] = a; |
| 33 | b[a] += e.X; |
| 34 | } |
| 35 | |
| 36 | // TCHECK: define {{.*}}void @__omp_offloading_{{.+}}([10 x float] addrspace(1)* noalias [[B_IN:%.+]], i{{[0-9]+}} [[A_IN:%.+]], [[TTII]]* noalias [[E_IN:%.+]]) |
| 37 | // TCHECK-NOT: alloca [[TTII]], |
| 38 | // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}}, |
| 39 | // TCHECK-NOT: alloca [[TTII]], |
| 40 | // TCHECK-NOT: alloca i{{[0-9]+}}, |
| 41 | // TCHECK-64: call void @llvm.dbg.declare(metadata [10 x float] addrspace(1)** %{{.+}}, metadata !{{[0-9]+}}, metadata !DIExpression()) |
| 42 | // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]], |
| 43 | // TCHECK: ret void |
| 44 | |
| 45 | #pragma omp target firstprivate(aa, b, c, d) |
| 46 | { |
| 47 | aa += 1; |
| 48 | b[2] = 1.0; |
| 49 | c[1][2] = 1.0; |
| 50 | d.X = 1; |
| 51 | d.Y = 1; |
| 52 | } |
| 53 | |
| 54 | // make sure that firstprivate variables are generated in all cases and that we use those instances for operations inside the |
| 55 | // target region |
| 56 | // TCHECK: define {{.*}}void @__omp_offloading_{{.+}}(i{{[0-9]+}}{{.*}} [[A2_IN:%.+]], [10 x float]*{{.*}} [[B_IN:%.+]], [5 x [10 x double]]*{{.*}} [[C_IN:%.+]], [[TT]]*{{.*}} [[D_IN:%.+]]) |
| 57 | // TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}}, |
| 58 | // TCHECK: [[B_ADDR:%.+]] = alloca [10 x float]*, |
| 59 | // TCHECK: [[C_ADDR:%.+]] = alloca [5 x [10 x double]]*, |
| 60 | // TCHECK: [[D_ADDR:%.+]] = alloca [[TT]]*, |
| 61 | // TCHECK-NOT: alloca i{{[0-9]+}}, |
| 62 | // TCHECK: [[B_PRIV:%.+]] = alloca [10 x float], |
| 63 | // TCHECK: [[C_PRIV:%.+]] = alloca [5 x [10 x double]], |
| 64 | // TCHECK: [[D_PRIV:%.+]] = alloca [[TT]], |
| 65 | // TCHECK: store i{{[0-9]+}} [[A2_IN]], i{{[0-9]+}}* [[A2_ADDR]], |
| 66 | // TCHECK: store [10 x float]* [[B_IN]], [10 x float]** [[B_ADDR]], |
| 67 | // TCHECK: store [5 x [10 x double]]* [[C_IN]], [5 x [10 x double]]** [[C_ADDR]], |
| 68 | // TCHECK: store [[TT]]* [[D_IN]], [[TT]]** [[D_ADDR]], |
| 69 | // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x float]*, [10 x float]** [[B_ADDR]], |
| 70 | // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x float]*, [10 x float]** % |
| 71 | // TCHECK: [[C_ADDR_REF:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** [[C_ADDR]], |
| 72 | // TCHECK: [[C_ADDR_REF:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** % |
| 73 | // TCHECK: [[D_ADDR_REF:%.+]] = load [[TT]]*, [[TT]]** [[D_ADDR]], |
| 74 | // TCHECK: [[D_ADDR_REF:%.+]] = load [[TT]]*, [[TT]]** % |
| 75 | |
| 76 | // firstprivate(aa): a_priv = a_in |
| 77 | |
| 78 | // firstprivate(b): memcpy(b_priv,b_in) |
| 79 | // TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x float]* [[B_PRIV]] to i8* |
| 80 | // TCHECK: [[B_ADDR_REF_BCAST:%.+]] = bitcast [10 x float]* [[B_ADDR_REF]] to i8* |
| 81 | // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[B_PRIV_BCAST]], i8* align {{[0-9]+}} [[B_ADDR_REF_BCAST]], {{.+}}) |
| 82 | |
| 83 | // firstprivate(c) |
| 84 | // TCHECK: [[C_PRIV_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C_PRIV]] to i8* |
| 85 | // TCHECK: [[C_IN_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C_ADDR_REF]] to i8* |
| 86 | // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[C_PRIV_BCAST]], i8* align {{[0-9]+}} [[C_IN_BCAST]],{{.+}}) |
| 87 | |
| 88 | // firstprivate(d) |
| 89 | // TCHECK: [[D_PRIV_BCAST:%.+]] = bitcast [[TT]]* [[D_PRIV]] to i8* |
| 90 | // TCHECK: [[D_IN_BCAST:%.+]] = bitcast [[TT]]* [[D_ADDR_REF]] to i8* |
| 91 | // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[D_PRIV_BCAST]], i8* align {{[0-9]+}} [[D_IN_BCAST]],{{.+}}) |
| 92 | |
| 93 | // TCHECK: load i16, i16* [[A2_ADDR]], |
| 94 | |
| 95 | #pragma omp target firstprivate(ptr) |
| 96 | { |
| 97 | ptr[0]++; |
| 98 | } |
| 99 | |
| 100 | // TCHECK: define weak void @__omp_offloading_{{.+}}(double* [[PTR_IN:%.+]]) |
| 101 | // TCHECK: [[PTR_ADDR:%.+]] = alloca double*, |
| 102 | // TCHECK-NOT: alloca double*, |
| 103 | // TCHECK: store double* [[PTR_IN]], double** [[PTR_ADDR]], |
| 104 | // TCHECK: [[PTR_IN_REF:%.+]] = load double*, double** [[PTR_ADDR]], |
| 105 | // TCHECK-NOT: store double* [[PTR_IN_REF]], double** [[PTR_PRIV]], |
| 106 | |
| 107 | return a; |
| 108 | } |
| 109 | |
| 110 | template <typename tx> |
| 111 | tx ftemplate(int n) { |
| 112 | tx a = 0; |
| 113 | tx b[10]; |
| 114 | |
| 115 | #pragma omp target firstprivate(a, b) |
| 116 | { |
| 117 | a += 1; |
| 118 | b[2] += 1; |
| 119 | } |
| 120 | |
| 121 | return a; |
| 122 | } |
| 123 | |
| 124 | static int fstatic(int n) { |
| 125 | int a = 0; |
| 126 | char aaa = 0; |
| 127 | int b[10]; |
| 128 | |
| 129 | #pragma omp target firstprivate(a, aaa, b) |
| 130 | { |
| 131 | a += 1; |
| 132 | aaa += 1; |
| 133 | b[2] += 1; |
| 134 | } |
| 135 | |
| 136 | return a; |
| 137 | } |
| 138 | |
| 139 | // TCHECK: define {{.*}}void @__omp_offloading_{{.+}}(i{{[0-9]+}}{{.*}} [[A_IN:%.+]], i{{[0-9]+}}{{.*}} [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]]) |
| 140 | // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}}, |
| 141 | // TCHECK: [[A3_ADDR:%.+]] = alloca i{{[0-9]+}}, |
| 142 | // TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*, |
| 143 | // TCHECK-NOT: alloca i{{[0-9]+}}, |
| 144 | // TCHECK: [[B_PRIV:%.+]] = alloca [10 x i{{[0-9]+}}], |
| 145 | // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]], |
| 146 | // TCHECK: store i{{[0-9]+}} [[A3_IN]], i{{[0-9]+}}* [[A3_ADDR]], |
| 147 | // TCHECK: store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]], |
| 148 | // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]], |
| 149 | // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** % |
| 150 | |
| 151 | // firstprivate(a): a_priv = a_in |
| 152 | |
| 153 | // firstprivate(aaa) |
| 154 | |
| 155 | // TCHECK-NOT: store i{{[0-9]+}} %{{.+}}, i{{[0-9]+}}* |
| 156 | |
| 157 | // firstprivate(b) |
| 158 | // TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_PRIV]] to i8* |
| 159 | // TCHECK: [[B_IN_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8* |
| 160 | // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[B_PRIV_BCAST]], i8* align {{[0-9]+}} [[B_IN_BCAST]],{{.+}}) |
| 161 | |
| 162 | // TCHECK: ret void |
| 163 | |
| 164 | struct S1 { |
| 165 | double a; |
| 166 | |
| 167 | int r1(int n) { |
| 168 | int b = n + 1; |
| 169 | |
| 170 | #pragma omp target firstprivate(b) |
| 171 | { |
| 172 | this->a = (double)b + 1.5; |
| 173 | } |
| 174 | |
| 175 | return (int)b; |
| 176 | } |
| 177 | |
| 178 | // TCHECK: define internal void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}} [[B_IN:%.+]]) |
| 179 | // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*, |
| 180 | // TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}}, |
| 181 | // TCHECK-NOT: alloca i{{[0-9]+}}, |
| 182 | |
| 183 | // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]], |
| 184 | // TCHECK: store i{{[0-9]+}} [[B_IN]], i{{[0-9]+}}* [[B_ADDR]], |
| 185 | // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]], |
| 186 | // TCHECK-64: [[B_ADDR_CONV:%.+]] = bitcast i{{[0-9]+}}* [[B_ADDR]] to i{{[0-9]+}}* |
| 187 | |
| 188 | // firstprivate(b) |
| 189 | // TCHECK-NOT: store i{{[0-9]+}} %{{.+}}, i{{[0-9]+}}* |
| 190 | |
| 191 | // TCHECK: ret void |
| 192 | }; |
| 193 | |
| 194 | int bar(int n, double *ptr) { |
| 195 | int a = 0; |
| 196 | a += foo(n, ptr); |
| 197 | S1 S; |
| 198 | a += S.r1(n); |
| 199 | a += fstatic(n); |
| 200 | a += ftemplate<int>(n); |
| 201 | |
| 202 | return a; |
| 203 | } |
| 204 | |
| 205 | // template |
| 206 | |
| 207 | // TCHECK: define internal void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]]) |
| 208 | // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}}, |
| 209 | // TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*, |
| 210 | // TCHECK-NOT: alloca i{{[0-9]+}}, |
| 211 | // TCHECK: [[B_PRIV:%.+]] = alloca [10 x i{{[0-9]+}}], |
| 212 | // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]], |
| 213 | // TCHECK: store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]], |
| 214 | // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]], |
| 215 | // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** % |
| 216 | |
| 217 | // firstprivate(a) |
| 218 | // TCHECK-NOT: store i{{[0-9]+}} %{{.+}}, i{{[0-9]+}}* |
| 219 | |
| 220 | // firstprivate(b) |
| 221 | // TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_PRIV]] to i8* |
| 222 | // TCHECK: [[B_IN_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8* |
| 223 | // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[B_PRIV_BCAST]], i8* align {{[0-9]+}} [[B_IN_BCAST]],{{.+}}) |
| 224 | |
| 225 | // TCHECK: ret void |
| 226 | |
| 227 | #endif |
| 228 | |
| 229 | // TCHECK-DAG: distinct !DISubprogram(linkageName: "__omp_offloading_{{.+}}_worker", |
| 230 | // TCHECK-DAG: distinct !DISubprogram(linkageName: "__omp_offloading_{{.+}}_worker", |
| 231 | // TCHECK-DAG: distinct !DISubprogram(linkageName: "__omp_offloading_{{.+}}_worker", |
| 232 | // TCHECK-DAG: distinct !DISubprogram(linkageName: "__omp_offloading_{{.+}}_worker", |
| 233 | // TCHECK-DAG: distinct !DISubprogram(linkageName: "__omp_offloading_{{.+}}_worker", |
| 234 | |