1 // Test target codegen - host bc file has to be created first.
2 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
3 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix SEQ
4 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -disable-llvm-optzns -fopenmp-cuda-parallel-target-regions | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix PAR
5 // expected-no-diagnostics
6 #ifndef HEADER
7 #define HEADER
8 
9 template<typename tx>
10 tx ftemplate(int n) {
11   tx b[10];
12 
13   #pragma omp target
14   {
15     tx d = n;
16     #pragma omp parallel for
17     for(int i=0; i<10; i++) {
18       b[i] += d;
19     }
20     b[3] += 1;
21   }
22 
23   return b[3];
24 }
25 
26 int bar(int n){
27   int a = 0;
28 
29   a += ftemplate<int>(n);
30 
31   return a;
32 }
33 
34 // SEQ: [[MEM_TY:%.+]] = type { [128 x i8] }
35 // SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = common addrspace(3) global [[MEM_TY]] zeroinitializer
36 // SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* null
37 // SEQ-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i{{64|32}} 4
38 // SEQ-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1
39 
40 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l13}}_worker()
41 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
42 // CHECK: call i1 @__kmpc_kernel_parallel(
43 // CHECK: call void @__omp_outlined___wrapper(
44 
45 // CHECK: define weak void @__omp_offloading_{{.*}}l13(
46 // CHECK: call void @__omp_offloading_{{.*}}l13_worker()
47 // CHECK: call void @__kmpc_kernel_init(
48 // CHECK: call void @__kmpc_data_sharing_init_stack()
49 // SEQ: [[IS_SHARED:%.+]] = load i16, i16* [[KERNEL_SHARED]],
50 // SEQ: [[SIZE:%.+]] = load i{{64|32}}, i{{64|32}}* [[KERNEL_SIZE]],
51 // SEQ: call void @__kmpc_get_team_static_memory(i16 0, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([[MEM_TY]], [[MEM_TY]] addrspace(3)* [[SHARED_GLOBAL_RD]], i32 0, i32 0, i32 0) to i8*), i64 %7, i16 %6, i8** addrspacecast (i8* addrspace(3)* [[KERNEL_PTR]] to i8**))
52 // SEQ: [[KERNEL_RD:%.+]] = load i8*, i8* addrspace(3)* [[KERNEL_PTR]],
53 // SEQ: [[STACK:%.+]] = getelementptr inbounds i8, i8* [[KERNEL_RD]], i{{64|32}} 0
54 // PAR: [[STACK:%.+]] = call i8* @__kmpc_data_sharing_push_stack(i{{32|64}} 4, i16 1)
55 // CHECK: call void @__kmpc_kernel_prepare_parallel(
56 // CHECK: call void @__kmpc_begin_sharing_variables({{.*}}, i64 2)
57 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
58 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
59 // CHECK: call void @__kmpc_end_sharing_variables()
60 // SEQ: [[IS_SHARED:%.+]] = load i16, i16* [[KERNEL_SHARED]],
61 // SEQ: call void @__kmpc_restore_team_static_memory(i16 0, i16 [[IS_SHARED]])
62 // PAR: call void @__kmpc_data_sharing_pop_stack(i8* [[STACK]])
63 // CHECK: call void @__kmpc_kernel_deinit(i16 1)
64 
65 // CHECK: define internal void @__omp_outlined__(
66 // CHECK: alloca
67 // CHECK: alloca
68 // CHECK: alloca
69 // CHECK: alloca
70 // CHECK: [[OMP_IV:%.*]] = alloca i32
71 // CHECK: store i32 0, {{.*}} [[OMP_LB:%.+]],
72 // CHECK: store i32 9, {{.*}} [[OMP_UB:%.+]],
73 // CHECK: store i32 1, {{.*}} [[OMP_ST:%.+]],
74 // CHECK: call void @__kmpc_for_static_init_4({{.*}} i32 33, {{.*}} [[OMP_LB]], {{.*}} [[OMP_UB]], {{.*}} [[OMP_ST]], i32 1, i32 1)
75 // CHECK: br label %[[OMP_DISPATCH_COND:.+]]
76 
77 // CHECK: [[OMP_DISPATCH_COND]]
78 // CHECK: [[OMP_UB_1:%.+]] = load {{.*}} [[OMP_UB]]
79 // CHECK: [[COMP_1:%.+]] = icmp sgt {{.*}} [[OMP_UB_1]]
80 // CHECK: br i1 [[COMP_1]], label %[[COND_TRUE:.+]], label %[[COND_FALSE:.+]]
81 
82 // CHECK: [[COND_TRUE]]
83 // CHECK: br label %[[COND_END:.+]]
84 
85 // CHECK: [[COND_FALSE]]
86 // CHECK: [[OMP_UB_2:%.+]] = load {{.*}}* [[OMP_UB]]
87 // CHECK: br label %[[COND_END]]
88 
89 // CHECK: [[COND_END]]
90 // CHECK: [[COND_RES:%.+]] = phi i32 [ 9, %[[COND_TRUE]] ], [ [[OMP_UB_2]], %[[COND_FALSE]] ]
91 // CHECK: store i32 [[COND_RES]], i32* [[OMP_UB]]
92 // CHECK: [[OMP_LB_1:%.+]] = load i32, i32* [[OMP_LB]]
93 // CHECK: store i32 [[OMP_LB_1]], i32* [[OMP_IV]]
94 // CHECK: [[OMP_IV_1:%.+]] = load i32, i32* [[OMP_IV]]
95 // CHECK: [[OMP_UB_3:%.+]] = load i32, i32* [[OMP_UB]]
96 // CHECK: [[COMP_2:%.+]] = icmp sle i32 [[OMP_IV_1]], [[OMP_UB_3]]
97 // CHECK: br i1 [[COMP_2]], label %[[DISPATCH_BODY:.+]], label %[[DISPATCH_END:.+]]
98 
99 // CHECK: [[DISPATCH_BODY]]
100 // CHECK: br label %[[OMP_INNER_FOR_COND:.+]]
101 
102 // CHECK: [[OMP_INNER_FOR_COND]]
103 // CHECK: [[OMP_IV_2:%.+]] = load i32, i32* [[OMP_IV]]
104 // CHECK: [[OMP_UB_4:%.+]] = load i32, i32* [[OMP_UB]]
105 // CHECK: [[COMP_3:%.+]] = icmp sle i32 [[OMP_IV_2]], [[OMP_UB_4]]
106 // CHECK: br i1 [[COMP_3]], label %[[OMP_INNER_FOR_BODY:.+]], label %[[OMP_INNER_FOR_END:.+]]
107 
108 // CHECK: [[OMP_INNER_FOR_BODY]]
109 // CHECK: br label %[[OMP_BODY_CONTINUE:.+]]
110 
111 // CHECK: [[OMP_BODY_CONTINUE]]
112 // CHECK: br label %[[OMP_INNER_FOR_INC:.+]]
113 
114 // CHECK: [[OMP_INNER_FOR_INC]]
115 // CHECK: [[OMP_IV_3:%.+]] = load i32, i32* [[OMP_IV]]
116 // CHECK: [[ADD_1:%.+]] = add nsw i32 [[OMP_IV_3]], 1
117 // CHECK: store i32 [[ADD_1]], i32* [[OMP_IV]]
118 // CHECK: br label %[[OMP_INNER_FOR_COND]]
119 
120 // CHECK: [[OMP_INNER_FOR_COND]]
121 // CHECK: br label %[[OMP_DISPATCH_INC:.+]]
122 
123 // CHECK: [[OMP_DISPATCH_INC]]
124 // CHECK: [[OMP_LB_2:%.+]] = load i32, i32* [[OMP_LB]]
125 // CHECK: [[OMP_ST_1:%.+]] = load i32, i32* [[OMP_ST]]
126 // CHECK: [[ADD_2:%.+]] = add nsw i32 [[OMP_LB_2]], [[OMP_ST_1]]
127 // CHECK: store i32 [[ADD_2]], i32* [[OMP_LB]]
128 // CHECK: [[OMP_UB_5:%.+]] = load i32, i32* [[OMP_UB]]
129 // CHECK: [[OMP_ST_2:%.+]] = load i32, i32* [[OMP_ST]]
130 // CHECK: [[ADD_3:%.+]] = add nsw i32 [[OMP_UB_5]], [[OMP_ST_2]]
131 // CHECK: store i32 [[ADD_3]], i32* [[OMP_UB]]
132 
133 // CHECK: [[DISPATCH_END]]
134 // CHECK: call void @__kmpc_for_static_fini(
135 // CHECK: ret void
136 
137 #endif
138