1 // Only test codegen on target side, as private clause does not require any action on the host side 2 // Test target codegen - host bc file has to be created first. 3 // RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc 4 // RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -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 5 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s 6 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64 7 // RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc 8 // RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -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 9 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s 10 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32 11 12 // RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc 13 // RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s 14 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s 15 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s 16 // RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc 17 // RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s 18 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s 19 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s 20 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} 21 22 // expected-no-diagnostics 23 #ifndef HEADER 24 #define HEADER 25 26 template<typename tx, typename ty> 27 struct TT{ 28 tx X; 29 ty Y; 30 }; 31 32 // TCHECK: [[TT:%.+]] = type { i64, i8 } 33 // TCHECK: [[S1:%.+]] = type { double } 34 35 int foo(int n) { 36 int a = 0; 37 short aa = 0; 38 float b[10]; 39 float bn[n]; 40 double c[5][10]; 41 double cn[5][n]; 42 TT<long long, char> d; 43 44 #pragma omp target private(a) 45 { 46 } 47 48 // TCHECK: define weak_odr void @__omp_offloading_{{.+}}() 49 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}, 50 // TCHECK-NOT: store {{.+}}, {{.+}} [[A]], 51 // TCHECK: ret void 52 53 #pragma omp target private(a) 54 { 55 a = 1; 56 } 57 58 // TCHECK: define weak_odr void @__omp_offloading_{{.+}}() 59 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}, 60 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]], 61 // TCHECK: ret void 62 63 #pragma omp target private(a, aa) 64 { 65 a = 1; 66 aa = 1; 67 } 68 69 // TCHECK: define weak_odr void @__omp_offloading_{{.+}}() 70 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}, 71 // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}}, 72 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]], 73 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2]], 74 // TCHECK: ret void 75 76 #pragma omp target private(a, b, bn, c, cn, d) 77 { 78 a = 1; 79 b[2] = 1.0; 80 bn[3] = 1.0; 81 c[1][2] = 1.0; 82 cn[1][3] = 1.0; 83 d.X = 1; 84 d.Y = 1; 85 } 86 // make sure that private variables are generated in all cases and that we use those instances for operations inside the 87 // target region 88 // TCHECK: define weak_odr void @__omp_offloading_{{.+}}(i{{[0-9]+}} noundef [[VLA:%.+]], i{{[0-9]+}} noundef [[VLA1:%.+]], i{{[0-9]+}} noundef [[VLA3:%.+]]) 89 // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}}, 90 // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}}, 91 // TCHECK: [[VLA_ADDR4:%.+]] = alloca i{{[0-9]+}}, 92 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}, 93 // TCHECK: [[B:%.+]] = alloca [10 x float], 94 // TCHECK: [[SSTACK:%.+]] = alloca i8*, 95 // TCHECK: [[C:%.+]] = alloca [5 x [10 x double]], 96 // TCHECK: [[D:%.+]] = alloca [[TT]], 97 // TCHECK: store i{{[0-9]+}} [[VLA]], i{{[0-9]+}}* [[VLA_ADDR]], 98 // TCHECK: store i{{[0-9]+}} [[VLA1]], i{{[0-9]+}}* [[VLA_ADDR2]], 99 // TCHECK: store i{{[0-9]+}} [[VLA3]], i{{[0-9]+}}* [[VLA_ADDR4]], 100 // TCHECK: [[VLA_ADDR_REF:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR]], 101 // TCHECK: [[VLA_ADDR_REF2:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR2]], 102 // TCHECK: [[VLA_ADDR_REF4:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR4]], 103 // TCHECK: [[RET_STACK:%.+]] = call i8* @llvm.stacksave() 104 // TCHECK: store i8* [[RET_STACK]], i8** [[SSTACK]], 105 // TCHECK: [[VLA5:%.+]] = alloca float, i{{[0-9]+}} [[VLA_ADDR_REF]], 106 // TCHECK: [[VLA6_SIZE:%.+]] = mul{{.+}} i{{[0-9]+}} [[VLA_ADDR_REF2]], [[VLA_ADDR_REF4]] 107 // TCHECK: [[VLA6:%.+]] = alloca double, i{{[0-9]+}} [[VLA6_SIZE]], 108 109 // a = 1 110 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]], 111 112 // b[2] = 1.0 113 // TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x float], [10 x float]* [[B]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 114 // TCHECK: store float 1.0{{.*}}, float* [[B_GEP]], 115 116 // bn[3] = 1.0 117 // TCHECK: [[BN_GEP:%.+]] = getelementptr inbounds float, float* [[VLA5]], i{{[0-9]+}} 3 118 // TCHECK: store float 1.0{{.*}}, float* [[BN_GEP]], 119 120 // c[1][2] = 1.0 121 // TCHECK: [[C_GEP1:%.+]] = getelementptr inbounds [5 x [10 x double]], [5 x [10 x double]]* [[C]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 122 // TCHECK: [[C_GEP2:%.+]] = getelementptr inbounds [10 x double], [10 x double]* [[C_GEP1]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 123 // TCHECK: store double 1.0{{.*}}, double* [[C_GEP2]], 124 125 // cn[1][3] = 1.0 126 // TCHECK: [[CN_IND:%.+]] = mul{{.+}} i{{[0-9]+}} 1, [[VLA_ADDR_REF4]] 127 // TCHECK: [[CN_GEP_IND:%.+]] = getelementptr inbounds double, double* [[VLA6]], i{{[0-9]+}} [[CN_IND]] 128 // TCHECK: [[CN_GEP_3:%.+]] = getelementptr inbounds double, double* [[CN_GEP_IND]], i{{[0-9]+}} 3 129 // TCHECK: store double 1.0{{.*}}, double* [[CN_GEP_3]], 130 131 // d.X = 1 132 // [[X_FIELD:%.+]] = getelementptr inbounds [[TT]] [[TT]]* [[D]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 133 // store i{{[0-9]+}} 1, i{{[0-9]+}}* [[X_FIELD]], 134 135 // d.Y = 1 136 // [[Y_FIELD:%.+]] = getelementptr inbounds [[TT]] [[TT]]* [[D]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 137 // store i{{[0-9]+}} 1, i{{[0-9]+}}* [[Y_FIELD]], 138 139 // finish 140 // [[RELOAD_SSTACK:%.+]] = load i8*, i8** [[SSTACK]], 141 // call ovid @llvm.stackrestore(i8* [[RELOAD_SSTACK]]) 142 // ret void 143 144 return a; 145 } 146 147 148 template<typename tx> 149 tx ftemplate(int n) { 150 tx a = 0; 151 short aa = 0; 152 tx b[10]; 153 154 #pragma omp target private(a,aa,b) 155 { 156 a = 1; 157 aa = 1; 158 b[2] = 1; 159 } 160 161 return a; 162 } 163 164 static 165 int fstatic(int n) { 166 int a = 0; 167 short aa = 0; 168 char aaa = 0; 169 int b[10]; 170 171 #pragma omp target private(a,aa,aaa,b) 172 { 173 a = 1; 174 aa = 1; 175 aaa = 1; 176 b[2] = 1; 177 } 178 179 return a; 180 } 181 182 // TCHECK: define weak_odr void @__omp_offloading_{{.+}}() 183 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}, 184 // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}}, 185 // TCHECK: [[A3:%.+]] = alloca i{{[0-9]+}}, 186 // TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}], 187 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]], 188 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2]], 189 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A3]], 190 // TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 191 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]], 192 // TCHECK: ret void 193 194 struct S1 { 195 double a; 196 197 int r1(int n){ 198 int b = n+1; 199 short int c[2][n]; 200 201 #pragma omp target private(b,c) 202 { 203 this->a = (double)b + 1.5; 204 c[1][1] = ++a; 205 } 206 207 return c[1][1] + (int)b; 208 } 209 210 // TCHECK: define weak_odr void @__omp_offloading_{{.+}}([[S1]]* noundef [[TH:%.+]], i{{[0-9]+}} noundef [[VLA:%.+]], i{{[0-9]+}} noundef [[VLA1:%.+]]) 211 // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*, 212 // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}}, 213 // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}}, 214 // TCHECK: [[B:%.+]] = alloca i{{[0-9]+}}, 215 // TCHECK: [[SSTACK:%.+]] = alloca i8*, 216 // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]], 217 // TCHECK: store i{{[0-9]+}} [[VLA]], i{{[0-9]+}}* [[VLA_ADDR]], 218 // TCHECK: store i{{[0-9]+}} [[VLA1]], i{{[0-9]+}}* [[VLA_ADDR2]], 219 // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]], 220 // TCHECK: [[VLA_ADDR_REF:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR]], 221 // TCHECK: [[VLA_ADDR_REF2:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR2]], 222 // TCHECK: [[RET_STACK:%.+]] = call i8* @llvm.stacksave() 223 // TCHECK: store i8* [[RET_STACK:%.+]], i8** [[SSTACK]], 224 225 // this->a = (double)b + 1.5; 226 // TCHECK: [[VLA_IND:%.+]] = mul{{.+}} i{{[0-9]+}} [[VLA_ADDR_REF]], [[VLA_ADDR_REF2]] 227 // TCHECK: [[VLA3:%.+]] = alloca i{{[0-9]+}}, i{{[0-9]+}} [[VLA_IND]], 228 // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B]], 229 // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double 230 // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00 231 // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 232 // TCHECK: store double [[NEW_A_VAL]], double* [[A_FIELD]], 233 234 // c[1][1] = ++a; 235 // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 236 // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, double* [[A_FIELD4]], 237 // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00 238 // TCHECK: store double [[A_FIELD_INC]], double* [[A_FIELD4]], 239 // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}} 240 // TCHECK: [[C_IND:%.+]] = mul{{.+}} i{{[0-9]+}} 1, [[VLA_ADDR_REF2]] 241 // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, i{{[0-9]+}}* [[VLA3]], i{{[0-9]+}} [[C_IND]] 242 // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, i{{[0-9]+}}* [[C_1_REF]], i{{[0-9]+}} 1 243 // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], i{{[0-9]+}}* [[C_1_1_REF]], 244 245 // finish 246 // TCHECK: [[RELOAD_SSTACK:%.+]] = load i8*, i8** [[SSTACK]], 247 // TCHECK: call void @llvm.stackrestore(i8* [[RELOAD_SSTACK]]) 248 // TCHECK: ret void 249 }; 250 251 252 int bar(int n){ 253 int a = 0; 254 a += foo(n); 255 S1 S; 256 a += S.r1(n); 257 a += fstatic(n); 258 a += ftemplate<int>(n); 259 260 return a; 261 } 262 263 // template 264 // TCHECK: define weak_odr void @__omp_offloading_{{.+}}() 265 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}, 266 // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}}, 267 // TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}], 268 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]], 269 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2]], 270 // TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 271 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]], 272 // TCHECK: ret void 273 274 #endif 275