1 // expected-no-diagnostics 2 #ifndef HEADER 3 #define HEADER 4 // Test host codegen. 5 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix HCK_NO_TGT 6 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64 --check-prefix HCK1 --check-prefix HCK1-64 7 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 8 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64 --check-prefix HCK1 --check-prefix HCK1-64 9 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32 --check-prefix HCK1 --check-prefix HCK1-64 10 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 11 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32 --check-prefix HCK1 --check-prefix HCK1-64 12 13 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY1 14 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 15 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY1 16 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY1 17 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 18 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY1 19 20 // Test target codegen - host bc file has to be created first. (no significant differences with host version of target region) 21 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc 22 // RUN: %clang_cc1 -DCK1 -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 CK1 --check-prefix CK1-64 --check-prefix TCK1 --check-prefix TCK1-64 23 // RUN: %clang_cc1 -DCK1 -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 24 // RUN: %clang_cc1 -DCK1 -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 CK1 --check-prefix CK1-64 --check-prefix TCK1 --check-prefix TCK1-64 25 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc 26 // RUN: %clang_cc1 -DCK1 -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 CK1 --check-prefix CK1-32 --check-prefix TCK1 --check-prefix TCK1-32 27 // RUN: %clang_cc1 -DCK1 -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 28 // RUN: %clang_cc1 -DCK1 -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 CK1 --check-prefix CK1-32 --check-prefix TCK1 --check-prefix TCK1-32 29 30 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc 31 // RUN: %clang_cc1 -DCK1 -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 %s --check-prefix SIMD-ONLY1 32 // RUN: %clang_cc1 -DCK1 -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 33 // RUN: %clang_cc1 -DCK1 -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 %s --check-prefix SIMD-ONLY1 34 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc 35 // RUN: %clang_cc1 -DCK1 -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 %s --check-prefix SIMD-ONLY1 36 // RUN: %clang_cc1 -DCK1 -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 37 // RUN: %clang_cc1 -DCK1 -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 %s --check-prefix SIMD-ONLY1 38 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}} 39 40 #ifdef CK1 41 42 // HCK_NO_TGT-NOT: @__kmpc_push_target_tripcount 43 44 // HCK1: define{{.*}} i32 @{{.+}}target_teams_fun{{.*}}( 45 int target_teams_fun(int *g){ 46 int n = 1000; 47 int a[1000]; 48 int te = n / 128; 49 int th = 128; 50 // discard n_addr 51 // HCK1: alloca i32, 52 // HCK1: [[TE:%.+]] alloca i32, 53 // HCK1: [[TH:%.+]] = alloca i32, 54 // discard capture expressions for te and th 55 // HCK1: = alloca i32, 56 // HCK1: = alloca i32, 57 // HCK1: = alloca i32, 58 // HCK1: = alloca i32, 59 // HCK1: = alloca i32, 60 // HCK1: [[N_CAST:%.+]] = alloca i{{32|64}}, 61 // HCK1: [[TE_CAST:%.+]] = alloca i{{32|64}}, 62 // HCK1: [[TH_CAST:%.+]] = alloca i{{32|64}}, 63 // HCK1: call void @__kmpc_push_target_tripcount(i64 -1, i64 %{{.+}}) 64 // HCK1: [[N_PAR:%.+]] = load{{.+}}, {{.+}} [[N_CAST]], 65 // HCK1: [[TE_PAR:%.+]] = load{{.+}}, {{.+}} [[TE_CAST]], 66 // HCK1: [[TH_PAR:%.+]] = load{{.+}}, {{.+}} [[TH_CAST]], 67 // HCK1: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 4, i8** %{{[^,]+}}, i8** %{{[^,]+}}, 68 69 // HCK1: call void @[[OFFL1:.+]](i{{32|64}} [[N_PAR]], {{.+}}, i{{32|64}} [[TE_PAR]], i{{32|64}} [[TH_PAR]]) 70 #pragma omp target teams distribute parallel for num_teams(te), thread_limit(th) 71 for(int i = 0; i < n; i++) { 72 a[i] = 0; 73 #pragma omp cancel for 74 } 75 76 // HCK1: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), 77 // HCK1: call void @[[OFFL2:.+]](i{{64|32}} %{{.+}}) 78 {{{ 79 #pragma omp target teams distribute parallel for is_device_ptr(g) 80 for(int i = 0; i < n; i++) { 81 a[i] = g[0]; 82 } 83 }}} 84 85 // outlined target regions 86 // HCK1: define internal void @[[OFFL1]](i{{32|64}} [[N_ARG:%.+]], i{{32|64}} [[TE_ARG:%.+]], i{{32|64}} [[TH_ARG:%.+]]) 87 // TCK1: define weak void @{{.+}}target_teams_fun{{.*}}(i{{32|64}} [[N_ARG:%.+]], {{.+}}, i{{32|64}} [[TE_ARG:%.+]], i{{32|64}} [[TH_ARG:%.+]]) 88 // CK1: [[N_ADDR:%.+]] = alloca i{{32|64}}, 89 // CK1: [[TE_ADDR:%.+]] = alloca i{{32|64}}, 90 // CK1: [[TH_ADDR:%.+]] = alloca i{{32|64}}, 91 // TCK1: store {{.+}} [[N_ARG]], {{.+}} [[N_ADDR]], 92 // CK1: store{{.+}} [[TE_ARG]], {{.+}} [[TE_ADDR]], 93 // CK1: store{{.+}} [[TH_ARG]], {{.+}} [[TH_ADDR]], 94 // CK1-64: [[TE_CONV:%.+]] = bitcast{{.+}} [[TE_ADDR]] to 95 // CK1-64: [[TH_CONV:%.+]] = bitcast{{.+}} [[TH_ADDR]] to 96 // CK1-64: [[TE_VAL:%.+]] = load i32, i32* [[TE_CONV]], 97 // CK1-64: [[TH_VAL:%.+]] = load i32, i32* [[TH_CONV]], 98 // CK1-32: [[TE_VAL:%.+]] = load i32, i32* [[TE_ADDR]], 99 // CK1-32: [[TH_VAL:%.+]] = load i32, i32* [[TH_ADDR]], 100 // CK1: {{%.+}} = call i32 @__kmpc_push_num_teams({{.+}}, {{.+}}, i32 [[TE_VAL]], i32 [[TH_VAL]]) 101 // CK1: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 2, {{.+}} @[[OUTL1:.+]] to {{.+}}, {{.+}}, {{.+}}) 102 // CK1: ret void 103 104 // CK1: define internal void @[[OUTL1]]({{.+}}) 105 // CK1: call void @__kmpc_for_static_init_4( 106 // CK1: call void {{.+}} @__kmpc_fork_call( 107 // CK1: call void @__kmpc_for_static_fini( 108 // CK1: ret void 109 110 // HCK1: define internal void @[[OFFL2]]( 111 // TCK1: define weak void @{{.+}}target_teams_fun{{.+}}( 112 // CK1: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 3, {{.+}} @[[OUTL2:.+]] to {{.+}}, {{.+}}, {{.+}}) 113 // CK1: ret void 114 115 // CK1: define internal void @[[OUTL2]]({{.+}}) 116 // CK1: call void @__kmpc_for_static_init_4( 117 // CK1: call void {{.+}} @__kmpc_fork_call( 118 // CK1: call void @__kmpc_for_static_fini( 119 // CK1: ret void 120 121 return a[0]; 122 } 123 124 #endif // CK1 125 #endif // HEADER 126