Kernel ID is always zero with CUDA Backend and no-sycl-libspirv flag

[MarkusBuettner]

Describe the bug
item.get_id() and id.get(0) always return 0 on the GPU if the program is compiled with -fno-sycl-libspirv

To Reproduce
Sample code: From ENCCS: Hello, SYCL!:

#include <iostream>
#include <sycl/sycl.hpp>

int main() {
  const std::string secret{"Ifmmp-!xpsme\"\012J(n!tpssz-!Ebwf/!"
                           "J(n!bgsbje!J!dbo(u!ep!uibu/!.!IBM\01"};

  const auto sz = secret.size();

  queue Q;

  std::cout << "Running on: " << Q.get_device().get_info<info::device::name>()
            << std::endl;

  char *result = malloc_shared<char>(sz, Q);
  std::memcpy(result, secret.data(), sz);

  // It does not matter whether I use tid.get_id(), tid, or sycl::id in the kernel
  Q.parallel_for(range<1>{sz}, [=](item<1> tid) { result[tid.get_id()] -= 1; }).wait();

  std::cout << result << "\n";

  free(result, Q);

  return 0;
}

The code is compiled with the command

clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fno-sycl-libspirv hello.cpp -o hello

The output on my machine is

Running on: NVIDIA Quadro RTX 6000
Hfmmp-!xpsme"
J(n!tpssz-!Ebwf/!J(n!bgsbje!J!dbo(u!ep!uibu/!.!IBM

but should be

Running on: NVIDIA Quadro RTX 6000
Hello, world!   I'm sorry, Dave. I'm afraid I can't do that. - HAL

When I compile the code above without the flag -fno-sycl-libspirv, everything works fine. And I'm getting the expected output.

Environment (please complete the following information):

• OS: Ubuntu 20.04
• NVIDIA GPU
• Clang compiled from sycl branch
• CUDA 11.3

** Additional information **
cuda-gdb shows that there are, in fact, 65 threads created on the device:

(cuda-gdb) info cuda threads
BlockIdx ThreadIdx To BlockIdx ThreadIdx Count         Virtual PC      Filename  Line 
Kernel 0
*  (0,0,0)   (0,0,0)     (0,0,0)  (64,0,0)    65 0x0000000000aa2050 sycl_test.cpp    22 

And printing tid shows the same object (?) on two different threads:

(cuda-gdb) cuda thread 30
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (30,0,0), device 0, sm 0, warp 0, lane 30]
22                        result[tid.get_id()] -= 1; 
(cuda-gdb) p tid
$1 = {MImpl = {MExtent = {sycl::_V1::detail::array<1> = {common_array = {65}}}, MIndex = {sycl::_V1::detail::array<1> = {common_array = {0}}}, MOffset = {
      sycl::_V1::detail::array<1> = {common_array = {0}}}}}
(cuda-gdb) cuda thread 0
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]
22                        result[tid.get_id()] -= 1; 
(cuda-gdb) p tid
$2 = {MImpl = {MExtent = {sycl::_V1::detail::array<1> = {common_array = {65}}}, MIndex = {sycl::_V1::detail::array<1> = {common_array = {0}}}, MOffset = {
      sycl::_V1::detail::array<1> = {common_array = {0}}}}}

[JackAKirk]

The CUDA and HIP backends are heavily reliant on general spirv interfaces to select appropriate backend dependant implementations of many functions including get_id() (and all sycl math functions for instance).
For this reason I'm not sure incorrect behaviour here can be easily resolved. Maybe we will have to consider whether the -fno-sycl-libspirv flag should be allowed to have an effect in these backends.

[MarkusBuettner]

I do not know much about the internals of the compiler, but I would have expected that SPIR-V is only needed during the compilation and/or the JIT compilation for certain backends.

Maybe it is possible for the time being to add a warning in the documentation/compiler, that these backends may not work properly if compiled with -fno-sycl-libspirv?

[JackAKirk]

I do not know much about the internals of the compiler, but I would have expected that SPIR-V is only needed during the compilation and/or the JIT compilation for certain backends.

Maybe it is possible for the time being to add a warning in the documentation/compiler, that these backends may not work properly if compiled with -fno-sycl-libspirv?

Hi. We have an internal ticket to investigate this further to find the best resolution.

[JackAKirk]

I do not know much about the internals of the compiler, but I would have expected that SPIR-V is only needed during the compilation and/or the JIT compilation for certain backends.

Maybe it is possible for the time being to add a warning in the documentation/compiler, that these backends may not work properly if compiled with -fno-sycl-libspirv?

See the description of this PR for a diagnosis of what is going on:

#7378

Basically there is no reason for you to use -fno-sycl-libspirv (unless you really want to omit linking libspirv libs (cuda/hip backends only) and the warnings added in #7378 now inform of this.

[rodburns]

This issue is resolved through PR #7378 so closing this issue.