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