class cupy.RawModule(code=None, *, path=None, options=(), backend=u'nvrtc', translate_cucomplex=False, 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.

  • 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 or
  • 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.


Each kernel in RawModule possesses independent function attributes.


get_function(self, 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_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