[CUDA] printfs results in ptxas error: Call has wrong number of parameters

[fwyzard]

Testing the latest iteration of #1091 (as of 44af06f2eb0) on top of the HEAD of the sycl branch (as of 7743e86) with a SYCLE kernel that calls printf, the compilation for the CUDA backend fails with

ptxas /tmp/fwyzard/main_oneapi-2cd1b3.s, line 516; error   : Call has wrong number of parameters
ptxas /tmp/fwyzard/main_oneapi-2cd1b3.s, line 721; error   : Call has wrong number of parameters
ptxas /tmp/fwyzard/main_oneapi-2cd1b3.s, line 890; error   : Call has wrong number of parameters
ptxas /tmp/fwyzard/main_oneapi-2cd1b3.s, line 1543; error   : Call has wrong number of parameters
ptxas /tmp/fwyzard/main_oneapi-2cd1b3.s, line 1643; error   : Call has wrong number of parameters
ptxas fatal   : Ptx assembly aborted due to errors
clang-11: error: ptxas command failed with exit code 255 (use -v to see invocation)

All errors point to similar PTX code:

        call.uni (retval0),
        _Z18__spirv_ocl_printfPU3AS4Kcz,
        (
        param0,
        param1
        );

where _Z18__spirv_ocl_printfPU3AS4Kcz demangles to __spirv_ocl_printf(char const AS4*, ...).

The same kernel compiles and runs with the OpenCL backend.

In case it's useful, the way I am using printf in kernel code is here and e.g. here.

The code compiles for the CUDA backend if I comment out the call to cl::sycl::intel::experimental::printf.

[fwyzard]

@Alexander-Johnston FYI

[Ruyk]

Thats an (undocumented?) intel extension, currently we are only implementing conformant SYCL features. We may do this later on if we have time.

[Ruyk]

@bader There is no extension document for this feature?

[bader]

@AlexeySachkov, could add documentation for the feature you added by #835 to https://github.com/intel/llvm/tree/sycl/sycl/doc/extensions, please?

[Ruyk]

Is this still a problem @fwyzard ?

[jeffhammond]

@Ruyk It's a problem for me today with the latest version.

[vrpascuzzi]

@Ruyk And me today with sycl-nightly/20210330. (@jeffhammond I swear, I'm not trolling your posts. Rather, seem to be some months behind you.)