'named symbol not found' error at runtime in cuda_piKernelCreate if static library with kernel is linked

[ivorobts]

A simple example showing the usage of static library with kernel which fails at runtime with CUDA BE:
Here is the content of file sycl_lib.cpp:

#include<CL/sycl.hpp>
using namespace sycl;
void func() {
  queue q;
  q.submit([&](sycl::handler &h) {
      sycl::stream os(1024, 768, h);
      h.parallel_for(32, [=](sycl::id<1> i) {
          os<<i<<"\n";
        });
    });
}

and a simple main code:

void func();
int main() {
  func();
  return 0;
}

Compilation of sycl_lib.cpp, static library creation and linking works fine:

clang++  -fsycl -fsycl-targets=nvptx64-nvidia-cuda-sycldevice -fsycl-unnamed-lambda -c sycl_lib.cpp
ar rvs sycl_lib.a sycl_lib.o
clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda-sycldevice main.cpp -o test-lib sycl_lib.a
SYCL_BE=PI_CUDA ./test-lib

The application crashes at runtime with following error:

PI CUDA ERROR:
        Value:           500
        Name:            CUDA_ERROR_NOT_FOUND
        Description:     named symbol not found
        Function:        cuda_piKernelCreate
        Source Location: /llvm/sycl/plugins/cuda/pi_cuda.cpp:2380

terminate called after throwing an instance of 'cl::sycl::runtime_error'
  what():  Native API failed. Native API returns: -999 (Unknown OpenCL error code) -999 (Unknown OpenCL error code)

Note that ahead of time compilation and linking works fine for CPU or Intel GPU devices. This looks like a limitation of CUDA BE.

[Michoumichmich]

Hello, I get also the same issue on the CUDA be, but this is not a limitation of the CUDA back-end. Could you try: https://github.com/Michoumichmich/llvm/commit/e3eb050b73bff24cb8335f84fa6a51f1ebaf05c5.diff, this is barely a temporary workaround to a bug that was introduced by f7ce532 in a pulldown

[ravil-mobile]

Thanks a lot. It saved my day.

[Michoumichmich]

Thanks a lot. It saved my day.

You're welcome!

[danchitnis]

It seems to happen on some versions but not others. It happens both on WSL and native Linux, I assumed it was related to GPU driver mismatch with CUDA. perhaps can you make the error info more useful with debugging traces?

[Michoumichmich]

It seems to happen on some versions but not others. It happens both on WSL and native Linux, I assumed it was related to GPU driver mismatch with CUDA. perhaps can you make the error info more useful with debugging traces?

The error is coming from the fact that the names of the generated device object files were changed, but clang-offload-bundler is still called on the old names it fails silently. If you compile with -v you won't even see ptxas compiling down the kernels during the final linking of the .so object or executable. It might seem to work if you did not clean your build directory, as you're probably using old files.
If your bug depends on the version of CUDA, it might not be related to that exact one

[danchitnis]

So far it is working for me, but I am not using the latest commit (a month old). I am aware that there is a recent bug with nvidia containers influencing WSL (and native Linux), so then these two bugs may not be related.

[Michoumichmich]

Hello @ivorobts can you try compiling with nvptx64-nvidia-cuda--sm_50 on the current branch to see if it's fixed ?

[npmiller]

This was fixed recently by this PR:

• [Driver] Properly set offload-deps target triple #4305

If you want to try it note that your SYCL code is missing a wait which might cause issues at the moment, something like this works fine though:

#include<CL/sycl.hpp>
using namespace sycl;
void func() {
  queue q;
  q.submit([&](sycl::handler &h) {
      sycl::stream os(1024, 768, h);
      h.parallel_for(32, [=](sycl::id<1> i) {
          os<<i<<"\n";
        });
    }).wait_and_throw();
}

I'm not entirely sure if it should work without the wait but either way that would be a separate issue to investigate, the original problem reported here should be fixed.

[npmiller]

Just as a side note, I double checked with the SYCL 2020 specification and the wait is indeed required in this case, the queue destructor doesn't not implicitly wait on kernel completion, section 3.9.8.1:

Note that the destructors of other SYCL objects (sycl::queue, sycl::context,...) do not block. Only a
sycl::buffer, sycl::sampled_image or sycl::unsampled_image destructor might block. The rationale is that
an object without any side effect on the host does not need to block on destruction as it would impact the
performance. So it is up to the programmer to use a member function to wait for completion in some
cases if this does not fit the goal.

[AerialMantis]

We believe this issue to be resolved now, so we are closing the ticket, @ivorobts if you continue to see any issue here please feel free to comment and we can re-open the ticket.