__builtin_printf not diagnosed but results in invalid SPIR-V

[hvdijk]

Describe the bug

Kernels are not permitted to call printf (see issue #487). This gets diagnosed correctly for regular calls, but calls to __builtin_printf go undiagnosed and result in invalid SPIR-V.

To Reproduce
Please describe the steps to reproduce the behavior:

1. Include code snippet as short as possible

#include <sycl/sycl.hpp>
int main() {
  sycl::queue queue;
  queue.submit([&](sycl::handler &cgh) {
    cgh.single_task([] {
      __builtin_printf("%s, %s!\n", "Hello", "world");
    });
  });
}

2. Specify the command which should be used to compile the program

clang++ -save-temps -fsycl sycl.cc -o sycl
for f in $(cat sycl-sycl-spir64-unknown-unknown-*.txt); do spirv-val $f; done

3. Specify the comment which should be used to launch the program

N/A

4. Indicate what is wrong and what was expected

This program should have either been rejected by the frontend as it would have been if __builtin_printf had been avoided and printf had been used instead:

sycl.cc:6:43: error: SYCL kernel cannot call a variadic function
    6 |       printf("%s, %s!\n", "Hello", "world");
      |                                           ^

Instead, SPIR-V is generated that declares printf as a function taking only a format string, but nonetheless calls it with three arguments, resulting in

error: line 169: OpFunctionCall Function <id>'s parameter count does not match the argument count.
  %call_i = OpFunctionCall %uint %printf %47 %49 %50

The precise results of actually running it depend on the driver used, but generally, it just does not work and cannot be expected to work.

Environment (please complete the following information):

• OS: Linux
• Target device and vendor: N/A
• DPC++ version: clang version 18.0.0 (https://github.com/intel/llvm 34a0635)
• Dependencies version: N/A

Additional context
Add any other context about the problem here.

[AlexeySachkov]

Hi @hvdijk, thanks for the report.

I would close this as not a bug, because the C++ program you have is essentially ill-formed due to use of reserved identifiers:

cppreference: Identifiers:

Identifiers that appear as a token or preprocessing token (i.e., not in user-defined-string-literal like operator ""id) (since C++11) of one of the following forms are reserved:

• identifiers with a double underscore anywhere;
• ...

"Reserved" here means that the standard library headers #define or declare such identifiers for their internal needs, the compiler may predefine non-standard identifiers of that kind, and that name mangling algorithm may assume that some of these identifiers are not in use. If the programmer uses such identifiers, the program is ill-formed, no diagnostic required.

[hvdijk]

Hi @hvdijk, thanks for the report.

I would close this as not a bug, because the C++ program you have is essentially ill-formed due to use of reserved identifiers:

As far as the C++ standard is concerned, you are right, but in the context of a specific compiler, if the use of the reserved identifiers is covered by a documented and fully supported extension, that is different. Imagine if DPC++ were to take the paragraph you quote as a basis for rejecting all programs that do #ifdef __SYCL_DEVICE_ONLY__. As far as the C++ standard is concerned, that might be valid. But it's clearly wrong.

But, actually, I am noticing something else now: in #7483, __builtin_printf was specifically added as an accepted extension in SYCL device code and a test for it was added. Despite the fact that in SPIR-V, it does not and cannot work.

[Naghasan]

Despite the fact that in SPIR-V, it does not and cannot work.

It can https://registry.khronos.org/SPIR-V/specs/unified1/OpenCL.ExtendedInstructionSet.100.html#printf

It is just improperly lowered by the translator.

Note: DPCPP is also using an extension because mapping the format string to the constant address space is problematic in SYCL.

[hvdijk]

It can https://registry.khronos.org/SPIR-V/specs/unified1/OpenCL.ExtendedInstructionSet.100.html#printf

I stand corrected! You can see a remnant of what I originally included in my report, "either ... or an extension should be used", I took out the "or an extension should be used" because I could not find an extension for variadic functions but left the "either" in by mistake. I had not imagined there was an extension for printf specifically. :) Should I update the original message to include that?

[LU-JOHN]

Compling with clang++ from (May 2, 2024) with the command:

clang++ -fsycl test.cpp

Produces the error message:

RequiresExtension: Feature requires the following SPIR-V extension:
Either 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).
%call.i = call spir_func i32 @Z18__spirv_ocl_printfPU3AS4cS0_S0(ptr addrspace(4) noundef %3, ptr addrspace(4) noundef %4, ptr addrspace(4) noundef %5) #6
llvm-foreach:
clang++: error: llvm-spirv command failed with exit code 19 (use -v to see invocation)

Compiling with:

clang++ -fsycl -Xspirv-translator --spirv-ext=+SPV_EXT_relaxed_printf_string_address_space test.cpp

produces an executable with no errors.

