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 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc 5 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 6 // RUN: %clang_cc1 -verify -fopenmp -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 7 // expected-no-diagnostics 8 #ifndef HEADER 9 #define HEADER 10 11 // CHECK: [[MEM_TY:%.+]] = type { [128 x i8] } 12 // CHECK-DAG: {{@__omp_offloading_.+}}_l19_exec_mode = weak constant i8 1 13 // CHECK-DAG: internal unnamed_addr constant i{{64|32}} 4 14 15 template<typename tx> 16 tx ftemplate(int n) { 17 int i; 18 19 #pragma omp target teams distribute 20 for (i = 0; i < 10; ++i) 21 { 22 #pragma omp parallel 23 ++i; 24 } 25 26 return i; 27 } 28 29 int bar(int n){ 30 int a = 0; 31 32 a += ftemplate<char>(n); 33 34 return a; 35 } 36 37 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l19}}_worker() 38 // CHECK: ret void 39 40 // CHECK: define {{.*}}void {{@__omp_offloading_.+template.+l19}}() 41 42 // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 43 // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 44 // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 45 // CHECK-DAG: [[TH_LIMIT:%.+]] = sub nuw i32 [[NTH]], [[WS]] 46 // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]] 47 // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]] 48 // 49 // CHECK: [[WORKER]] 50 // CHECK: {{call|invoke}} void {{@__omp_offloading_.+template.+l19}}_worker() 51 // CHECK: br label {{%?}}[[EXIT:.+]] 52 // 53 // CHECK: [[CHECK_MASTER]] 54 // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 55 // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 56 // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 57 // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]], 58 // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]] 59 // 60 // CHECK: [[MASTER]] 61 // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 62 // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 63 // CHECK: [[MTMP1:%.+]] = sub nuw i32 [[MNTH]], [[MWS]] 64 // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]] 65 // CHECK: call void @__kmpc_get_team_static_memory(i16 0, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([[MEM_TY]], [[MEM_TY]] addrspace(3)* @{{.+}}, i32 0, i32 0, i32 0) to i8*), i{{64|32}} 4, i16 1, i8** addrspacecast (i8* addrspace(3)* [[BUF:@.+]] to i8**)) 66 // CHECK: [[PTR:%.+]] = load i8*, i8* addrspace(3)* [[BUF]], 67 // CHECK: [[RD:%.+]] = bitcast i8* [[PTR]] to [[GLOB_TY:%.+]]* 68 // CHECK: [[I_ADDR:%.+]] = getelementptr inbounds [[GLOB_TY]], [[GLOB_TY]]* [[RD]], i32 0, i32 0 69 // 70 // CHECK: call void @__kmpc_for_static_init_4( 71 // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* @{{.+}} to i8*), i16 1) 72 // CHECK: call void @__kmpc_begin_sharing_variables(i8*** [[SHARED_VARS_PTR:%.+]], i{{64|32}} 1) 73 // CHECK: [[SHARED_VARS_BUF:%.+]] = load i8**, i8*** [[SHARED_VARS_PTR]], 74 // CHECK: [[I_ADDR_BC:%.+]] = bitcast i32* [[I_ADDR]] to i8* 75 // CHECK: store i8* [[I_ADDR_BC]], i8** [[SHARED_VARS_BUF]], 76 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 77 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 78 // CHECK: call void @__kmpc_end_sharing_variables() 79 // CHECK: call void @__kmpc_for_static_fini( 80 // CHECK: br label {{%?}}[[TERMINATE:.+]] 81 // 82 // CHECK: [[TERMINATE]] 83 // CHECK: call void @__kmpc_kernel_deinit( 84 // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) 85 // CHECK: br label {{%?}}[[EXIT]] 86 // 87 // CHECK: [[EXIT]] 88 // CHECK: ret void 89 90 #endif 91