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