@hvdijk is this behavior okay?

[hvdijk]

@hvdijk is this behavior okay?

Having it use an extension is fine, but it seems like the result still does not pass validation:

error: line 165: OpenCL.std printf: expected Format storage class to be UniformConstant
%call_i = OpExtInst %uint %1 printf %48 %49 %50

Is this an extension that is not yet supported in SPIRV-Tools, or is there something else going on?

[LU-JOHN]

Having it use an extension is fine, but it seems like the result still does not pass validation:

error: line 165: OpenCL.std printf: expected Format storage class to be UniformConstant
%call_i = OpExtInst %uint %1 printf %48 %49 %50

At what point do you see this error message? What version of the backend tools do you have?
I am able to compile and run the test program fine:

lujohn@scsel-tl-03:~/exp$ SYCL_PI_TRACE=1 ./a.out
SYCL_PI_TRACE[basic]: Plugin found and successfully loaded: libpi_opencl.so [ PluginVersion: 15.48.1 ]
SYCL_PI_TRACE[basic]: Plugin found and successfully loaded: libpi_level_zero.so [ PluginVersion: 15.48.1 ]
SYCL_PI_TRACE[basic]: Plugin found and successfully loaded: libpi_unified_runtime.so [ PluginVersion: 15.48.1 ]
SYCL_PI_TRACE[all]: Requested device_type: info::device_type::automatic
SYCL_PI_TRACE[all]: Requested device_type: info::device_type::automatic
SYCL_PI_TRACE[all]: Requested device_type: info::device_type::automatic
SYCL_PI_TRACE[all]: Requested device_type: info::device_type::automatic
SYCL_PI_TRACE[all]: Selected device: -> final score = 1550
SYCL_PI_TRACE[all]: platform: Intel(R) Level-Zero
SYCL_PI_TRACE[all]: device: Intel(R) Iris(R) Xe Graphics
Hello, world!

[hvdijk]

At what point do you see this error message?

When I run spirv-val like in my original message, using a fresh clone from current https://github.com/KhronosGroup/SPIRV-Tools

[LU-JOHN]

I can reproduce a different spirv-val error:

error: line 169: OpFunctionCall Function 's parameter count does not match the argument count.
%call_i = OpFunctionCall %uint %printf %47 %49 %50

I'll investigate if this is an issue with SPIRV-Tools.

[LU-JOHN]

PR to update spirv-val to validate printf correctly made in:

KhronosGroup/SPIRV-Tools#5667

to fix incorrect validation message:

error: line 169: OpFunctionCall Function 's parameter count does not match the argument count.
%call_i = OpFunctionCall %uint %printf %47 %49 %50

[github-actions]

Hi! There have been no updates for at least the last 60 days, though the issue has assignee(s).

@LU-JOHN, could you please take one of the following actions:

• provide an update if you have any
• unassign yourself if you're not looking / going to look into this issue
• mark this issue with the 'confirmed' label if you have confirmed the problem/request and our team should work on it
• close the issue if it has been resolved
• take any other suitable action.

Thanks!

[dm-vodopyanov]

PR to update spirv-val to validate printf correctly made in:

KhronosGroup/SPIRV-Tools#5667

to fix incorrect validation message:

error: line 169: OpFunctionCall Function 's parameter count does not match the argument count.
%call_i = OpFunctionCall %uint %printf %47 %49 %50

@LU-JOHN KhronosGroup/SPIRV-Tools#5667 is closed, not merged. Could you please provide what are the next steps required to resolve this issue? Or if it's already resolved, could you please close it?

[LU-JOHN]

llvm-spirv updated to use printf instruction from OpenCL.std in KhronosGroup/SPIRV-LLVM-Translator#2581 and has been merged.

spirv-val updated to allow printf calls with non-constant format strings in KhronosGroup/SPIRV-Tools#5677 and is awaiting merge.

[dm-vodopyanov]

@LU-JOHN thanks! If KhronosGroup/SPIRV-Tools#5677 finally fixes this issue, could you please add

Fixes https://github.com/intel/llvm/issues/11733

to the description of KhronosGroup/SPIRV-Tools#5677?

[dm-vodopyanov]

KhronosGroup/SPIRV-Tools#5677 and is awaiting merge

Merged. @LU-JOHN can the issue be closed now?

[hvdijk]

This looks like it's fixed to me, though the previously generated invalid SPIR-V was (IIRC) accepted by the Intel OpenCL driver, the newly generated valid SPIR-V results in errors there. With the original test program:

$ dpcppllvm/build/x86_64-linux/install/bin/clang++ -fsycl sycl.cc -Xspirv-translator --spirv-ext=+SPV_EXT_relaxed_printf_string_address_space -o sycl
$ LD_LIBRARY_PATH=dpcppllvm/build/x86_64-linux/install/lib ./sycl
terminate called after throwing an instance of 'sycl::_V1::exception'
  what():  The program was built for 1 devices
