1// REQUIRES: amdgpu-registered-target 2// 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 %s 4// 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 %s 6// 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 %s 8// 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 %s 10 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_begin 18 // HOSTCALL-AMDGCNSPIRV: call addrspace(4) i64 @__ockl_printf_begin 19 // BUFFERED: call ptr addrspace(1) @__printf_alloc 20 // BUFFERED-AMDGCNSPIRV: call addrspace(4) ptr addrspace(1) @__printf_alloc 21 // CHECK-NOT: call i32 (ptr, ...) @printf 22 // CHECK-AMDGCNSPIRV-NOT: call i32 (ptr, ...) @printf 23 return __builtin_printf("Hello World\n"); 24} 25 26// CHECK-LABEL: @_Z4foo2v() 27__device__ int foo2() { 28 // CHECK: call i32 (ptr, ...) @printf 29 // CHECK-AMDGCNSPIRV: call spir_func addrspace(4) i32 (ptr addrspace(4), ...) @printf 30 return printf("Hello World\n"); 31} 32