Skip to content

Commit 78d80a1

Browse files
AlexeySachkovromanovvlad
AlexeySachkov
authored and
romanovvlad
committed
[SYCL] Add 'cl::sycl::intel::experimental::printf' (#835)
New built-in is mapped to OpenCL 'printf' built-in by using SPIR-V friendly IR. Signed-off-by: Alexey Sachkov <alexey.sachkov@intel.com>
1 parent f607520 commit 78d80a1

File tree

3 files changed

+209
-0
lines changed

3 files changed

+209
-0
lines changed

‎sycl/include/CL/sycl.hpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,7 @@
2121
#include <CL/sycl/handler.hpp>
2222
#include <CL/sycl/id.hpp>
2323
#include <CL/sycl/image.hpp>
24+
#include <CL/sycl/intel/builtins.hpp>
2425
#include <CL/sycl/intel/function_pointer.hpp>
2526
#include <CL/sycl/intel/sub_group.hpp>
2627
#include <CL/sycl/item.hpp>

‎sycl/include/CL/sycl/intel/builtins.hpp

Lines changed: 77 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,77 @@
1+
//==------ builtins.hpp - Non-standard SYCL built-in functions -------------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#pragma once
10+
11+
12+
#ifdef __SYCL_DEVICE_ONLY__
13+
#define CONSTANT_AS __attribute__((opencl_constant))
14+
// Note: __format string is declared in constant address space to be compatible
15+
// with OpenCL C
16+
extern int __spirv_ocl_printf(const CONSTANT_AS char *__format, ...);
17+
#else
18+
#define CONSTANT_AS
19+
#endif
20+
21+
namespace cl {
22+
namespace sycl {
23+
namespace intel {
24+
namespace experimental {
25+
26+
// Provides functionality to print data from kernels in a C way:
27+
// - On non-host devices this function is directly mapped to printf from
28+
// OpenCL C
29+
// - On host device, this function should be equivalent to standard printf
30+
// function from C/C++.
31+
//
32+
// Please refer to corresponding section in OpenCL C specification to find
33+
// information about format string and its differences from standard C rules.
34+
//
35+
// This function is placed under 'experimental' namespace on purpose, because it
36+
// has too much caveats you need to be aware of before using it. Please find
37+
// them below and read carefully before using it:
38+
//
39+
// - According to the OpenCL spec, the format string must be
40+
// resolvable at compile time i.e. cannot be dynamically created by the
41+
// executing program.
42+
//
43+
// - According to the OpenCL spec, the format string must reside in constant
44+
// address space. This requires to perform "tricky" declarations of them, see
45+
// test/built-ins/printf.cpp for examples
46+
// FIXME: this potentially can be done on SYCL FE side automatically
47+
//
48+
// - The format string is interpreted according to the OpenCL C spec, where all
49+
// data types has fixed size, opposed to C++ types which doesn't guarantee
50+
// the exact width of particular data types (except, may be, char). This might
51+
// lead to unexpected result, for example: %ld in OpenCL C means that printed
52+
// argument has 'long' type which is 64-bit wide by the OpenCL C spec. However,
53+
// by C++ spec long is just at least 32-bit wide, so, you need to ensure (by
54+
// performing a cast, for example) that if you use %ld specifier, you pass
55+
// 64-bit argument to the cl::sycl::experimental::printf
56+
//
57+
// - OpenCL spec defines several additional features, like, for example, 'v'
58+
// modifier which allows to print OpenCL vectors: note that these features are
59+
// not available on host device and therefore their usage should be either
60+
// guarded using __SYCL_DEVICE_ONLY__ preprocessor macro or avoided in favor
61+
// of more portable solutions if needed
62+
//
63+
template <typename... Args>
64+
int printf(const CONSTANT_AS char *__format, Args... args) {
65+
#ifdef __SYCL_DEVICE_ONLY__
66+
return __spirv_ocl_printf(__format, args...);
67+
#else
68+
return ::printf(__format, args...);
69+
#endif
70+
}
71+
72+
} // namespace experimental
73+
} // namespace intel
74+
} // namespace sycl
75+
} // namespace cl
76+
77+
#undef CONSTANT_AS

‎sycl/test/built-ins/printf.cpp

Lines changed: 131 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,131 @@
1+
// RUN: %clangxx -fsycl %s -o %t.out
2+
// RUN: env SYCL_DEVICE_TYPE=HOST %t.out | FileCheck %s
3+
// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
4+
// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
5+
// RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
6+
7+
#include <CL/sycl.hpp>
8+
9+
#include <cstdint>
10+
#include <iostream>
11+
12+
using namespace cl::sycl;
13+
14+
// According to OpenCL C spec, the format string must be in constant address
15+
// space
16+
#ifdef __SYCL_DEVICE_ONLY__
17+
#define CONSTANT __attribute__((opencl_constant))
18+
#else
19+
#define CONSTANT
20+
#endif
21+
22+
// This is one of the possible ways to define a format string in a correct
23+
// address space
24+
static const CONSTANT char format_hello_world[] = "Hello, World!\n";
25+
26+
// Static isn't really needed if you define it in global scope
27+
const CONSTANT char format_int[] = "%d\n";
28+
29+
static const CONSTANT char format_vec[] = "%d,%d,%d,%d\n";
30+
31+
const CONSTANT char format_hello_world_2[] = "%lu: Hello, World!\n";
32+
33+
int main() {
34+
{
35+
default_selector Selector;
36+
queue Queue(Selector);
37+
38+
Queue.submit([&](handler &CGH) {
39+
CGH.single_task<class integral>([=]() {
40+
// String
41+
intel::experimental::printf(format_hello_world);
42+
// Due to a bug in Intel CPU Runtime for OpenCL on Windows, information
43+
// printed using such format strings (without %-specifiers) might
44+
// appear in different order if output is redirected to a file or
45+
// another app
46+
// FIXME: strictly check output order once the bug is fixed
47+
// CHECK: {{(Hello, World!)?}}
48+
49+
// Integral types
50+
intel::experimental::printf(format_int, (int32_t)123);
51+
intel::experimental::printf(format_int, (int32_t)-123);
52+
// CHECK: 123
53+
// CHECK-NEXT: -123
54+
55+
// Floating point types
56+
{
57+
// You can declare format string in non-global scope, but in this case
58+
// static keyword is required
59+
static const CONSTANT char format[] = "%f\n";
60+
intel::experimental::printf(format, 33.4f);
61+
intel::experimental::printf(format, -33.4f);
62+
}
63+
// CHECK-NEXT: 33.4
64+
// CHECK-NEXT: -33.4
65+
66+
// Vectors
67+
cl::sycl::vec<int, 4> v4{5, 6, 7, 8};
68+
#ifdef __SYCL_DEVICE_ONLY__
69+
// On device side, vectors can be printed via native OpenCL types:
70+
using ocl_int4 = cl::sycl::vec<int, 4>::vector_t;
71+
{
72+
static const CONSTANT char format[] = "%v4d\n";
73+
intel::experimental::printf(format, (ocl_int4)v4);
74+
}
75+
76+
// However, you are still able to print them by-element:
77+
{
78+
intel::experimental::printf(format_vec, (int32_t)v4.w(),
79+
(int32_t)v4.z(), (int32_t)v4.y(),
80+
(int32_t)v4.x());
81+
}
82+
#else
83+
// On host side you always have to print them by-element:
84+
intel::experimental::printf(format_vec, (int32_t)v4.x(),
85+
(int32_t)v4.y(), (int32_t)v4.z(),
86+
(int32_t)v4.w());
87+
intel::experimental::printf(format_vec, (int32_t)v4.w(),
88+
(int32_t)v4.z(), (int32_t)v4.y(),
89+
(int32_t)v4.x());
90+
#endif // __SYCL_DEVICE_ONLY__
91+
// CHECK-NEXT: 5,6,7,8
92+
// CHECK-NEXT: 8,7,6,5
93+
94+
// Pointers
95+
int a = 5;
96+
int *Ptr = &a;
97+
// According to OpenCL spec, argument should be a void pointer
98+
{
99+
static const CONSTANT char format[] = "%p\n";
100+
intel::experimental::printf(format, (void *)Ptr);
101+
}
102+
// CHECK-NEXT: {{(0x)?[0-9a-fA-F]+$}}
103+
});
104+
});
105+
Queue.wait();
106+
107+
// printf in parallel_for
108+
Queue.submit([&](handler &CGH) {
109+
CGH.parallel_for<class stream_string>(range<1>(10), [=](id<1> i) {
110+
// cast to uint64_t to be sure that we pass 64-bit unsigned value
111+
intel::experimental::printf(format_hello_world_2, (uint64_t)i.get(0));
112+
});
113+
});
114+
Queue.wait();
115+
// CHECK-NEXT: {{[0-9]+}}: Hello, World!
116+
// CHECK-NEXT: {{[0-9]+}}: Hello, World!
117+
// CHECK-NEXT: {{[0-9]+}}: Hello, World!
118+
// CHECK-NEXT: {{[0-9]+}}: Hello, World!
119+
// CHECK-NEXT: {{[0-9]+}}: Hello, World!
120+
// CHECK-NEXT: {{[0-9]+}}: Hello, World!
121+
// CHECK-NEXT: {{[0-9]+}}: Hello, World!
122+
// CHECK-NEXT: {{[0-9]+}}: Hello, World!
123+
// CHECK-NEXT: {{[0-9]+}}: Hello, World!
124+
// CHECK-NEXT: {{[0-9]+}}: Hello, World!
125+
}
126+
127+
// FIXME: strictly check output order once the bug mentioned above is fixed
128+
// CHECK: {{(Hello, World!)?}}
129+
130+
return 0;
131+
}

0 commit comments

Comments
 (0)