| 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 | |
|---|