Build program log for 'Intel(R) Core(TM) i7-8700 CPU @ 3.20GHz':
Compilation started
Compilation done
Linking started
Linking done
Device build started
Options used by backend compiler:
Failed to build device program
CompilerException Failed to lookup symbol _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_
JIT session error: Symbols not found: [ _Z18__spirv_ocl_printfPU3AS4PcS1_S1_ ]
Failed to materialize symbols: { (main, { _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_ }) }


Aborted (core dumped)
$ LD_LIBRARY_PATH=dpcppllvm/build/x86_64-linux/install/lib dpcppllvm/build/x86_64-linux/install/bin/sycl-ls --verbose
[opencl:cpu][opencl:0] Intel(R) OpenCL, Intel(R) Core(TM) i7-8700 CPU @ 3.20GHz OpenCL 3.0 (Build 0) [2024.18.6.0.02_160000]
[opencl:cpu][opencl:1] Portable Computing Language, cpu-haswell-Intel(R) Core(TM) i7-8700 CPU @ 3.20GHz OpenCL 3.0 PoCL HSTR: cpu-x86_64-pc-linux-gnu-haswell [5.0+debian]
[opencl:fpga][opencl:2] Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device OpenCL 1.2  [2024.18.6.0.02_160000]
[native_cpu:cpu][native_cpu:0] SYCL_NATIVE_CPU, SYCL Native CPU 0.1 [0.0.0]

Platforms: 4
Platform [#1]:
    Version  : OpenCL 3.0 LINUX
    Name     : Intel(R) OpenCL
    Vendor   : Intel(R) Corporation
    Devices  : 1
        Device [#0]:
        Type              : cpu
        Version           : OpenCL 3.0 (Build 0)
        Name              : Intel(R) Core(TM) i7-8700 CPU @ 3.20GHz
        Vendor            : Intel(R) Corporation
        Driver            : 2024.18.6.0.02_160000
        Num SubDevices    : 0
        Num SubSubDevices : 0
        Aspects           : cpu fp16 fp64 online_compiler online_linker queue_profiling usm_device_allocations usm_host_allocations usm_shared_allocations usm_system_allocations usm_atomic_host_allocations usm_atomic_shared_allocations atomic64 ext_oneapi_srgb ext_oneapi_native_assert ext_intel_legacy_image ext_oneapi_ballot_group ext_oneapi_fixed_size_group ext_oneapi_opportunistic_group ext_oneapi_tangle_group ext_oneapi_private_alloca
        info::device::sub_group_sizes: 4 8 16 32 64
        Architecture: x86_64
Platform [#2]:
    Version  : OpenCL 3.0 PoCL 5.0+debian  Linux, None+Asserts, RELOC, SPIR, LLVM 16.0.6, SLEEF, DISTRO, POCL_DEBUG
    Name     : Portable Computing Language
    Vendor   : The pocl project
    Devices  : 1
        Device [#1]:
        Type              : cpu
        Version           : OpenCL 3.0 PoCL HSTR: cpu-x86_64-pc-linux-gnu-haswell
        Name              : cpu-haswell-Intel(R) Core(TM) i7-8700 CPU @ 3.20GHz
        Vendor            : GenuineIntel
        Driver            : 5.0+debian
        Num SubDevices    : 0
        Num SubSubDevices : 0
        Aspects           : cpu fp64 online_compiler online_linker queue_profiling usm_device_allocations usm_host_allocations usm_shared_allocations atomic64 ext_oneapi_srgb ext_intel_legacy_image ext_oneapi_ballot_group ext_oneapi_fixed_size_group ext_oneapi_opportunistic_group ext_oneapi_tangle_group ext_oneapi_limited_graph ext_oneapi_private_alloca
        info::device::sub_group_sizes: 1 2 4 8 16 32 64 128 256 512
        Architecture: SYCL Exception encountered: Native API failed. Native API returns: -30 (PI_ERROR_INVALID_VALUE)

I'm not sure what the right place for reporting that is. (It does not work with other implementations either, but those are not Intel's responsibility.)

[dm-vodopyanov]

@hvdijk thanks for confirming this!

@LU-JOHN: is this a regression caused by the patches mentioned above? If yes, it should be fixed as part of this GH issue, if not, @hvdijk, could you please create a separate GH issue on that?

[LU-JOHN]

The OpenCL runtime team will be investigating the missing symbol issue.

[LU-JOHN]

@LU-JOHN thanks! If KhronosGroup/SPIRV-Tools#5677 finally fixes this issue, could you please add

Fixes https://github.com/intel/llvm/issues/11733

to the description of KhronosGroup/SPIRV-Tools#5677?

Updated description.