From ba3b741eccbc16a71454a0fa90c5055385e4c080 Mon Sep 17 00:00:00 2001 From: LU-JOHN Date: Wed, 29 May 2024 04:37:31 -0500 Subject: [PATCH] Ensure 'printf' is translated to a printf extended insn (#2581) Calls to printf can generated by a user call to __builtin_printf. Represent this in SPIRV with a printf extended instruction instead of a SPIRV call to "printf". A SPIRV call to a variadic function will fail in spirv-val. Signed-off-by: Lu, John Original commit: https://github.com/KhronosGroup/SPIRV-LLVM-Translator/commit/d4098cdd72bce73 --- llvm-spirv/lib/SPIRV/SPIRVUtil.cpp | 2 +- .../builtin_printf.ll | 42 +++++++++++++++++++ llvm-spirv/test/transcoding/BuiltinPrintf.cl | 18 ++++++++ llvm-spirv/test/transcoding/Printf.cl | 19 +++++++++ 4 files changed, 80 insertions(+), 1 deletion(-) create mode 100644 llvm-spirv/test/extensions/EXT/SPV_EXT_relaxed_printf_string_address_space/builtin_printf.ll create mode 100644 llvm-spirv/test/transcoding/BuiltinPrintf.cl create mode 100644 llvm-spirv/test/transcoding/Printf.cl diff --git a/llvm-spirv/lib/SPIRV/SPIRVUtil.cpp b/llvm-spirv/lib/SPIRV/SPIRVUtil.cpp index f25175655772b..f10536cc892ac 100644 --- a/llvm-spirv/lib/SPIRV/SPIRVUtil.cpp +++ b/llvm-spirv/lib/SPIRV/SPIRVUtil.cpp @@ -442,7 +442,7 @@ bool getSPIRVBuiltin(const std::string &OrigName, spv::BuiltIn &B) { // if true is returned bool oclIsBuiltin(StringRef Name, StringRef &DemangledName, bool IsCpp) { if (Name == "printf") { - DemangledName = Name; + DemangledName = "__spirv_ocl_printf"; return true; } if (isNonMangledOCLBuiltin(Name)) { diff --git a/llvm-spirv/test/extensions/EXT/SPV_EXT_relaxed_printf_string_address_space/builtin_printf.ll b/llvm-spirv/test/extensions/EXT/SPV_EXT_relaxed_printf_string_address_space/builtin_printf.ll new file mode 100644 index 0000000000000..346d2e39e9e5d --- /dev/null +++ b/llvm-spirv/test/extensions/EXT/SPV_EXT_relaxed_printf_string_address_space/builtin_printf.ll @@ -0,0 +1,42 @@ +; Test that calls to "printf" are mapped to OpenCL Extended instruction "printf" +; Also ensure that spirv-val can validate format strings in non-constant space +; +; Testcase derived from: +; #include +; int main() { +; sycl::queue queue; +; queue.submit([&](sycl::handler &cgh) { +; cgh.single_task([] { +; __builtin_printf("%s, %s %d %d %d %s!\n", "Hello", "world", 1, 2, 3, "Bam"); +; }); +; }); +; } + +; RUN: llvm-as %s -o %t.bc +; RUN: not llvm-spirv %t.bc -o %t.spv 2>&1 | FileCheck %s --check-prefix=CHECK-WO-EXT + +; RUN: llvm-spirv -spirv-text %t.bc -o %t.spt --spirv-ext=+SPV_EXT_relaxed_printf_string_address_space +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV +; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_EXT_relaxed_printf_string_address_space +; Change TODO to RUN when spirv-val allows non-constant printf formats +; TODO: spirv-val %t.spv + + +; CHECK-WO-EXT: RequiresExtension: Feature requires the following SPIR-V extension: +; CHECK-WO-EXT: SPV_EXT_relaxed_printf_string_address_space extension should be allowed to translate this module, because this LLVM module contains the printf function with format string, whose address space is not equal to 2 (constant). + +; CHECK-SPIRV: Extension "SPV_EXT_relaxed_printf_string_address_space" +; CHECK-SPIRV: ExtInst [[#]] [[#]] [[#]] printf [[#]] + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1" +target triple = "spir64-unknown-unknown" + +@.str = external addrspace(1) constant [21 x i8] + +define spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_() { +entry: + %call.i = tail call spir_func i32 (ptr addrspace(4), ...) @printf(ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str to ptr addrspace(4)), ptr addrspace(4) null, ptr addrspace(4) null, i32 0, i32 0, i32 0, ptr addrspace(4) null) + ret void +} + +declare spir_func i32 @printf(ptr addrspace(4), ...) diff --git a/llvm-spirv/test/transcoding/BuiltinPrintf.cl b/llvm-spirv/test/transcoding/BuiltinPrintf.cl new file mode 100644 index 0000000000000..31016a85b6a1b --- /dev/null +++ b/llvm-spirv/test/transcoding/BuiltinPrintf.cl @@ -0,0 +1,18 @@ +// Ensure __builtin_printf is translated to a SPIRV printf instruction + +// RUN: %clang_cc1 -triple spir-unknown-unknown -emit-llvm-bc %s -o %t.bc +// RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV +// RUN: llvm-spirv %t.bc -o %t.spv +// Change TODO to RUN when spirv-val allows array of 8-bit ints for format +// TODO: spirv-val %t.spv +// RUN: llvm-spirv -r %t.spv -o %t.rev.bc +// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM + +// CHECK-SPIRV: ExtInst [[#]] [[#]] [[#]] printf [[#]] +// CHECK-LLVM: call spir_func i32 (ptr addrspace(2), ...) @printf(ptr addrspace(2) {{.*}}) + +kernel void BuiltinPrintf() { + __builtin_printf("Hello World"); +} + + diff --git a/llvm-spirv/test/transcoding/Printf.cl b/llvm-spirv/test/transcoding/Printf.cl new file mode 100644 index 0000000000000..b36585303bb37 --- /dev/null +++ b/llvm-spirv/test/transcoding/Printf.cl @@ -0,0 +1,19 @@ +// Ensure printf is translated to a SPIRV printf instruction + +// RUN: %clang_cc1 -triple spir-unknown-unknown -emit-llvm-bc %s -o %t.bc -finclude-default-header +// RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV +// RUN: llvm-spirv %t.bc -o %t.spv +// Change TODO to RUN when spirv-val allows array of 8-bit ints for format +// TODO: spirv-val %t.spv +// RUN: llvm-spirv -r %t.spv -o %t.rev.bc +// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM + +// CHECK-SPIRV: ExtInst [[#]] [[#]] [[#]] printf [[#]] +// CHECK-LLVM: call spir_func i32 (ptr addrspace(2), ...) @printf(ptr addrspace(2) {{.*}}) + + +kernel void Printf() { + printf("Hello World"); +} + +