Alexey Bataev | 3f96fe6 | 2017-12-13 17:31:39 +0000 | [diff] [blame] | 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 | |
Alexey Bataev | a8a9153a | 2017-12-29 18:07:07 +0000 | [diff] [blame] | 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 | |
Alexey Bataev | 3f96fe6 | 2017-12-13 17:31:39 +0000 | [diff] [blame] | 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 | |
Alexey Bataev | 9a70017 | 2018-05-08 14:16:57 +0000 | [diff] [blame] | 48 | // TCHECK: define weak void @__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}}) |
Alexey Bataev | 3f96fe6 | 2017-12-13 17:31:39 +0000 | [diff] [blame] | 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 | |
Alexey Bataev | 9a70017 | 2018-05-08 14:16:57 +0000 | [diff] [blame] | 59 | // TCHECK: define weak void @__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}}) |
Alexey Bataev | 3f96fe6 | 2017-12-13 17:31:39 +0000 | [diff] [blame] | 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 | |
Alexey Bataev | 9a70017 | 2018-05-08 14:16:57 +0000 | [diff] [blame] | 72 | // TCHECK: define weak void @__omp_offloading_{{.+}}(i32*{{.+}} [[A:%.+]], i16*{{.+}} [[AA:%.+]]) |
Alexey Bataev | 3f96fe6 | 2017-12-13 17:31:39 +0000 | [diff] [blame] | 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 | |
Alexey Bataev | 9a70017 | 2018-05-08 14:16:57 +0000 | [diff] [blame] | 121 | // TCHECK: define weak void @__omp_offloading_{{.+}}(i32*{{.+}}, i16*{{.+}}, i8*{{.+}}, [10 x i32]*{{.+}}) |
Alexey Bataev | 3f96fe6 | 2017-12-13 17:31:39 +0000 | [diff] [blame] | 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 | |
Alexey Bataev | 9a70017 | 2018-05-08 14:16:57 +0000 | [diff] [blame] | 157 | // TCHECK: define weak void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i32*{{.+}}, i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i16*{{.+}}) |
Alexey Bataev | 3f96fe6 | 2017-12-13 17:31:39 +0000 | [diff] [blame] | 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 |
Alexey Bataev | 9a70017 | 2018-05-08 14:16:57 +0000 | [diff] [blame] | 209 | // TCHECK: define weak void @__omp_offloading_{{.+}}(i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}}, [10 x i32]*{{.+}}) |
Alexey Bataev | 3f96fe6 | 2017-12-13 17:31:39 +0000 | [diff] [blame] | 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 |