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 // RUN: %clang_cc1 -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 -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 -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 -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 -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 -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 -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 -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 TT<tx, ty> operator*(const TT<tx, ty> &) { return *this; } 31 }; 32 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 reduction(*:a) 45 { 46 } 47 48 // TCHECK: define weak void @__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}}) 49 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*, 50 // TCHECK: store {{.+}}, {{.+}} [[A]], 51 // TCHECK: load i32*, i32** [[A]], 52 // TCHECK: ret void 53 54 #pragma omp target reduction(+:a) 55 { 56 a = 1; 57 } 58 59 // TCHECK: define weak void @__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}}) 60 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*, 61 // TCHECK: store {{.+}}, {{.+}} [[A]], 62 // TCHECK: [[REF:%.+]] = load i32*, i32** [[A]], 63 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[REF]], 64 // TCHECK: ret void 65 66 #pragma omp target reduction(-:a, aa) 67 { 68 a = 1; 69 aa = 1; 70 } 71 72 // TCHECK: define weak void @__omp_offloading_{{.+}}(i32*{{.+}} [[A:%.+]], i16*{{.+}} [[AA:%.+]]) 73 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*, 74 // TCHECK: [[AA:%.+]] = alloca i{{[0-9]+}}*, 75 // TCHECK: store {{.+}}, {{.+}} [[A]], 76 // TCHECK: store {{.+}}, {{.+}} [[AA]], 77 // TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]], 78 // TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]], 79 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]], 80 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]], 81 // TCHECK: ret void 82 83 return a; 84 } 85 86 87 template<typename tx> 88 tx ftemplate(int n) { 89 tx a = 0; 90 short aa = 0; 91 tx b[10]; 92 93 #pragma omp target reduction(+:a,aa,b) 94 { 95 a = 1; 96 aa = 1; 97 b[2] = 1; 98 } 99 100 return a; 101 } 102 103 static 104 int fstatic(int n) { 105 int a = 0; 106 short aa = 0; 107 char aaa = 0; 108 int b[10]; 109 110 #pragma omp target reduction(-:a,aa,aaa,b) 111 { 112 a = 1; 113 aa = 1; 114 aaa = 1; 115 b[2] = 1; 116 } 117 118 return a; 119 } 120 121 // TCHECK: define weak void @__omp_offloading_{{.+}}(i32*{{.+}}, i16*{{.+}}, i8*{{.+}}, [10 x i32]*{{.+}}) 122 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*, 123 // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}}*, 124 // TCHECK: [[A3:%.+]] = alloca i{{[0-9]+}}*, 125 // TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}]*, 126 // TCHECK: store {{.+}}, {{.+}} [[A]], 127 // TCHECK: store {{.+}}, {{.+}} [[A2]], 128 // TCHECK: store {{.+}}, {{.+}} [[A3]], 129 // TCHECK: store {{.+}}, {{.+}} [[B]], 130 // TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]], 131 // TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]], 132 // TCHECK: [[A3_REF:%.+]] = load i8*, i8** [[A3]], 133 // TCHECK: [[B_REF:%.+]] = load {{.+}}*, {{.+}}** [[B]], 134 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]], 135 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]], 136 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A3_REF]], 137 // TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 138 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]], 139 // TCHECK: ret void 140 141 struct S1 { 142 double a; 143 144 int r1(int n){ 145 int b = n+1; 146 short int c[2][n]; 147 148 #pragma omp target reduction(max:b,c) 149 { 150 this->a = (double)b + 1.5; 151 c[1][1] = ++a; 152 } 153 154 return c[1][1] + (int)b; 155 } 156 157 // TCHECK: define weak void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i32*{{.+}}, i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i16*{{.+}}) 158 // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*, 159 // TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}}*, 160 // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}}, 161 // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}}, 162 // TCHECK: [[C_ADDR:%.+]] = alloca i{{[0-9]+}}*, 163 // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]], 164 // TCHECK: store i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}** [[B_ADDR]], 165 // TCHECK: store i{{[0-9]+}} [[VLA]], i{{[0-9]+}}* [[VLA_ADDR]], 166 // TCHECK: store i{{[0-9]+}} [[VLA1]], i{{[0-9]+}}* [[VLA_ADDR2]], 167 // TCHECK: store i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}** [[C_ADDR]], 168 // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]], 169 // TCHECK: [[B_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[B_ADDR]], 170 // TCHECK: [[VLA_ADDR_REF:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR]], 171 // TCHECK: [[VLA_ADDR_REF2:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR2]], 172 // TCHECK: [[C_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[C_ADDR]], 173 174 // this->a = (double)b + 1.5; 175 // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_REF]], 176 // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double 177 // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00 178 // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 179 // TCHECK: store double [[NEW_A_VAL]], double* [[A_FIELD]], 180 181 // c[1][1] = ++a; 182 // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 183 // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, double* [[A_FIELD4]], 184 // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00 185 // TCHECK: store double [[A_FIELD_INC]], double* [[A_FIELD4]], 186 // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}} 187 // TCHECK: [[C_IND:%.+]] = mul{{.+}} i{{[0-9]+}} 1, [[VLA_ADDR_REF2]] 188 // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, i{{[0-9]+}}* [[C_REF]], i{{[0-9]+}} [[C_IND]] 189 // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, i{{[0-9]+}}* [[C_1_REF]], i{{[0-9]+}} 1 190 // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], i{{[0-9]+}}* [[C_1_1_REF]], 191 192 // finish 193 // TCHECK: ret void 194 }; 195 196 197 int bar(int n){ 198 int a = 0; 199 a += foo(n); 200 S1 S; 201 a += S.r1(n); 202 a += fstatic(n); 203 a += ftemplate<int>(n); 204 205 return a; 206 } 207 208 // template 209 // TCHECK: define weak void @__omp_offloading_{{.+}}(i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}}, [10 x i32]*{{.+}}) 210 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*, 211 // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}}*, 212 // TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}]*, 213 // TCHECK: store {{.+}}, {{.+}} [[A]], 214 // TCHECK: store {{.+}}, {{.+}} [[A2]], 215 // TCHECK: store {{.+}}, {{.+}} [[B]], 216 // TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]], 217 // TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]], 218 // TCHECK: [[B_REF:%.+]] = load {{.+}}*, {{.+}}** [[B]], 219 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]], 220 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]], 221 // TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 222 // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]], 223 // TCHECK: ret void 224 225 #endif 226