cupy.RawModule using name_expressions and nvcc and/or path

510 views Asked by At

I am using CuPy for testing cuda kernels from a library. More specifically I use the cupy.RawModule to exploit the kernels in python. However, the kernels are templated and enclosed in a namespace. Before the name_expressions parameter to RawModule in CuPy 8.0.0, I had to copy the c++-mangled names into the get_function() method manually of the RawModule. Using name_expressions I thought that this should be possible, nevertheless, this requires the code to be compiled from source using the code parameter in combination with backend='nvrtc'.

Should it be possible to enable (any of the below)?:

  1. name_expressions in conjunction with path
  2. name_expressions in conjunction with backend='nvcc'
1

There are 1 answers

13
Leo Fang On BEST ANSWER

Should it be possible to enable (any of the below)?:

  1. 'name_expressions' in conjunction with 'path'
  2. 'name_expressions' in conjunction with 'backend'='nvcc'

The answer is no for both questions.

The name_expressions feature requires the source code for just-in-time (JIT) compilation of your C++ template kernels using NVRTC, whereas the path argument is for loading external cubin, fatbin, or ptx code. If you want to compile an external source code, you can do so by loading it in Python first, and then pass it as the code argument:

with open('my_cuda_cpp_code.cu') as f:
    code = f.read()

mod = cp.RawModule(code=code, name_expressions=(...), ...)

Unfortunately unlike NVRTC, NVCC does not provide an API to return mangled names, so using NVCC is not possible. If you pass backend='nvcc' to RawModule, it'd raise an error.