| 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 -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 |
| 6 | // 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 |
| 7 | // 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 |
| 8 | // 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 |
| 9 | // 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 |
| 10 | // 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 |
| 11 | |
| 12 | // 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 |
| 13 | // 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 |
| 14 | // 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 |
| 15 | // 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 |
| 16 | // 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 |
| 17 | // 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 |
| 18 | |
| 19 | // Test target codegen - host bc file has to be created first. (no significant differences with host version of target region) |
| 20 | // 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 |
| 21 | // 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 |
| 22 | // 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 |
| 23 | // 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 |
| 24 | // 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 |
| 25 | // 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 |
| 26 | // 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 |
| 27 | // 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 |
| 28 | |
| 29 | // 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 |
| 30 | // 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 |
| 31 | // 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 |
| 32 | // 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 |
| 33 | // 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 |
| 34 | // 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 |
| 35 | // 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 |
| 36 | // 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 |
| 37 | // SIMD-ONLY1-NOT: {{__kmpc|__tgt}} |
| 38 | |
| 39 | #ifdef CK1 |
| 40 | |
| 41 | // HCK1: define{{.*}} i32 @{{.+}}target_teams_fun{{.*}}( |
| 42 | int target_teams_fun(int *g){ |
| 43 | int n = 1000; |
| 44 | int a[1000]; |
| 45 | int te = n / 128; |
| 46 | int th = 128; |
| 47 | // discard n_addr |
| 48 | // HCK1: alloca i32, |
| 49 | // HCK1: [[TE:%.+]] alloca i32, |
| 50 | // HCK1: [[TH:%.+]] = alloca i32, |
| 51 | // HCK1: [[I:%.+]] = alloca i32, |
| 52 | // discard capture expressions for te and th |
| 53 | // HCK1: = alloca i32, |
| 54 | // HCK1: = alloca i32, |
| 55 | // HCK1: = alloca i32, |
| 56 | // HCK1: = alloca i32, |
| 57 | // HCK1: = alloca i32, |
| 58 | // HCK1: [[I_CAST:%.+]] = alloca i{{32|64}}, |
| 59 | // HCK1: [[N_CAST:%.+]] = alloca i{{32|64}}, |
| 60 | // HCK1: [[TE_CAST:%.+]] = alloca i{{32|64}}, |
| 61 | // HCK1: [[TH_CAST:%.+]] = alloca i{{32|64}}, |
| 62 | // HCK1: call void @__kmpc_push_target_tripcount(i64 -1, i64 %{{.+}}) |
| 63 | // HCK1: [[I_PAR:%.+]] = load{{.+}}, {{.+}} [[I_CAST]], |
| 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 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, |
| 68 | |
| 69 | // HCK1: call void @[[OFFL1:.+]](i{{32|64}} [[I_PAR]], i{{32|64}} [[N_PAR]], {{.+}}, i{{32|64}} [[TE_PAR]], i{{32|64}} [[TH_PAR]]) |
| 70 | int i; |
| 71 | #pragma omp target teams distribute parallel for simd num_teams(te), thread_limit(th) aligned(a : 8) safelen(16) simdlen(4) linear(i : n) |
| 72 | for(i = 0; i < n; i++) { |
| 73 | a[i] = 0; |
| 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 simd is_device_ptr(g) simdlen(8) |
| 80 | for( |
| 81 | int i = 0; i < n; i++) { |
| 82 | a[i] = g[0]; |
| 83 | } |
| 84 | }}} |
| 85 | |
| 86 | // outlined target regions |
| 87 | // HCK1: define internal void @[[OFFL1]](i{{32|64}} [[I_ARG:%.+]], i{{32|64}} [[N_ARG:%.+]], {{.+}}, i{{32|64}} [[TE_ARG:%.+]], i{{32|64}} [[TH_ARG:%.+]]) |
| 88 | // TCK1: define weak void @{{.+}}target_teams_fun{{.*}}(i{{32|64}} [[I_ARG:%.+]], i{{32|64}} [[N_ARG:%.+]], {{.+}}, i{{32|64}} [[TE_ARG:%.+]], i{{32|64}} [[TH_ARG:%.+]]) |
| 89 | // CK1: [[I_ADDR:%.+]] = alloca i{{32|64}}, |
| 90 | // CK1: [[N_ADDR:%.+]] = alloca i{{32|64}}, |
| 91 | // CK1: [[TE_ADDR:%.+]] = alloca i{{32|64}}, |
| 92 | // CK1: [[TH_ADDR:%.+]] = alloca i{{32|64}}, |
| 93 | // TCK1: store {{.+}} [[N_ARG]], {{.+}} [[N_ADDR]], |
| 94 | // CK1: store{{.+}} [[TE_ARG]], {{.+}} [[TE_ADDR]], |
| 95 | // CK1: store{{.+}} [[TH_ARG]], {{.+}} [[TH_ADDR]], |
| 96 | // CK1-64: [[TE_CONV:%.+]] = bitcast{{.+}} [[TE_ADDR]] to |
| 97 | // CK1-64: [[TH_CONV:%.+]] = bitcast{{.+}} [[TH_ADDR]] to |
| 98 | // CK1-64: [[TE_VAL:%.+]] = load i32, i32* [[TE_CONV]], |
| 99 | // CK1-64: [[TH_VAL:%.+]] = load i32, i32* [[TH_CONV]], |
| 100 | // CK1-32: [[TE_VAL:%.+]] = load i32, i32* [[TE_ADDR]], |
| 101 | // CK1-32: [[TH_VAL:%.+]] = load i32, i32* [[TH_ADDR]], |
| 102 | // CK1: {{%.+}} = call i32 @__kmpc_push_num_teams({{.+}}, {{.+}}, i32 [[TE_VAL]], i32 [[TH_VAL]]) |
| 103 | // CK1: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 3, {{.+}} @[[OUTL1:.+]] to {{.+}}, {{.+}}, {{.+}}) |
| 104 | // CK1: ret void |
| 105 | |
| 106 | // CK1: define internal void @[[OUTL1]]({{.+}}) |
| 107 | // CK1: [[ARRDECAY:%.+]] = getelementptr inbounds [1000 x i32], [1000 x i32]* %{{.+}}, i{{32|64}} 0, i{{32|64}} 0 |
| 108 | // CK1: [[ARR_CAST:%.+]] = ptrtoint i32* [[ARRDECAY]] to i{{32|64}} |
| 109 | // CK1: [[MASKED_PTR:%.+]] = and i{{32|64}} [[ARR_CAST]], 7 |
| 110 | // CK1: [[COND:%.+]] = icmp eq i{{32|64}} [[MASKED_PTR]], 0 |
| 111 | // CK1: call void @llvm.assume(i1 [[COND]]) |
| 112 | // CK1: call void @__kmpc_for_static_init_4( |
| 113 | // CK1: call void {{.+}} @__kmpc_fork_call( |
| 114 | // CK1: call void @__kmpc_for_static_fini( |
| 115 | // CK1: ret void |
| 116 | |
| 117 | // HCK1: define internal void @[[OFFL2]]( |
| 118 | // TCK1: define weak void @{{.+}}target_teams_fun{{.+}}( |
| 119 | // CK1: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 3, {{.+}} @[[OUTL2:.+]] to {{.+}}, {{.+}}, {{.+}}) |
| 120 | // CK1: ret void |
| 121 | |
| 122 | // CK1: define internal void @[[OUTL2]]({{.+}}) |
| 123 | // CK1: call void @__kmpc_for_static_init_4( |
| 124 | // CK1: call void {{.+}} @__kmpc_fork_call( |
| 125 | // CK1: call void @__kmpc_for_static_fini( |
| 126 | // CK1: ret void |
| 127 | |
| 128 | return a[0]; |
| 129 | } |
| 130 | |
| 131 | // CK1-DAG: !{!"llvm.loop.vectorize.width", i32 4} |
| 132 | // CK1-DAG: !{!"llvm.loop.vectorize.enable", i1 true} |
| 133 | // CK1-DAG: !{!"llvm.loop.vectorize.width", i32 8} |
| 134 | |
| 135 | #endif // CK1 |
| 136 | #endif // HEADER |
| 137 | |