__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.