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
11 #define __device__ __attribute__((device))
13 extern "C" __device__ int printf(const char *format, ...);
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");
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");