cupy.RawModule

class cupy.RawModule(unicode code=None, *, unicode path=None, tuple options=(), unicode backend=u'nvrtc', bool translate_cucomplex=False, bool enable_cooperative_groups=False)

User-defined custom module.

This class can be used to either compile raw CUDA sources or load CUDA modules (*.cubin, *.ptx). This class is useful when a number of CUDA kernels in the same source need to be retrieved.

For the former case, the CUDA source code is compiled when initializing a new instance of this class, and the kernels can be retrieved by calling get_function(), which will return an instance of RawKernel. (Same as in RawKernel, the generated binary is also cached.)

For the latter case, an existing CUDA binary (*.cubin) or a PTX file can be loaded by providing its path, and kernels therein can be retrieved similarly.

Parameters
  • code (str) – CUDA source code. Mutually exclusive with path.

  • path (str) – Path to cubin/ptx. Mutually exclusive with code.

  • options (tuple of str) – Compiler options passed to the backend (NVRTC or NVCC). For details, see https://docs.nvidia.com/cuda/nvrtc/index.html#group__options or https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#command-option-description

  • backend (str) – Either nvrtc or nvcc. Defaults to nvrtc

  • translate_cucomplex (bool) – Whether the CUDA source includes the header cuComplex.h or not. If set to True, any code that uses the functions from cuComplex.h will be translated to its Thrust counterpart. Defaults to False.

  • enable_cooperative_groups (bool) – Whether to enable cooperative groups in the CUDA source. If set to True, compile options are configured properly and the kernel is launched with cuLaunchCooperativeKernel so that cooperative groups can be used from the CUDA source. This feature is only supported in CUDA 9 or later.

Note

Each kernel in RawModule possesses independent function attributes.

Methods

get_function(self, unicode name)

Retrieve a CUDA kernel by its name from the module.

Parameters

name (str) – Name of the kernel function.

Returns

An RawKernel instance.

Return type

RawKernel

get_global(self, name)

Retrieve a pointer to a global symbol by its name from the module.

Parameters

name (str) – Name of the global symbol.

Returns

A handle to the global symbol.

Return type

MemoryPointer

Note

This method can be used to access, for example, constant memory:

# to get a pointer to "arr" declared in the source like this:
# __constant__ float arr[10];
memptr = mod.get_global("arr")
# ...wrap it using cupy.ndarray with a known shape
arr_ndarray = cp.ndarray((10,), cp.float32, memptr)
# ...perform data transfer to initialize it
arr_ndarray[...] = cp.random.random((10,), dtype=cp.float32)
# ...and arr is ready to be accessed by RawKernels
get_texref(self, name)

Retrieve a texture reference by its name from the module.

Parameters

name (str) – Name of the texture reference.

Returns

A CUtexref handle, to be passed to TextureReference.

Return type

intptr_t

Attributes

backend
code
enable_cooperative_groups
file_path
module
options