282 lines · cpp
1// Only test codegen on target side, as private clause does not require any action on the host side2// Test target codegen - host bc file has to be created first.3// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc4// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-645// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s6// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-647// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc8// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-329// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s10// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-3211 12// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc13// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s14// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s15// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s16// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc17// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s18// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s19// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s20// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}21 22// expected-no-diagnostics23#ifndef HEADER24#define HEADER25 26template<typename tx, typename ty>27struct TT{28 tx X;29 ty Y;30};31 32// TCHECK: [[TT:%.+]] = type { i64, i8 }33// TCHECK: [[S1:%.+]] = type { double }34 35int 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 protected void @__omp_offloading_{{.+}}(ptr {{[^,]+}})49 // TCHECK: [[DYN_PTR:%.+]] = alloca ptr50 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},51 // TCHECK-NOT: store {{.+}}, {{.+}} [[A]],52 // TCHECK: ret void53 54#pragma omp target private(a)55 {56 a = 1;57 }58 59 // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(ptr {{[^,]+}})60 // TCHECK: [[DYN_PTR:%.+]] = alloca ptr61 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},62 // TCHECK: store i{{[0-9]+}} 1, ptr [[A]],63 // TCHECK: ret void64 65#pragma omp target private(a, aa)66 {67 a = 1;68 aa = 1;69 }70 71 // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(ptr {{[^,]+}})72 // TCHECK: [[DYN_PTR:%.+]] = alloca ptr73 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},74 // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},75 // TCHECK: store i{{[0-9]+}} 1, ptr [[A]],76 // TCHECK: store i{{[0-9]+}} 1, ptr [[A2]],77 // TCHECK: ret void78 79 #pragma omp target private(a, b, bn, c, cn, d)80 {81 a = 1;82 b[2] = 1.0;83 bn[3] = 1.0;84 c[1][2] = 1.0;85 cn[1][3] = 1.0;86 d.X = 1;87 d.Y = 1;88 }89 // make sure that private variables are generated in all cases and that we use those instances for operations inside the90 // target region91 // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(ptr {{[^,]+}}, i{{[0-9]+}} noundef [[VLA:%.+]], i{{[0-9]+}} noundef [[VLA1:%.+]], i{{[0-9]+}} noundef [[VLA3:%.+]])92 // TCHECK: [[DYN_PTR:%.+]] = alloca ptr93 // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},94 // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}},95 // TCHECK: [[VLA_ADDR4:%.+]] = alloca i{{[0-9]+}},96 // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},97 // TCHECK: [[B:%.+]] = alloca [10 x float],98 // TCHECK: [[SSTACK:%.+]] = alloca ptr,99 // TCHECK: [[C:%.+]] = alloca [5 x [10 x double]],100 // TCHECK: [[D:%.+]] = alloca [[TT]],101 // TCHECK: store i{{[0-9]+}} [[VLA]], ptr [[VLA_ADDR]],102 // TCHECK: store i{{[0-9]+}} [[VLA1]], ptr [[VLA_ADDR2]],103 // TCHECK: store i{{[0-9]+}} [[VLA3]], ptr [[VLA_ADDR4]],104 // TCHECK: [[VLA_ADDR_REF:%.+]] = load i{{[0-9]+}}, ptr [[VLA_ADDR]],105 // TCHECK: [[VLA_ADDR_REF2:%.+]] = load i{{[0-9]+}}, ptr [[VLA_ADDR2]],106 // TCHECK: [[VLA_ADDR_REF4:%.+]] = load i{{[0-9]+}}, ptr [[VLA_ADDR4]],107 // TCHECK: [[RET_STACK:%.+]] = call ptr @llvm.stacksave.p0()108 // TCHECK: store ptr [[RET_STACK]], ptr [[SSTACK]],109 // TCHECK: [[VLA5:%.+]] = alloca float, i{{[0-9]+}} [[VLA_ADDR_REF]],110 // TCHECK: [[VLA6_SIZE:%.+]] = mul{{.+}} i{{[0-9]+}} [[VLA_ADDR_REF2]], [[VLA_ADDR_REF4]]111 // TCHECK: [[VLA6:%.+]] = alloca double, i{{[0-9]+}} [[VLA6_SIZE]],112 113 // a = 1114 // TCHECK: store i{{[0-9]+}} 1, ptr [[A]],115 116 // b[2] = 1.0117 // TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x float], ptr [[B]], i{{[0-9]+}} 0, i{{[0-9]+}} 2118 // TCHECK: store float 1.0{{.*}}, ptr [[B_GEP]],119 120 // bn[3] = 1.0121 // TCHECK: [[BN_GEP:%.+]] = getelementptr inbounds float, ptr [[VLA5]], i{{[0-9]+}} 3122 // TCHECK: store float 1.0{{.*}}, ptr [[BN_GEP]],123 124 // c[1][2] = 1.0125 // TCHECK: [[C_GEP1:%.+]] = getelementptr inbounds [5 x [10 x double]], ptr [[C]], i{{[0-9]+}} 0, i{{[0-9]+}} 1126 // TCHECK: [[C_GEP2:%.+]] = getelementptr inbounds [10 x double], ptr [[C_GEP1]], i{{[0-9]+}} 0, i{{[0-9]+}} 2127 // TCHECK: store double 1.0{{.*}}, ptr [[C_GEP2]],128 129 // cn[1][3] = 1.0130 // TCHECK: [[CN_IND:%.+]] = mul{{.+}} i{{[0-9]+}} 1, [[VLA_ADDR_REF4]]131 // TCHECK: [[CN_GEP_IND:%.+]] = getelementptr inbounds double, ptr [[VLA6]], i{{[0-9]+}} [[CN_IND]]132 // TCHECK: [[CN_GEP_3:%.+]] = getelementptr inbounds double, ptr [[CN_GEP_IND]], i{{[0-9]+}} 3133 // TCHECK: store double 1.0{{.*}}, ptr [[CN_GEP_3]],134 135 // d.X = 1136 // [[X_FIELD:%.+]] = getelementptr inbounds nuw [[TT]] ptr [[D]], i{{[0-9]+}} 0, i{{[0-9]+}} 0137 // store i{{[0-9]+}} 1, ptr [[X_FIELD]],138 139 // d.Y = 1140 // [[Y_FIELD:%.+]] = getelementptr inbounds nuw [[TT]] ptr [[D]], i{{[0-9]+}} 0, i{{[0-9]+}} 1141 // store i{{[0-9]+}} 1, ptr [[Y_FIELD]],142 143 // finish144 // [[RELOAD_SSTACK:%.+]] = load ptr, ptr [[SSTACK]],145 // call ovid @llvm.stackrestore.p0(ptr [[RELOAD_SSTACK]])146 // ret void147 148 return a;149}150 151 152template<typename tx>153tx ftemplate(int n) {154 tx a = 0;155 short aa = 0;156 tx b[10];157 158#pragma omp target private(a,aa,b)159 {160 a = 1;161 aa = 1;162 b[2] = 1;163 }164 165 return a;166}167 168static169int fstatic(int n) {170 int a = 0;171 short aa = 0;172 char aaa = 0;173 int b[10];174 175#pragma omp target private(a,aa,aaa,b)176 {177 a = 1;178 aa = 1;179 aaa = 1;180 b[2] = 1;181 }182 183 return a;184}185 186// TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(ptr {{[^,]+}})187// TCHECK: [[DYN_PTR:%.+]] = alloca ptr188// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},189// TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},190// TCHECK: [[A3:%.+]] = alloca i{{[0-9]+}},191// TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}],192// TCHECK: store i{{[0-9]+}} 1, ptr [[A]],193// TCHECK: store i{{[0-9]+}} 1, ptr [[A2]],194// TCHECK: store i{{[0-9]+}} 1, ptr [[A3]],195// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], ptr [[B]], i{{[0-9]+}} 0, i{{[0-9]+}} 2196// TCHECK: store i{{[0-9]+}} 1, ptr [[B_GEP]],197// TCHECK: ret void198 199struct S1 {200 double a;201 202 int r1(int n){203 int b = n+1;204 short int c[2][n];205 206#pragma omp target private(b,c)207 {208 this->a = (double)b + 1.5;209 c[1][1] = ++a;210 }211 212 return c[1][1] + (int)b;213 }214 215 // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(ptr {{[^,]+}}, ptr noundef [[TH:%.+]], i{{[0-9]+}} noundef [[VLA:%.+]], i{{[0-9]+}} noundef [[VLA1:%.+]])216 // TCHECK: [[DYN_PTR:%.+]] = alloca ptr217 // TCHECK: [[TH_ADDR:%.+]] = alloca ptr,218 // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},219 // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}},220 // TCHECK: [[B:%.+]] = alloca i{{[0-9]+}},221 // TCHECK: [[SSTACK:%.+]] = alloca ptr,222 // TCHECK: store ptr [[TH]], ptr [[TH_ADDR]],223 // TCHECK: store i{{[0-9]+}} [[VLA]], ptr [[VLA_ADDR]],224 // TCHECK: store i{{[0-9]+}} [[VLA1]], ptr [[VLA_ADDR2]],225 // TCHECK: [[TH_ADDR_REF:%.+]] = load ptr, ptr [[TH_ADDR]],226 // TCHECK: [[VLA_ADDR_REF:%.+]] = load i{{[0-9]+}}, ptr [[VLA_ADDR]],227 // TCHECK: [[VLA_ADDR_REF2:%.+]] = load i{{[0-9]+}}, ptr [[VLA_ADDR2]],228 // TCHECK: [[RET_STACK:%.+]] = call ptr @llvm.stacksave.p0()229 // TCHECK: store ptr [[RET_STACK:%.+]], ptr [[SSTACK]],230 231 // this->a = (double)b + 1.5;232 // TCHECK: [[VLA_IND:%.+]] = mul{{.+}} i{{[0-9]+}} [[VLA_ADDR_REF]], [[VLA_ADDR_REF2]]233 // TCHECK: [[VLA3:%.+]] = alloca i{{[0-9]+}}, i{{[0-9]+}} [[VLA_IND]],234 // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, ptr [[B]],235 // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double236 // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00237 // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds nuw [[S1]], ptr [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0238 // TCHECK: store double [[NEW_A_VAL]], ptr [[A_FIELD]],239 240 // c[1][1] = ++a;241 // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds nuw [[S1]], ptr [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0242 // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, ptr [[A_FIELD4]],243 // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00244 // TCHECK: store double [[A_FIELD_INC]], ptr [[A_FIELD4]],245 // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}}246 // TCHECK: [[C_IND:%.+]] = mul{{.+}} i{{[0-9]+}} 1, [[VLA_ADDR_REF2]]247 // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, ptr [[VLA3]], i{{[0-9]+}} [[C_IND]]248 // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, ptr [[C_1_REF]], i{{[0-9]+}} 1249 // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], ptr [[C_1_1_REF]],250 251 // finish252 // TCHECK: [[RELOAD_SSTACK:%.+]] = load ptr, ptr [[SSTACK]],253 // TCHECK: call void @llvm.stackrestore.p0(ptr [[RELOAD_SSTACK]])254 // TCHECK: ret void255};256 257 258int bar(int n){259 int a = 0;260 a += foo(n);261 S1 S;262 a += S.r1(n);263 a += fstatic(n);264 a += ftemplate<int>(n);265 266 return a;267}268 269// template270// TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(ptr {{[^,]+}})271// TCHECK: [[DYN_PTR:%.+]] = alloca ptr272// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},273// TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},274// TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}],275// TCHECK: store i{{[0-9]+}} 1, ptr [[A]],276// TCHECK: store i{{[0-9]+}} 1, ptr [[A2]],277// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], ptr [[B]], i{{[0-9]+}} 0, i{{[0-9]+}} 2278// TCHECK: store i{{[0-9]+}} 1, ptr [[B_GEP]],279// TCHECK: ret void280 281#endif282