[CUDA] 'PI_ERROR_INVALID_KERNEL_NAME' when dynamic library with kernel is loaded by dlopen secondly

[DongBaiYue]

Describe the bug 'PI_ERROR_INVALID_KERNEL_NAME' error when dynamic library with kernel is loaded by dlopen secondly on linux with CUDA backend.

To Reproduce Error message:

AAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAterminate called after throwing an instance of 'sycl::_V1::exception'
  what():  Native API failed. Native API returns: -46 (PI_ERROR_INVALID_KERNEL_NAME)
Aborted (core dumped)

The small reproducer:

// main.cpp
#include <sycl/sycl.hpp>
#include <dlfcn.h>

void dlexe(std::string shared_lib_path, std::string func_name){
    void * so_handler = dlopen(shared_lib_path.c_str(), RTLD_LAZY);
    void (*kernel_func)() = (void (*)())dlsym(so_handler, func_name.c_str());
    kernel_func();
    dlclose(so_handler);
}

int main(){
    dlexe("sycl_libA.so", "funcA");
    dlexe("sycl_libB.so", "funcB");
    return 0;
}

// sycl_libA.cpp
#include <CL/sycl.hpp>
using namespace sycl;

#ifdef __cplusplus
extern "C"
#endif
void funcA() {
    queue q;
    q.submit([&](sycl::handler &h) {
        sycl::stream os(1024, 768, h);
        h.parallel_for(32, [=](sycl::id<1> i) {
            os<<"A";
        });
    }).wait();
}

// sycl_libB.cpp
#include <CL/sycl.hpp>
using namespace sycl;

#ifdef __cplusplus
extern "C"
#endif
void funcB() {
    queue q;
    q.submit([&](sycl::handler &h) {
        sycl::stream os(1024, 768, h);
        h.parallel_for(32, [=](sycl::id<1> i) {
            os<<"B";
        });
    }).wait();
}

// command
clang++ -std=c++17 -O3 -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fPIC -shared sycl_libA.cpp -o sycl_libA.so
clang++ -std=c++17 -O3 -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fPIC -shared sycl_libB.cpp -o sycl_libB.so
clang++ -std=c++17 -O3 -fsycl -fsycl-targets=nvptx64-nvidia-cuda main.cpp -ldl

Environment (please complete the following information):

• OS: Linux
• Target device and vendor: Nvidia GPU
• DPC++ version: clang version 18.0.0 (https://github.com/intel/llvm dbd9b67cd5c097b41f20c8a61f37c711888f2a20)
• CUDA version:

  Cuda compilation tools, release 11.7, V11.7.99
  Build cuda_11.7.r11.7/compiler.31442593_0

Additional context if DPC++ version is 2022-12，Error message is:

AAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAA
PI CUDA ERROR:
        Value:           500
        Name:            CUDA_ERROR_NOT_FOUND
        Description:     named symbol not found
        Function:        cuda_piKernelCreate
        Source Location: /home/ly/sycl_workspace/llvm-2022-12/sycl/plugins/cuda/pi_cuda.cpp:2872

terminate called after throwing an instance of 'sycl::_V1::exception'
  what():  Native API failed. Native API returns: -999 (Unknown PI error)
Aborted (core dumped)

[DongBaiYue]

Can you reproduce this bug?

@npmiller @mdtoguchi @AerialMantis

[mdtoguchi]

This command clang++ -std=c++17 -O3 -fsycl -fsycl-targets=nvptx64-nvidia-cuda main.cpp -ldl does not have any references to the generated shared objects, is that purposeful? The executable built from main.cpp needs to have some kind of reference to the shared objects to or it won't know what to look for to resolve.

[DongBaiYue]

This command clang++ -std=c++17 -O3 -fsycl -fsycl-targets=nvptx64-nvidia-cuda main.cpp -ldl does not have any references to the generated shared objects, is that purposeful? The executable built from main.cpp needs to have some kind of reference to the shared objects to or it won't know what to look for to resolve.

dlfcn allows runtime dynamic library loading, no need to link at compile time。

• run-time loading: dlopen()
• access the symbols from run-time loaded shared library: dlsym()

[DongBaiYue]

For the above reproducer

//expected result
AAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAABBBBBBBBBBBBBBBBBBBBBBBBBBBBBBBB
//actual result
AAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAterminate called after throwing an instance of 'sycl::_V1::exception'
  what():  Native API failed. Native API returns: -46 (PI_ERROR_INVALID_KERNEL_NAME)
Aborted (core dumped)

The error occurs when funcB is executed after funcA is successfully executed.