Does DPC++ expose driver level api?

[horrorChen]

Hello, I am trying to migrate a program from cuda to sycl, while it takes advantage of some cuda driver apis like cuModuleGetFunction. It seems that sycl only exposes runtime level api and I wonder if sycl provides similar api to read a function object from a external file.

I have tried Dynamic linking of device code, but the feature only works on intel devices. Also, the feature provides the ability to call a external function rather than read a function object.

I also tried to use internal function of kernel and kernel_bundle classes in DPC++. There does exist a get function which returns a OpenCL kernel handle, but it calls for the definition of __SYCL_INTERNAL_API macro. If I want to use it, I need to simulate my program as intraware of DPC++ and it might cause other problem.

Thanks.

[AerialMantis]

Hi @horrorChen, thanks for the question. So the SYCL API doesn't map specifically to the CUDA runtime or driver APIs, it can be implemented on top of either, for example hipSYCL uses the runtime API whereas DPC++ uses the driver API. Saying that you are correct there is not a direct mapping for the cuModuleGetFunction function in SYCL, the equivalent of a CUmodule in SYCL would be a kernel_bundle (https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_the_kernel_bundle_class) and the equivalent of cuModuleGetFunction in SYCL would be kernel_bundle::get_kernel_ids which returns a kernel_id (https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_kernel_identifiers). The main difference being that I believe cuModuleGetFunction will also return non-kernel functions, whereas get_kernel_ids will not.

Another option could be CUDA backend interoperability, where you could retrieve a CUmodule from a kernel_bundle, however, this is still work in progress (KhronosGroup/SYCL-Docs#197, #6162).

If what you're looking to do is call a CUDA function in a different translation unit from a SYCL kernel, this is possible in DPC++ with the CUDA backend. The way to do this is mark the __device__ qualified CUDA function with SYCL_EXTERNAL and make sure there is a declaration of this function in the translation unit it's being called from with SYCL_EXTERNAL. There's an example of this in the SYCL CTS; https://github.com/KhronosGroup/SYCL-CTS/blob/SYCL-2020/tests/cuda_interop/cuda_interop_device_function.cu, this has the CUDA device function and the SYCL kernel in the same translation unit, but it works much the same.

I hope this help.

[horrorChen]

Thanks for your reply @AerialMantis! Actually, the key code I want to substitute is below:

string module_file = path + "/" + kernel_name + ".cubin";
CUDA_DRIVER_CALL(cuModuleLoad(&module, module_file.c_str()));
CUDA_DRIVER_CALL(cuModuleGetFunction(&function, module, kernel_name.c_str()));
return function;

In this case, the first and third method you mentioned might not work in my program, since a kernel_id do not offer a get function, not to mention that in DPC++ SYCL_EXTERNAL only work on Intel device (Test plan for dynamic linking of device code).

I will pay attention to the progress of CUDA backend interoperability, and I am glad to find any else methods to resolve the question.

Thanks.