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 // expected-no-diagnostics 7 extern int printf(const char *, ...); 8 9 // Check a simple call to printf end-to-end. 10 // CHECK: [[SIMPLE_PRINTF_TY:%[a-zA-Z0-9_]+]] = type { i32, i64, double } 11 int CheckSimple() { 12 // CHECK: define {{.*}}void [[T1:@__omp_offloading_.+CheckSimple.+]]_worker() 13 #pragma omp target 14 { 15 // Entry point. 16 // CHECK: define {{.*}}void [[T1]]() 17 // Alloca in entry block. 18 // CHECK: [[BUF:%[a-zA-Z0-9_]+]] = alloca [[SIMPLE_PRINTF_TY]] 19 20 // CHECK: {{call|invoke}} void [[T1]]_worker() 21 // CHECK: br label {{%?}}[[EXIT:.+]] 22 // 23 // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 24 // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 25 // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 26 // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]], 27 // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]] 28 // 29 // CHECK: [[MASTER]] 30 // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 31 // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 32 // CHECK: [[MTMP1:%.+]] = sub nuw i32 [[MNTH]], [[MWS]] 33 // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]] 34 35 // printf in master-only basic block. 36 // CHECK: [[FMT:%[0-9]+]] = load{{.*}}%fmt 37 const char* fmt = "%d %lld %f"; 38 // CHECK: [[PTR0:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]* [[BUF]], i32 0, i32 0 39 // CHECK: store i32 1, i32* [[PTR0]], align 4 40 // CHECK: [[PTR1:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]* [[BUF]], i32 0, i32 1 41 // CHECK: store i64 2, i64* [[PTR1]], align 8 42 // CHECK: [[PTR2:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]* [[BUF]], i32 0, i32 2 43 44 // CHECK: store double 3.0{{[^,]*}}, double* [[PTR2]], align 8 45 // CHECK: [[BUF_CAST:%[0-9]+]] = bitcast [[SIMPLE_PRINTF_TY]]* [[BUF]] to i8* 46 // CHECK: [[RET:%[0-9]+]] = call i32 @vprintf(i8* [[FMT]], i8* [[BUF_CAST]]) 47 printf(fmt, 1, 2ll, 3.0); 48 } 49 50 return 0; 51 } 52 53 void CheckNoArgs() { 54 // CHECK: define {{.*}}void [[T2:@__omp_offloading_.+CheckNoArgs.+]]_worker() 55 #pragma omp target 56 { 57 // Entry point. 58 // CHECK: define {{.*}}void [[T2]]() 59 60 // CHECK: {{call|invoke}} void [[T2]]_worker() 61 // CHECK: br label {{%?}}[[EXIT:.+]] 62 // 63 // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 64 // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 65 // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 66 // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]], 67 // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]] 68 // 69 // CHECK: [[MASTER]] 70 // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 71 // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 72 // CHECK: [[MTMP1:%.+]] = sub nuw i32 [[MNTH]], [[MWS]] 73 // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]] 74 75 // printf in master-only basic block. 76 // CHECK: call i32 @vprintf({{.*}}, i8* null){{$}} 77 printf("hello, world!"); 78 } 79 } 80 81 // Check that printf's alloca happens in the entry block, not inside the if 82 // statement. 83 int foo; 84 void CheckAllocaIsInEntryBlock() { 85 // CHECK: define {{.*}}void [[T3:@__omp_offloading_.+CheckAllocaIsInEntryBlock.+]]_worker() 86 #pragma omp target 87 { 88 // Entry point. 89 // CHECK: define {{.*}}void [[T3]]( 90 // Alloca in entry block. 91 // CHECK: alloca %printf_args 92 93 // CHECK: {{call|invoke}} void [[T3]]_worker() 94 // CHECK: br label {{%?}}[[EXIT:.+]] 95 // 96 // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() 97 // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 98 // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 99 // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]], 100 // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]] 101 // 102 // CHECK: [[MASTER]] 103 // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() 104 // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() 105 // CHECK: [[MTMP1:%.+]] = sub nuw i32 [[MNTH]], [[MWS]] 106 // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]] 107 108 if (foo) { 109 printf("%d", 42); 110 } 111 } 112 } 113