Skip to content

Commit ba3b741

Browse files
LU-JOHNsys-ce-bb
authored andcommitted
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
1 parent 82607c6 commit ba3b741

File tree

4 files changed

+80
-1
lines changed

4 files changed

+80
-1
lines changed

llvm-spirv/lib/SPIRV/SPIRVUtil.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -442,7 +442,7 @@ bool getSPIRVBuiltin(const std::string &OrigName, spv::BuiltIn &B) {
442442
// if true is returned
443443
bool oclIsBuiltin(StringRef Name, StringRef &DemangledName, bool IsCpp) {
444444
if (Name == "printf") {
445-
DemangledName = Name;
445+
DemangledName = "__spirv_ocl_printf";
446446
return true;
447447
}
448448
if (isNonMangledOCLBuiltin(Name)) {
Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,42 @@
1+
; Test that calls to "printf" are mapped to OpenCL Extended instruction "printf"
2+
; Also ensure that spirv-val can validate format strings in non-constant space
3+
;
4+
; Testcase derived from:
5+
; #include <sycl/sycl.hpp>
6+
; int main() {
7+
; sycl::queue queue;
8+
; queue.submit([&](sycl::handler &cgh) {
9+
; cgh.single_task([] {
10+
; __builtin_printf("%s, %s %d %d %d %s!\n", "Hello", "world", 1, 2, 3, "Bam");
11+
; });
12+
; });
13+
; }
14+
15+
; RUN: llvm-as %s -o %t.bc
16+
; RUN: not llvm-spirv %t.bc -o %t.spv 2>&1 | FileCheck %s --check-prefix=CHECK-WO-EXT
17+
18+
; RUN: llvm-spirv -spirv-text %t.bc -o %t.spt --spirv-ext=+SPV_EXT_relaxed_printf_string_address_space
19+
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
20+
; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_EXT_relaxed_printf_string_address_space
21+
; Change TODO to RUN when spirv-val allows non-constant printf formats
22+
; TODO: spirv-val %t.spv
23+
24+
25+
; CHECK-WO-EXT: RequiresExtension: Feature requires the following SPIR-V extension:
26+
; 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).
27+
28+
; CHECK-SPIRV: Extension "SPV_EXT_relaxed_printf_string_address_space"
29+
; CHECK-SPIRV: ExtInst [[#]] [[#]] [[#]] printf [[#]]
30+
31+
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"
32+
target triple = "spir64-unknown-unknown"
33+
34+
@.str = external addrspace(1) constant [21 x i8]
35+
36+
define spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_() {
37+
entry:
38+
%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)
39+
ret void
40+
}
41+
42+
declare spir_func i32 @printf(ptr addrspace(4), ...)
Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,18 @@
1+
// Ensure __builtin_printf is translated to a SPIRV printf instruction
2+
3+
// RUN: %clang_cc1 -triple spir-unknown-unknown -emit-llvm-bc %s -o %t.bc
4+
// RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
5+
// RUN: llvm-spirv %t.bc -o %t.spv
6+
// Change TODO to RUN when spirv-val allows array of 8-bit ints for format
7+
// TODO: spirv-val %t.spv
8+
// RUN: llvm-spirv -r %t.spv -o %t.rev.bc
9+
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
10+
11+
// CHECK-SPIRV: ExtInst [[#]] [[#]] [[#]] printf [[#]]
12+
// CHECK-LLVM: call spir_func i32 (ptr addrspace(2), ...) @printf(ptr addrspace(2) {{.*}})
13+
14+
kernel void BuiltinPrintf() {
15+
__builtin_printf("Hello World");
16+
}
17+
18+

llvm-spirv/test/transcoding/Printf.cl

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
// Ensure printf is translated to a SPIRV printf instruction
2+
3+
// RUN: %clang_cc1 -triple spir-unknown-unknown -emit-llvm-bc %s -o %t.bc -finclude-default-header
4+
// RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
5+
// RUN: llvm-spirv %t.bc -o %t.spv
6+
// Change TODO to RUN when spirv-val allows array of 8-bit ints for format
7+
// TODO: spirv-val %t.spv
8+
// RUN: llvm-spirv -r %t.spv -o %t.rev.bc
9+
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
10+
11+
// CHECK-SPIRV: ExtInst [[#]] [[#]] [[#]] printf [[#]]
12+
// CHECK-LLVM: call spir_func i32 (ptr addrspace(2), ...) @printf(ptr addrspace(2) {{.*}})
13+
14+
15+
kernel void Printf() {
16+
printf("Hello World");
17+
}
18+
19+

0 commit comments

Comments
 (0)