2 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP45
3 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
4 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP45
5 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP45
6 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
7 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP45
8 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -DOMP5 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP51
9 // RUN: %clang_cc1 -fopenmp -DOMP5 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
10 // RUN: %clang_cc1 -fopenmp -DOMP5 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP51
11 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -DOMP5 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP51
12 // RUN: %clang_cc1 -fopenmp -DOMP5 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
13 // RUN: %clang_cc1 -fopenmp -DOMP5 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP51
15 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
16 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
17 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
18 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
19 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
20 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
21 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -DOMP5 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
22 // RUN: %clang_cc1 -fopenmp-simd -DOMP5 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
23 // RUN: %clang_cc1 -fopenmp-simd -DOMP5 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
24 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -DOMP5 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
25 // RUN: %clang_cc1 -fopenmp-simd -DOMP5 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
26 // RUN: %clang_cc1 -fopenmp-simd -DOMP5 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
27 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
29 // Test target codegen - host bc file has to be created first.
30 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -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 -verify -Wno-vla -fopenmp -fopenmp-version=45 -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-64 --check-prefix TOMP45
32 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 %s
33 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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-64 --check-prefix TOMP45
34 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -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 -verify -Wno-vla -fopenmp -fopenmp-version=45 -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-32 --check-prefix TOMP45
36 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 %s
37 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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-32 --check-prefix TOMP45
38 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -DOMP5 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
39 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -DOMP5 -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-64 --check-prefix TOMP51
40 // RUN: %clang_cc1 -fopenmp -DOMP5 -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 %s
41 // RUN: %clang_cc1 -fopenmp -DOMP5 -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-64 --check-prefix TOMP51
42 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -DOMP5 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
43 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -DOMP5 -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-32 --check-prefix TOMP51
44 // RUN: %clang_cc1 -fopenmp -DOMP5 -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 %s
45 // RUN: %clang_cc1 -fopenmp -DOMP5 -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-32 --check-prefix TOMP51
47 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
48 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-version=45 -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-ONLY1 %s
49 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -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 %s
50 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -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-ONLY1 %s
51 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
52 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-version=45 -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-ONLY1 %s
53 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -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 %s
54 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -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-ONLY1 %s
55 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -DOMP5 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
56 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -DOMP5 -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-ONLY1 %s
57 // RUN: %clang_cc1 -fopenmp-simd -DOMP5 -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 %s
58 // RUN: %clang_cc1 -fopenmp-simd -DOMP5 -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-ONLY1 %s
59 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -DOMP5 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
60 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -DOMP5 -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-ONLY1 %s
61 // RUN: %clang_cc1 -fopenmp-simd -DOMP5 -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 %s
62 // RUN: %clang_cc1 -fopenmp-simd -DOMP5 -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-ONLY1 %s
63 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
65 // expected-no-diagnostics
69 // CHECK-DAG: [[IDENT_T:%.+]] = type { i32, i32, i32, i32, ptr }
70 // CHECK-DAG: [[KMP_TASK_T_WITH_PRIVATES:%.+]] = type { [[KMP_TASK_T:%.+]] }
71 // CHECK-DAG: [[KMP_TASK_T]] = type { ptr, ptr, i32, %{{[^,]+}}, %{{[^,]+}} }
72 // CHECK-DAG: [[TT:%.+]] = type { i64, i8 }
73 // CHECK-DAG: [[S1:%.+]] = type { double }
74 // CHECK-DAG: [[ENTTY:%.+]] = type { ptr, ptr, i[[SZ:32|64]], i32, i32 }
76 // TCHECK: [[ENTTY:%.+]] = type { ptr, ptr, i{{32|64}}, i32, i32 }
78 // We have 8 target regions, but only 7 that actually will generate offloading
79 // code, only 6 will have mapped arguments, and only 4 have all-constant map
82 // CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [3 x i64] [i64 2, i64 4, i64 4]
83 // CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 800]
84 // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 2]
85 // CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800]
86 // CHECK-DAG: [[SIZET4:@.+]] = private unnamed_addr constant [9 x i64] [i64 4, i64 40, i64 {{8|4}}, i64 0, i64 400, i64 {{8|4}}, i64 {{8|4}}, i64 0, i64 {{16|12}}]
87 // CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i64] [i64 800, i64 547, i64 800, i64 547, i64 547, i64 800, i64 800, i64 547, i64 547]
88 // CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i64] [i64 4, i64 2, i64 40]
89 // CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 547]
90 // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i64] [i64 4, i64 2, i64 1, i64 40]
91 // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 800, i64 800, i64 800, i64 547]
92 // OMP45-DAG: [[SIZET7:@.+]] = private unnamed_addr constant [5 x i64] [i64 {{8|4}}, i64 4, i64 {{8|4}}, i64 {{8|4}}, i64 0]
93 // OMP51-DAG: [[SIZET7:@.+]] = private unnamed_addr constant [6 x i64] [i64 {{8|4}}, i64 4, i64 {{8|4}}, i64 {{8|4}}, i64 0, i64 1]
94 // OMP45-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547]
95 // OMP51-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [6 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547, i64 800]
96 // CHECK-DAG: @{{.*}} = weak constant i8 0
97 // CHECK-DAG: @{{.*}} = weak constant i8 0
98 // CHECK-DAG: @{{.*}} = weak constant i8 0
99 // CHECK-DAG: @{{.*}} = weak constant i8 0
100 // CHECK-DAG: @{{.*}} = weak constant i8 0
101 // CHECK-DAG: @{{.*}} = weak constant i8 0
102 // CHECK-DAG: @{{.*}} = weak constant i8 0
104 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
105 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
106 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
107 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
108 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
109 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
110 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
111 // TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
113 // Check target registration is registered as a Ctor.
114 // CHECK: appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 0, ptr @.omp_offloading.requires_reg, ptr null }]
117 template<typename tx
, typename ty
>
123 // CHECK-LABEL: get_val
124 long long get_val() { return 0; }
126 // CHECK: define {{.*}}[[FOO:@.+]](
134 TT
<long long, char> d
;
136 // CHECK-32: [[TASK:%.+]] = call ptr @__kmpc_omp_target_task_alloc(ptr @{{[^,]+}}, i32 %{{[^,]+}}, i32 1, i32 20, i32 1, ptr [[OMP_TASK_ENTRY:@[^,]+]], i64 -1)
137 // CHECK-64: [[TASK:%.+]] = call ptr @__kmpc_omp_target_task_alloc(ptr @{{[^,]+}}, i32 %{{[^,]+}}, i32 1, i64 40, i64 1, ptr [[OMP_TASK_ENTRY:@[^,]+]], i64 -1)
138 // CHECK: call i32 @__kmpc_omp_task(ptr @{{[^,]+}}, i32 %{{[^,]+}}, ptr [[TASK]])
139 #pragma omp target simd nowait
140 for (int i
= 3; i
< 32; i
+= 5) {
143 // CHECK: call void [[HVT1:@.+]](i[[SZ]] {{[^,]+}}, {{[^)]+}})
144 long long k
= get_val();
145 #pragma omp target simd if(target: 0) linear(k : 3)
146 for (int i
= 10; i
> 1; i
--) {
150 // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
151 // CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
152 // CHECK-DAG: store ptr [[BP:%.+]], ptr [[BPARG]]
153 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
154 // CHECK-DAG: store ptr [[P:%.+]], ptr [[PARG]]
155 // CHECK-DAG: [[BP]] = getelementptr inbounds [3 x ptr], ptr [[BPR:%[^,]+]], i32 0, i32 0
156 // CHECK-DAG: [[P]] = getelementptr inbounds [3 x ptr], ptr [[PR:%[^,]+]], i32 0, i32 0
157 // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BPR]], i32 0, i32 0
158 // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [3 x ptr], ptr [[PR]], i32 0, i32 0
159 // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], ptr [[BPADDR0]],
160 // CHECK-DAG: store i[[SZ]] [[VAL0]], ptr [[PADDR0]],
161 // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BPR]], i32 0, i32 1
162 // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[PR]], i32 0, i32 1
163 // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], ptr [[BPADDR1]],
164 // CHECK-DAG: store i[[SZ]] [[VAL1]], ptr [[PADDR1]],
165 // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BPR]], i32 0, i32 2
166 // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[PR]], i32 0, i32 2
167 // CHECK-DAG: store i[[SZ]] [[VAL2:%.+]], ptr [[BPADDR2]],
168 // CHECK-DAG: store i[[SZ]] [[VAL2]], ptr [[PADDR2]],
170 // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
171 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
173 // CHECK: call void [[HVT2:@.+]](i[[SZ]] {{[^,]+}}, i[[SZ]] {{[^)]+}})
174 // CHECK-NEXT: br label %[[END]]
177 #pragma omp target simd if(target: 1) linear(lin, a : get_val())
178 for (unsigned long long it
= 2000; it
>= 600; it
-=400) {
182 // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 10
183 // CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
185 // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
186 // CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
187 // CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]
188 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
189 // CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]
190 // CHECK-DAG: [[BPR]] = getelementptr inbounds [2 x ptr], ptr [[BP:%[^,]+]], i32 0, i32 0
191 // CHECK-DAG: [[PR]] = getelementptr inbounds [2 x ptr], ptr [[P:%[^,]+]], i32 0, i32 0
193 // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP]], i32 0, i32 0
194 // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 0
195 // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], ptr [[BPADDR0]],
196 // CHECK-DAG: store i[[SZ]] [[VAL0]], ptr [[PADDR0]],
198 // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP]], i32 0, i32 1
199 // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 1
200 // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], ptr [[BPADDR1]],
201 // CHECK-DAG: store i[[SZ]] [[VAL1]], ptr [[PADDR1]],
202 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
203 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
205 // CHECK: call void [[HVT3:@.+]]({{[^,]+}}, {{[^,]+}})
206 // CHECK-NEXT: br label %[[END]]
208 // CHECK-NEXT: br label %[[IFEND:.+]]
210 // CHECK: call void [[HVT3]]({{[^,]+}}, {{[^,]+}})
211 // CHECK-NEXT: br label %[[IFEND]]
214 #pragma omp target simd if(target: n>10)
215 for (short it
= 6; it
<= 20; it
-=-4) {
220 // We capture 3 VLA sizes in this target region
221 // CHECK-64: [[A_VAL:%.+]] = load i32, ptr %{{.+}},
222 // CHECK-64: store i32 [[A_VAL]], ptr [[A_CADDR:%.+]],
223 // CHECK-64: [[A_CVAL:%.+]] = load i[[SZ]], ptr [[A_CADDR]],
225 // CHECK-32: [[A_VAL:%.+]] = load i32, ptr %{{.+}},
226 // CHECK-32: store i32 [[A_VAL]], ptr [[A_CADDR:%.+]],
227 // CHECK-32: [[A_CVAL:%.+]] = load i[[SZ]], ptr [[A_CADDR]],
229 // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 20
230 // CHECK: br i1 [[IF]], label %[[TRY:[^,]+]], label %[[FAIL:[^,]+]]
232 // CHECK-64: [[BNSIZE:%.+]] = mul nuw i64 [[VLA0:%.+]], 4
233 // CHECK-32: [[BNSZSIZE:%.+]] = mul nuw i32 [[VLA0:%.+]], 4
234 // CHECK-32: [[BNSIZE:%.+]] = sext i32 [[BNSZSIZE]] to i64
235 // CHECK: [[CNELEMSIZE2:%.+]] = mul nuw i[[SZ]] 5, [[VLA1:%.+]]
236 // CHECK-64: [[CNSIZE:%.+]] = mul nuw i64 [[CNELEMSIZE2]], 8
237 // CHECK-32: [[CNSZSIZE:%.+]] = mul nuw i32 [[CNELEMSIZE2]], 8
238 // CHECK-32: [[CNSIZE:%.+]] = sext i32 [[CNSZSIZE]] to i64
240 // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
241 // CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
242 // CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]
243 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
244 // CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]
245 // CHECK-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
246 // CHECK-DAG: store ptr [[SR:%.+]], ptr [[SARG]]
247 // CHECK-DAG: [[BPR]] = getelementptr inbounds [9 x ptr], ptr [[BP:%[^,]+]], i32 0, i32 0
248 // CHECK-DAG: [[PR]] = getelementptr inbounds [9 x ptr], ptr [[P:%[^,]+]], i32 0, i32 0
249 // CHECK-DAG: [[SR]] = getelementptr inbounds [9 x i64], ptr [[S:%[^,]+]], i32 0, i32 0
251 // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX0:[0-9]+]]
252 // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX0]]
253 // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX1:[0-9]+]]
254 // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX1]]
255 // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX2:[0-9]+]]
256 // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX2]]
257 // CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [9 x i64], ptr [[S]], i32 0, i32 [[IDX3:[0-9]+]]
258 // CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX3]]
259 // CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX3]]
260 // CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX4:[0-9]+]]
261 // CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX4]]
262 // CHECK-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX5:[0-9]+]]
263 // CHECK-DAG: [[PADDR5:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX5]]
264 // CHECK-DAG: [[BPADDR6:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX6:[0-9]+]]
265 // CHECK-DAG: [[PADDR6:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX6]]
266 // CHECK-DAG: [[SADDR7:%.+]] = getelementptr inbounds [9 x i64], ptr [[S]], i32 0, i32 [[IDX7:[0-9]+]]
267 // CHECK-DAG: [[BPADDR7:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX7]]
268 // CHECK-DAG: [[PADDR7:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX7]]
269 // CHECK-DAG: [[BPADDR8:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX8:[0-9]+]]
270 // CHECK-DAG: [[PADDR8:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX8]]
272 // The names below are not necessarily consistent with the names used for the
273 // addresses above as some are repeated.
274 // CHECK-DAG: store i[[SZ]] [[A_CVAL]], ptr {{%[^,]+}},
275 // CHECK-DAG: store i[[SZ]] [[A_CVAL]], ptr {{%[^,]+}},
277 // CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},
278 // CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},
280 // CHECK-DAG: store i[[SZ]] %{{.+}}, ptr {{%[^,]+}},
281 // CHECK-DAG: store i[[SZ]] %{{.+}}, ptr {{%[^,]+}},
283 // CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},
284 // CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},
285 // CHECK-DAG: store i64 [[BNSIZE]], ptr {{%[^,]+}}
287 // CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},
288 // CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},
290 // CHECK-DAG: store i[[SZ]] 5, ptr {{%[^,]+}},
291 // CHECK-DAG: store i[[SZ]] 5, ptr {{%[^,]+}},
293 // CHECK-DAG: store i[[SZ]] %{{.+}}, ptr {{%[^,]+}},
294 // CHECK-DAG: store i[[SZ]] %{{.+}}, ptr {{%[^,]+}},
296 // CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},
297 // CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},
298 // CHECK-DAG: store i64 [[CNSIZE]], ptr {{%[^,]+}}
300 // CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},
301 // CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},
303 // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
304 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
307 // CHECK: call void [[HVT4:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
308 // CHECK-NEXT: br label %[[END]]
310 #pragma omp target simd if (target \
312 for (unsigned char it
= 'z'; it
>= 'a'; it
+=-1) {
325 // Check that the offloading functions are emitted and that the arguments are
326 // correct and loaded correctly for the target regions in foo().
328 // CHECK: define internal void [[HVT0:@.+]]()
333 // CHECK: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%0, ptr noalias noundef %1)
334 // CHECK: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
335 // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
336 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
338 // CHECK: call void [[HVT0]]()
339 // CHECK-NEXT: br label %[[END]]
342 // CHECK: define internal void [[HVT1]](i[[SZ]] noundef %{{.+}}, {{.+}})
343 // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align
344 // CHECK: store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align
345 // CHECK-64: [[AA:%.+]] = load i32, ptr [[AA_ADDR]], align
346 // CHECK-32: [[AA:%.+]] = load i32, ptr [[AA_ADDR]], align
347 // CHECK: !llvm.access.group
352 // CHECK: define internal void [[HVT2]](i[[SZ]] noundef %{{.+}}, i[[SZ]] noundef %{{.+}}, i[[SZ]] noundef %{{.+}})
353 // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align
354 // CHECK: store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align
355 // CHECK: [[AA:%.+]] = load i16, ptr [[AA_ADDR]], align
360 // CHECK: define internal void [[HVT3]]
361 // CHECK: [[A_ADDR:%.+]] = alloca i[[SZ]], align
362 // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align
363 // CHECK-DAG: store i[[SZ]] %{{.+}}, ptr [[A_ADDR]], align
364 // CHECK-DAG: store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align
369 // CHECK: define internal void [[HVT4]]
370 // Create local storage for each capture.
371 // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
372 // CHECK: [[LOCAL_B:%.+]] = alloca ptr
373 // CHECK: [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]
374 // CHECK: [[LOCAL_BN:%.+]] = alloca ptr
375 // CHECK: [[LOCAL_C:%.+]] = alloca ptr
376 // CHECK: [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]
377 // CHECK: [[LOCAL_VLA3:%.+]] = alloca i[[SZ]]
378 // CHECK: [[LOCAL_CN:%.+]] = alloca ptr
379 // CHECK: [[LOCAL_D:%.+]] = alloca ptr
380 // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], ptr [[LOCAL_A]]
381 // CHECK-DAG: store ptr [[ARG_B:%.+]], ptr [[LOCAL_B]]
382 // CHECK-DAG: store i[[SZ]] [[ARG_VLA1:%.+]], ptr [[LOCAL_VLA1]]
383 // CHECK-DAG: store ptr [[ARG_BN:%.+]], ptr [[LOCAL_BN]]
384 // CHECK-DAG: store ptr [[ARG_C:%.+]], ptr [[LOCAL_C]]
385 // CHECK-DAG: store i[[SZ]] [[ARG_VLA2:%.+]], ptr [[LOCAL_VLA2]]
386 // CHECK-DAG: store i[[SZ]] [[ARG_VLA3:%.+]], ptr [[LOCAL_VLA3]]
387 // CHECK-DAG: store ptr [[ARG_CN:%.+]], ptr [[LOCAL_CN]]
388 // CHECK-DAG: store ptr [[ARG_D:%.+]], ptr [[LOCAL_D]]
390 // CHECK-DAG: [[REF_B:%.+]] = load ptr, ptr [[LOCAL_B]],
391 // CHECK-DAG: [[VAL_VLA1:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA1]],
392 // CHECK-DAG: [[REF_BN:%.+]] = load ptr, ptr [[LOCAL_BN]],
393 // CHECK-DAG: [[REF_C:%.+]] = load ptr, ptr [[LOCAL_C]],
394 // CHECK-DAG: [[VAL_VLA2:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA2]],
395 // CHECK-DAG: [[VAL_VLA3:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA3]],
396 // CHECK-DAG: [[REF_CN:%.+]] = load ptr, ptr [[LOCAL_CN]],
397 // CHECK-DAG: [[REF_D:%.+]] = load ptr, ptr [[LOCAL_D]],
400 template<typename tx
>
401 tx
ftemplate(int n
) {
406 #pragma omp target simd if(target: n>40)
407 for (long long i
= -10; i
< 10; i
+= 3) {
423 #pragma omp target simd if(target: n>50)
424 for (unsigned i
=100; i
<10; i
+=10) {
442 #pragma omp target simd if(n>60) nontemporal(a) private(a)
444 #pragma omp target simd if(n>60) private(a)
446 for (unsigned long long it
= 2000; it
>= 600; it
-= 400) {
447 this->a
= (double)b
+ 1.5;
451 return c
[1][1] + (int)b
;
455 // CHECK: define {{.*}}@{{.*}}bar{{.*}}
459 // CHECK: call {{.*}}i32 [[FOO]](i32 {{.*}})
463 // CHECK: call {{.*}}i32 [[FS1:@.+]](ptr {{.*}}, i32 {{.*}})
466 // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}})
469 // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}})
470 a
+= ftemplate
<int>(n
);
476 // CHECK: define {{.*}}[[FS1]]
478 // CHECK: ptr @llvm.stacksave.p0()
479 // CHECK-32: store i32 %{{.+}}, ptr %__vla_expr
480 // OMP51: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 60
481 // CHECK-64: store i32 %{{.+}}, ptr [[B_ADDR:%.+]],
482 // CHECK-64: [[B_CVAL:%.+]] = load i[[SZ]], ptr [[B_ADDR]],
484 // CHECK-32: store i32 %{{.+}}, ptr [[B_ADDR:%.+]],
485 // CHECK-32: [[B_CVAL:%.+]] = load i[[SZ]], ptr [[B_ADDR]],
487 // OMP45: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 60
488 // OMP51: [[TOBOOL:%.+]] = trunc i8 %{{.+}} to i1
489 // OMP51: [[FROMBOOL:%.+]] = zext i1 [[TOBOOL]] to i8
490 // OMP51: store i8 [[FROMBOOL]], ptr [[CAP:%.+]],
491 // OMP51: [[SIMD_COND:%.+]] = load i[[SZ]], ptr [[CAP]],
492 // OMP51: [[IF:%.+]] = trunc i8 %{{.+}} to i1
493 // CHECK: br i1 [[IF]], label %[[TRY:[^,]+]], label %[[FAIL:[^,]+]]
495 // We capture 2 VLA sizes in this target region
496 // CHECK: [[CELEMSIZE2:%.+]] = mul nuw i[[SZ]] 2, [[VLA0:%.+]]
497 // CHECK-64: [[CSIZE:%.+]] = mul nuw i64 [[CELEMSIZE2]], 2
498 // CHECK-32: [[CSZSIZE:%.+]] = mul nuw i32 [[CELEMSIZE2]], 2
499 // CHECK-32: [[CSIZE:%.+]] = sext i32 [[CSZSIZE]] to i64
501 // OMP45-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
502 // OMP45-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
503 // OMP45-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]
504 // OMP45-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
505 // OMP45-DAG: store ptr [[PR:%.+]], ptr [[PARG]]
506 // OMP45-DAG: [[BPR]] = getelementptr inbounds [5 x ptr], ptr [[BP:%.+]], i32 0, i32 0
507 // OMP45-DAG: [[PR]] = getelementptr inbounds [5 x ptr], ptr [[P:%.+]], i32 0, i32 0
508 // OMP45-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 [[IDX0:[0-9]+]]
509 // OMP45-DAG: [[PADDR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 [[IDX0]]
510 // OMP45-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 [[IDX1:[0-9]+]]
511 // OMP45-DAG: [[PADDR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 [[IDX1]]
512 // OMP45-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 [[IDX2:[0-9]+]]
513 // OMP45-DAG: [[PADDR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 [[IDX2]]
514 // OMP45-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 [[IDX3:[0-9]+]]
515 // OMP45-DAG: [[PADDR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 [[IDX3]]
516 // OMP45-DAG: [[SADDR4:%.+]] = getelementptr inbounds [5 x i64], ptr [[S]], i32 [[IDX4:[0-9]+]]
517 // OMP45-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 [[IDX4]]
518 // OMP45-DAG: [[PADDR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 [[IDX4]]
519 // OMP51-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
520 // OMP51-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
521 // OMP51-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]
522 // OMP51-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
523 // OMP51-DAG: store ptr [[PR:%.+]], ptr [[PARG]]
524 // OMP51-DAG: [[BPR]] = getelementptr inbounds [6 x ptr], ptr [[BP:%.+]], i32 0, i32 0
525 // OMP51-DAG: [[PR]] = getelementptr inbounds [6 x ptr], ptr [[P:%.+]], i32 0, i32 0
526 // OMP51-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BP]], i32 [[IDX0:[0-9]+]]
527 // OMP51-DAG: [[PADDR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[P]], i32 [[IDX0]]
528 // OMP51-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BP]], i32 [[IDX1:[0-9]+]]
529 // OMP51-DAG: [[PADDR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[P]], i32 [[IDX1]]
530 // OMP51-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BP]], i32 [[IDX2:[0-9]+]]
531 // OMP51-DAG: [[PADDR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[P]], i32 [[IDX2]]
532 // OMP51-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BP]], i32 [[IDX3:[0-9]+]]
533 // OMP51-DAG: [[PADDR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[P]], i32 [[IDX3]]
534 // OMP51-DAG: [[SADDR4:%.+]] = getelementptr inbounds [6 x i64], ptr [[S]], i32 [[IDX4:[0-9]+]]
535 // OMP51-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BP]], i32 [[IDX4]]
536 // OMP51-DAG: [[PADDR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[P]], i32 [[IDX4]]
537 // OMP51-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BP]], i32 [[IDX5:[0-9]+]]
538 // OMP51-DAG: [[PADDR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[P]], i32 [[IDX5]]
540 // The names below are not necessarily consistent with the names used for the
541 // addresses above as some are repeated.
542 // CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},
543 // CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},
545 // CHECK-DAG: store i[[SZ]] [[B_CVAL]], ptr {{%[^,]+}},
546 // CHECK-DAG: store i[[SZ]] [[B_CVAL]], ptr {{%[^,]+}},
548 // CHECK-DAG: store i[[SZ]] 2, ptr {{%[^,]+}},
549 // CHECK-DAG: store i[[SZ]] 2, ptr {{%[^,]+}},
551 // CHECK-DAG: store i[[SZ]] %{{.+}}, ptr {{%[^,]+}},
552 // CHECK-DAG: store i[[SZ]] %{{.+}}, ptr {{%[^,]+}},
554 // CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},
555 // CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},
556 // CHECK-DAG: store i64 [[CSIZE]], ptr {{%[^,]+}}
558 // OMP51-DAG: store i[[SZ]] [[SIMD_COND]], ptr {{%[^,]+}}
559 // OMP51-DAG: store i[[SZ]] [[SIMD_COND]], ptr {{%[^,]+}}
561 // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
562 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
565 // OMP45: call void [[HVT7:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
566 // OMP51: call void [[HVT7:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
567 // CHECK-NEXT: br label %[[END]]
571 // CHECK: define {{.*}}[[FSTATIC]]
573 // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 50
574 // CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
576 // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
577 // CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
578 // CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]
579 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
580 // CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]
581 // CHECK-DAG: [[BPR]] = getelementptr inbounds [4 x ptr], ptr [[BP:%.+]], i32 0, i32 0
582 // CHECK-DAG: [[PR]] = getelementptr inbounds [4 x ptr], ptr [[P:%.+]], i32 0, i32 0
584 // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 0
585 // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 0
586 // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], ptr [[BPADDR0]],
587 // CHECK-DAG: store i[[SZ]] [[VAL0]], ptr [[PADDR0]],
589 // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 1
590 // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 1
591 // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], ptr [[BPADDR1]],
592 // CHECK-DAG: store i[[SZ]] [[VAL1]], ptr [[PADDR1]],
594 // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 2
595 // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 2
596 // CHECK-DAG: store i[[SZ]] [[VAL2:%.+]], ptr [[BPADDR2]],
597 // CHECK-DAG: store i[[SZ]] [[VAL2]], ptr [[PADDR2]],
599 // CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 3
600 // CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 3
601 // CHECK-DAG: store ptr [[VAL3:%.+]], ptr [[BPADDR3]],
602 // CHECK-DAG: store ptr [[VAL3]], ptr [[PADDR3]],
604 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
605 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
607 // CHECK: call void [[HVT6:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
608 // CHECK-NEXT: br label %[[END]]
610 // CHECK-NEXT: br label %[[IFEND:.+]]
612 // CHECK: call void [[HVT6]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
613 // CHECK-NEXT: br label %[[IFEND]]
617 // CHECK: define {{.*}}[[FTEMPLATE]]
619 // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 40
620 // CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
622 // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
623 // CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
624 // CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]
625 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
626 // CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]
627 // CHECK-DAG: [[BPR]] = getelementptr inbounds [3 x ptr], ptr [[BP:%.+]], i32 0, i32 0
628 // CHECK-DAG: [[PR]] = getelementptr inbounds [3 x ptr], ptr [[P:%.+]], i32 0, i32 0
630 // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 0
631 // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 0
632 // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], ptr [[BPADDR0]],
633 // CHECK-DAG: store i[[SZ]] [[VAL0]], ptr [[PADDR0]],
635 // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 1
636 // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 1
637 // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], ptr [[BPADDR1]],
638 // CHECK-DAG: store i[[SZ]] [[VAL1]], ptr [[PADDR1]],
640 // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 2
641 // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 2
642 // CHECK-DAG: store ptr [[VAL2:%.+]], ptr [[BPADDR2]],
643 // CHECK-DAG: store ptr [[VAL2]], ptr [[PADDR2]],
645 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
646 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
648 // CHECK: call void [[HVT5:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}})
649 // CHECK-NEXT: br label %[[END]]
651 // CHECK-NEXT: br label %[[IFEND:.+]]
653 // CHECK: call void [[HVT:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}})
654 // CHECK-NEXT: br label %[[IFEND]]
657 // Check that the offloading functions are emitted and that the arguments are
658 // correct and loaded correctly for the target regions of the callees of bar().
660 // CHECK: define internal void [[HVT7]]
661 // Create local storage for each capture.
662 // CHECK: [[LOCAL_THIS:%.+]] = alloca ptr
663 // CHECK: [[LOCAL_B:%.+]] = alloca i[[SZ]]
664 // CHECK: [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]
665 // CHECK: [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]
666 // CHECK: [[LOCAL_C:%.+]] = alloca ptr
667 // OMP51: [[LOCAL_SIMD_COND_CASTED:%.+]] = alloca i[[SZ]],
668 // CHECK-DAG: store ptr [[ARG_THIS:%.+]], ptr [[LOCAL_THIS]]
669 // CHECK-DAG: store i[[SZ]] [[ARG_B:%.+]], ptr [[LOCAL_B]]
670 // CHECK-DAG: store i[[SZ]] [[ARG_VLA1:%.+]], ptr [[LOCAL_VLA1]]
671 // CHECK-DAG: store i[[SZ]] [[ARG_VLA2:%.+]], ptr [[LOCAL_VLA2]]
672 // CHECK-DAG: store ptr [[ARG_C:%.+]], ptr [[LOCAL_C]]
673 // Store captures in the context.
674 // CHECK-DAG: [[REF_THIS:%.+]] = load ptr, ptr [[LOCAL_THIS]],
675 // CHECK-DAG: [[VAL_VLA1:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA1]],
676 // CHECK-DAG: [[VAL_VLA2:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA2]],
677 // CHECK-DAG: [[REF_C:%.+]] = load ptr, ptr [[LOCAL_C]],
678 // OMP51-DAG: [[SIMD_COND:%.+]] = load i8, ptr [[LOCAL_SIMD_COND_CASTED]],
679 // OMP51-DAG: trunc i8 [[SIMD_COND]] to i1
680 // OMP45-NOT: !nontemporal
681 // OMP51: store double {{.*}}!nontemporal
682 // OMP51: load double, {{.*}}!nontemporal
683 // OMP51: store double {{.*}}!nontemporal
685 // CHECK: define internal void [[HVT6]]
686 // Create local storage for each capture.
687 // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
688 // CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]
689 // CHECK: [[LOCAL_AAA:%.+]] = alloca i[[SZ]]
690 // CHECK: [[LOCAL_B:%.+]] = alloca ptr
691 // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], ptr [[LOCAL_A]]
692 // CHECK-DAG: store i[[SZ]] [[ARG_AA:%.+]], ptr [[LOCAL_AA]]
693 // CHECK-DAG: store i[[SZ]] [[ARG_AAA:%.+]], ptr [[LOCAL_AAA]]
694 // CHECK-DAG: store ptr [[ARG_B:%.+]], ptr [[LOCAL_B]]
695 // Store captures in the context.
696 // CHECK-DAG: [[REF_B:%.+]] = load ptr, ptr [[LOCAL_B]],
698 // CHECK: define internal void [[HVT5]]
699 // Create local storage for each capture.
700 // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
701 // CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]
702 // CHECK: [[LOCAL_B:%.+]] = alloca ptr
703 // CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], ptr [[LOCAL_A]]
704 // CHECK-DAG: store i[[SZ]] [[ARG_AA:%.+]], ptr [[LOCAL_AA]]
705 // CHECK-DAG: store ptr [[ARG_B:%.+]], ptr [[LOCAL_B]]
706 // Store captures in the context.
707 // CHECK-DAG: [[REF_B:%.+]] = load ptr, ptr [[LOCAL_B]],
709 // OMP45-NOT: !{!"llvm.loop.vectorize.enable", i1 false}
710 // TOMP45-NOT: !{!"llvm.loop.vectorize.enable", i1 false}
711 // OMP51: !{!"llvm.loop.vectorize.enable", i1 false}
712 // TOMP51: !{!"llvm.loop.vectorize.enable", i1 false}