1 ; RUN: opt -S -passes=openmp-opt -openmp-ir-builder-optimistic-attributes -pass-remarks=openmp-opt -openmp-print-gpu-kernels < %s | FileCheck %s
2 ; RUN: opt -S -passes=openmp-opt -pass-remarks=openmp-opt -openmp-print-gpu-kernels < %s | FileCheck %s
4 ; C input used for this test:
11 ; #pragma omp target teams
13 ; #pragma omp parallel
17 ; #pragma omp parallel
22 ; Verify we replace the function pointer uses for the first and last outlined
23 ; region (1 and 3) but not for the middle one (2) because it could be called from
26 ; CHECK-DAG: @__omp_outlined__1_wrapper.ID = private constant i8 undef
27 ; CHECK-DAG: @__omp_outlined__2_wrapper.ID = private constant i8 undef
29 ; CHECK-DAG: icmp eq ptr %worker.work_fn.addr_cast, @__omp_outlined__1_wrapper.ID
30 ; CHECK-DAG: icmp eq ptr %worker.work_fn.addr_cast, @__omp_outlined__2_wrapper.ID
33 ; CHECK-DAG: call void @__kmpc_parallel_51(ptr @1, i32 %{{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_outlined__1, ptr @__omp_outlined__1_wrapper.ID, ptr %{{.*}}, i64 0)
34 ; CHECK-DAG: call void @__kmpc_parallel_51(ptr @1, i32 %{{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_outlined__2, ptr @__omp_outlined__2_wrapper.ID, ptr %{{.*}}, i64 0)
35 ; CHECK-DAG: call void @__kmpc_parallel_51(ptr @2, i32 %{{.*}}, i32 1, i32 -1, i32 -1, ptr @__omp_outlined__3, ptr @__omp_outlined__3_wrapper, ptr %{{.*}}, i64 0)
38 %struct.ident_t = type { i32, i32, i32, i32, ptr }
39 %struct.KernelEnvironmentTy = type { %struct.ConfigurationEnvironmentTy, ptr, ptr }
40 %struct.ConfigurationEnvironmentTy = type { i8, i8, i8, i32, i32, i32, i32 }
42 @0 = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1
43 @1 = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 0, i32 0, ptr @0 }, align 8
44 @2 = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 2, i32 0, ptr @0 }, align 8
45 @__omp_offloading_10301_87b2c_foo_l7_kernel_environment = local_unnamed_addr constant %struct.KernelEnvironmentTy { %struct.ConfigurationEnvironmentTy { i8 1, i8 0, i8 1, i32 0, i32 0, i32 0, i32 0 }, ptr @1, ptr null }
47 define weak void @__omp_offloading_10301_87b2c_foo_l7() "kernel" {
49 %.zero.addr = alloca i32, align 4
50 %.threadid_temp. = alloca i32, align 4
51 store i32 0, ptr %.zero.addr, align 4
52 %0 = call i32 @__kmpc_target_init(ptr @__omp_offloading_10301_87b2c_foo_l7_kernel_environment, ptr null)
53 %exec_user_code = icmp eq i32 %0, -1
54 br i1 %exec_user_code, label %user_code.entry, label %worker.exit
56 user_code.entry: ; preds = %entry
57 %1 = call i32 @__kmpc_global_thread_num(ptr @1)
58 store i32 %1, ptr %.threadid_temp., align 4
59 call void @__omp_outlined__(ptr %.threadid_temp., ptr %.zero.addr)
60 call void @__kmpc_target_deinit()
63 worker.exit: ; preds = %entry
67 define weak i32 @__kmpc_target_init(ptr %0, ptr) {
71 declare void @unknown()
73 define internal void @__omp_outlined__(ptr noalias %.global_tid., ptr noalias %.bound_tid.) {
75 %.global_tid..addr = alloca ptr, align 8
76 %.bound_tid..addr = alloca ptr, align 8
77 %captured_vars_addrs = alloca [0 x ptr], align 8
78 %captured_vars_addrs1 = alloca [0 x ptr], align 8
79 store ptr %.global_tid., ptr %.global_tid..addr, align 8
80 store ptr %.bound_tid., ptr %.bound_tid..addr, align 8
81 %0 = load ptr, ptr %.global_tid..addr, align 8
82 %1 = load i32, ptr %0, align 4
83 call void @__kmpc_parallel_51(ptr @1, i32 %1, i32 1, i32 -1, i32 -1, ptr @__omp_outlined__1, ptr @__omp_outlined__1_wrapper, ptr %captured_vars_addrs, i64 0)
86 call void @__kmpc_parallel_51(ptr @1, i32 %1, i32 1, i32 -1, i32 -1, ptr @__omp_outlined__2, ptr @__omp_outlined__2_wrapper, ptr %captured_vars_addrs1, i64 0)
90 define internal void @__omp_outlined__1(ptr noalias %.global_tid., ptr noalias %.bound_tid.) {
92 %.global_tid..addr = alloca ptr, align 8
93 %.bound_tid..addr = alloca ptr, align 8
94 store ptr %.global_tid., ptr %.global_tid..addr, align 8
95 store ptr %.bound_tid., ptr %.bound_tid..addr, align 8
99 define internal void @__omp_outlined__1_wrapper(i16 zeroext %0, i32 %1) {
101 %.addr = alloca i16, align 2
102 %.addr1 = alloca i32, align 4
103 %.zero.addr = alloca i32, align 4
104 %global_args = alloca ptr, align 8
105 store i32 0, ptr %.zero.addr, align 4
106 store i16 %0, ptr %.addr, align 2
107 store i32 %1, ptr %.addr1, align 4
108 call void @__kmpc_get_shared_variables(ptr %global_args)
109 call void @__omp_outlined__1(ptr %.addr1, ptr %.zero.addr)
113 declare void @__kmpc_get_shared_variables(ptr)
115 declare void @__kmpc_parallel_51(ptr, i32, i32, i32, i32, ptr, ptr, ptr, i64)
117 define hidden void @bar() {
119 %captured_vars_addrs = alloca [0 x ptr], align 8
120 %0 = call i32 @__kmpc_global_thread_num(ptr @2)
121 call void @__kmpc_parallel_51(ptr @2, i32 %0, i32 1, i32 -1, i32 -1, ptr @__omp_outlined__3, ptr @__omp_outlined__3_wrapper, ptr %captured_vars_addrs, i64 0)
125 define internal void @__omp_outlined__2(ptr noalias %.global_tid., ptr noalias %.bound_tid.) {
127 %.global_tid..addr = alloca ptr, align 8
128 %.bound_tid..addr = alloca ptr, align 8
129 store ptr %.global_tid., ptr %.global_tid..addr, align 8
130 store ptr %.bound_tid., ptr %.bound_tid..addr, align 8
134 define internal void @__omp_outlined__2_wrapper(i16 zeroext %0, i32 %1) {
136 %.addr = alloca i16, align 2
137 %.addr1 = alloca i32, align 4
138 %.zero.addr = alloca i32, align 4
139 %global_args = alloca ptr, align 8
140 store i32 0, ptr %.zero.addr, align 4
141 store i16 %0, ptr %.addr, align 2
142 store i32 %1, ptr %.addr1, align 4
143 call void @__kmpc_get_shared_variables(ptr %global_args)
144 call void @__omp_outlined__2(ptr %.addr1, ptr %.zero.addr)
148 declare i32 @__kmpc_global_thread_num(ptr)
150 declare void @__kmpc_target_deinit()
152 define internal void @__omp_outlined__3(ptr noalias %.global_tid., ptr noalias %.bound_tid.) {
154 %.global_tid..addr = alloca ptr, align 8
155 %.bound_tid..addr = alloca ptr, align 8
156 store ptr %.global_tid., ptr %.global_tid..addr, align 8
157 store ptr %.bound_tid., ptr %.bound_tid..addr, align 8
161 define internal void @__omp_outlined__3_wrapper(i16 zeroext %0, i32 %1) {
163 %.addr = alloca i16, align 2
164 %.addr1 = alloca i32, align 4
165 %.zero.addr = alloca i32, align 4
166 %global_args = alloca ptr, align 8
167 store i32 0, ptr %.zero.addr, align 4
168 store i16 %0, ptr %.addr, align 2
169 store i32 %1, ptr %.addr1, align 4
170 call void @__kmpc_get_shared_variables(ptr %global_args)
171 call void @__omp_outlined__3(ptr %.addr1, ptr %.zero.addr)
175 !omp_offload.info = !{!0}
176 !nvvm.annotations = !{!1}
177 !llvm.module.flags = !{!2, !3}
179 !0 = !{i32 0, i32 66305, i32 555956, !"foo", i32 7, i32 0}
180 !1 = !{ptr @__omp_offloading_10301_87b2c_foo_l7, !"kernel", i32 1}
181 !2 = !{i32 7, !"openmp", i32 50}
182 !3 = !{i32 7, !"openmp-device", i32 50}