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 <[email protected]>

Original commit:
KhronosGroup/SPIRV-LLVM-Translator@d4098cdd72bce73

Author: LU-JOHN

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)) {

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 <sycl/sycl.hpp>
+;   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), ...)

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");
+}
+
+

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");
+}
+
+