32 lines · plain
1// REQUIRES: amdgpu-registered-target2// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm -disable-llvm-optzns -mprintf-kind=hostcall -fno-builtin-printf -fcuda-is-device \3// RUN: -o - %s | FileCheck --check-prefixes=CHECK,HOSTCALL %s4// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -emit-llvm -disable-llvm-optzns -mprintf-kind=hostcall -fno-builtin-printf -fcuda-is-device \5// RUN: -o - %s | FileCheck --check-prefixes=CHECK-AMDGCNSPIRV,HOSTCALL-AMDGCNSPIRV %s6// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm -disable-llvm-optzns -mprintf-kind=buffered -fno-builtin-printf -fcuda-is-device \7// RUN: -o - %s | FileCheck --check-prefixes=CHECK,BUFFERED %s8// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -emit-llvm -disable-llvm-optzns -mprintf-kind=buffered -fno-builtin-printf -fcuda-is-device \9// RUN: -o - %s | FileCheck --check-prefixes=CHECK-AMDGCNSPIRV,BUFFERED-AMDGCNSPIRV %s10 11#define __device__ __attribute__((device))12 13extern "C" __device__ int printf(const char *format, ...);14 15// CHECK-LABEL: @_Z4foo1v()16__device__ int foo1() {17 // HOSTCALL: call i64 @__ockl_printf_begin18 // HOSTCALL-AMDGCNSPIRV: call addrspace(4) i64 @__ockl_printf_begin19 // BUFFERED: call ptr addrspace(1) @__printf_alloc20 // BUFFERED-AMDGCNSPIRV: call addrspace(4) ptr addrspace(1) @__printf_alloc21 // CHECK-NOT: call i32 (ptr, ...) @printf22 // CHECK-AMDGCNSPIRV-NOT: call i32 (ptr, ...) @printf23 return __builtin_printf("Hello World\n");24}25 26// CHECK-LABEL: @_Z4foo2v()27__device__ int foo2() {28 // CHECK: call i32 (ptr, ...) @printf29 // CHECK-AMDGCNSPIRV: call spir_func addrspace(4) i32 (ptr addrspace(4), ...) @printf30 return printf("Hello World\n");31}32