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 { [4 x i8] } 34 // CHECK-DAG: [[GLOBAL_RD:@.+]] = weak global [{{[0-9]+}} x [{{[0-9]+}} x [[MEM_TY]]]] zeroinitializer 35 // CHECK-DAG: [[GLOBAL_RD_PTR:@.+]] = weak unnamed_addr constant i8* getelementptr inbounds ([{{[0-9]+}} x [{{[0-9]+}} x [[MEM_TY]]]], [{{[0-9]+}} x [{{[0-9]+}} x [[MEM_TY]]]]* [[GLOBAL_RD]], i{{[0-9]+}} 0, i{{[0-9]+}} 0, i{{[0-9]+}} 0, i{{[0-9]+}} 0, i{{[0-9]+}} 0) 36 // CHECK-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* null 37 // CHECK-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i{{64|32}} 4 38 39 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l12}}_worker() 40 // CHECK: call void @llvm.nvvm.barrier0() 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: [[GLOBAL_RD:%.+]] = load i8*, i8** [[GLOBAL_RD_PTR]], 49 // CHECK: [[SIZE:%.+]] = load i{{64|32}}, i{{64|32}}* [[KERNEL_SIZE]], 50 // CHECK: call void @__kmpc_get_team_static_memory(i8* [[GLOBAL_RD]], i{{64|32}} [[SIZE]], i16 0, 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 @llvm.nvvm.barrier0() 56 // CHECK: call void @llvm.nvvm.barrier0() 57 // CHECK: call void @__kmpc_end_sharing_variables() 58 // CHECK: call void @__kmpc_restore_team_static_memory(i16 0) 59 // CHECK: call void @__kmpc_kernel_deinit(i16 1) 60 61 // CHECK: define internal void @__omp_outlined__( 62 // CHECK: alloca 63 // CHECK: alloca 64 // CHECK: alloca 65 // CHECK: alloca 66 // CHECK: [[OMP_IV:%.*]] = alloca i32 67 // CHECK: store i32 0, {{.*}} [[OMP_LB:%.+]], 68 // CHECK: store i32 9, {{.*}} [[OMP_UB:%.+]], 69 // CHECK: store i32 1, {{.*}} [[OMP_ST:%.+]], 70 // CHECK: call void @__kmpc_for_static_init_4({{.*}} i32 33, {{.*}} [[OMP_LB]], {{.*}} [[OMP_UB]], {{.*}} [[OMP_ST]], i32 1, i32 1) 71 // CHECK: br label %[[OMP_DISPATCH_COND:.+]] 72 73 // CHECK: [[OMP_DISPATCH_COND]] 74 // CHECK: [[OMP_UB_1:%.+]] = load {{.*}} [[OMP_UB]] 75 // CHECK: [[COMP_1:%.+]] = icmp sgt {{.*}} [[OMP_UB_1]] 76 // CHECK: br i1 [[COMP_1]], label %[[COND_TRUE:.+]], label %[[COND_FALSE:.+]] 77 78 // CHECK: [[COND_TRUE]] 79 // CHECK: br label %[[COND_END:.+]] 80 81 // CHECK: [[COND_FALSE]] 82 // CHECK: [[OMP_UB_2:%.+]] = load {{.*}}* [[OMP_UB]] 83 // CHECK: br label %[[COND_END]] 84 85 // CHECK: [[COND_END]] 86 // CHECK: [[COND_RES:%.+]] = phi i32 [ 9, %[[COND_TRUE]] ], [ [[OMP_UB_2]], %[[COND_FALSE]] ] 87 // CHECK: store i32 [[COND_RES]], i32* [[OMP_UB]] 88 // CHECK: [[OMP_LB_1:%.+]] = load i32, i32* [[OMP_LB]] 89 // CHECK: store i32 [[OMP_LB_1]], i32* [[OMP_IV]] 90 // CHECK: [[OMP_IV_1:%.+]] = load i32, i32* [[OMP_IV]] 91 // CHECK: [[OMP_UB_3:%.+]] = load i32, i32* [[OMP_UB]] 92 // CHECK: [[COMP_2:%.+]] = icmp sle i32 [[OMP_IV_1]], [[OMP_UB_3]] 93 // CHECK: br i1 [[COMP_2]], label %[[DISPATCH_BODY:.+]], label %[[DISPATCH_END:.+]] 94 95 // CHECK: [[DISPATCH_BODY]] 96 // CHECK: br label %[[OMP_INNER_FOR_COND:.+]] 97 98 // CHECK: [[OMP_INNER_FOR_COND]] 99 // CHECK: [[OMP_IV_2:%.+]] = load i32, i32* [[OMP_IV]] 100 // CHECK: [[OMP_UB_4:%.+]] = load i32, i32* [[OMP_UB]] 101 // CHECK: [[COMP_3:%.+]] = icmp sle i32 [[OMP_IV_2]], [[OMP_UB_4]] 102 // CHECK: br i1 [[COMP_3]], label %[[OMP_INNER_FOR_BODY:.+]], label %[[OMP_INNER_FOR_END:.+]] 103 104 // CHECK: [[OMP_INNER_FOR_BODY]] 105 // CHECK: br label %[[OMP_BODY_CONTINUE:.+]] 106 107 // CHECK: [[OMP_BODY_CONTINUE]] 108 // CHECK: br label %[[OMP_INNER_FOR_INC:.+]] 109 110 // CHECK: [[OMP_INNER_FOR_INC]] 111 // CHECK: [[OMP_IV_3:%.+]] = load i32, i32* [[OMP_IV]] 112 // CHECK: [[ADD_1:%.+]] = add nsw i32 [[OMP_IV_3]], 1 113 // CHECK: store i32 [[ADD_1]], i32* [[OMP_IV]] 114 // CHECK: br label %[[OMP_INNER_FOR_COND]] 115 116 // CHECK: [[OMP_INNER_FOR_COND]] 117 // CHECK: br label %[[OMP_DISPATCH_INC:.+]] 118 119 // CHECK: [[OMP_DISPATCH_INC]] 120 // CHECK: [[OMP_LB_2:%.+]] = load i32, i32* [[OMP_LB]] 121 // CHECK: [[OMP_ST_1:%.+]] = load i32, i32* [[OMP_ST]] 122 // CHECK: [[ADD_2:%.+]] = add nsw i32 [[OMP_LB_2]], [[OMP_ST_1]] 123 // CHECK: store i32 [[ADD_2]], i32* [[OMP_LB]] 124 // CHECK: [[OMP_UB_5:%.+]] = load i32, i32* [[OMP_UB]] 125 // CHECK: [[OMP_ST_2:%.+]] = load i32, i32* [[OMP_ST]] 126 // CHECK: [[ADD_3:%.+]] = add nsw i32 [[OMP_UB_5]], [[OMP_ST_2]] 127 // CHECK: store i32 [[ADD_3]], i32* [[OMP_UB]] 128 129 // CHECK: [[DISPATCH_END]] 130 // CHECK: call void @__kmpc_for_static_fini( 131 // CHECK: ret void 132 133 #endif 134