cupy.RawKernel#

class cupy.RawKernel(unicode code, unicode name, tuple options=(), unicode backend=u'nvrtc', bool translate_cucomplex=False, *, bool enable_cooperative_groups=False, bool jitify=False)[source]#

用户自定义核函数。

这个类可以用来使用原始 CUDA 源代码定义自定义核函数。

核函数在调用 __call__() 方法时编译，并为每个设备缓存。编译后的二进制文件也缓存在 $HOME/.cupy/kernel_cache/ 目录下的一个哈希文件名的文件中。缓存的二进制文件可以被其他进程重用。

参数:

  • code (str) – CUDA 源代码。

  • name (str) – 核函数的名称。

  • options (str 的元组) – 传递给后端（NVRTC 或 NVCC）的编译器选项。详情请参阅 https://docs.nvda.net.cn/cuda/nvrtc/index.html#group__options 或 https://docs.nvda.net.cn/cuda/cuda-compiler-driver-nvcc/index.html#command-option-description

  • backend (str) – nvrtc 或 nvcc。默认为 nvrtc

  • translate_cucomplex (bool) – CUDA 源代码是否包含头文件 cuComplex.h。如果设置为 True，任何使用 cuComplex.h 中函数的代码都将转换为其对应的 Thrust 函数。默认为 False。

  • enable_cooperative_groups (bool) – 是否在 CUDA 源代码中启用协作组 (cooperative groups)。如果设置为 True，则会正确配置编译选项，并使用 cuLaunchCooperativeKernel 启动核函数，以便可以在 CUDA 源代码中使用协作组。此功能仅在 CUDA 9 或更高版本中受支持。

  • jitify (bool) – 是否使用 Jitify 来帮助 NVRTC 编译 C++ 核函数。默认为 False。

注意

从 CuPy v13.0.0 开始，如果未在 options 中指定，RawKernel 默认使用 C++11 标准 (-std=c++11) 进行编译。

方法

__call__(self, grid, block, args, *, shared_mem=0)#

编译并调用核函数。

仅在核函数未缓存时运行编译。

参数:

  • grid (元组) – 块网格的大小。

  • block (元组) – 每个线程块的维度。

  • args (元组) – 核函数的参数。

  • shared_mem (int) – 每个线程块的动态共享内存大小（字节）。

compile(self, log_stream=None)#

编译当前核函数。

通常，您不需要调用此方法；核函数在首次调用时会自动编译。

参数:

log_stream (object) – 传递 sys.stdout 或一个文件对象，编译器输出将写入其中。默认为 None。

__eq__(value, /)#

返回 self==value。

__ne__(value, /)#

返回 self!=value。

__lt__(value, /)#

返回 self<value。

__le__(value, /)#

返回 self<=value。

__gt__(value, /)#

返回 self>value。

__ge__(value, /)#

返回 self>=value。

属性

attributes#

返回一个包含运行时核函数属性的字典。这是一个只读属性；要覆盖属性，请使用

kernel = RawKernel(...)  # arguments omitted
kernel.max_dynamic_shared_size_bytes = ...
kernel.preferred_shared_memory_carveout = ...

请注意，上面示例中显示的两个属性是当前 CUDA 中唯一可设置的两个属性。

任何在当前 CUDA 工具包版本中不存在的属性将具有值 -1。

返回:

一个包含核函数属性的字典。

返回类型:

dict

backend#
binary_version#

编译期间使用的二进制架构版本，格式为：10*主版本号 + 次版本号。

cache_mode_ca#

指示编译期间是否设置了选项 “-Xptxas –dlcm=ca”。

code#
const_size_bytes#

函数使用的常量内存大小（字节）。

enable_cooperative_groups#
file_path#
kernel#
local_size_bytes#

函数使用的局部内存大小（字节）。

max_dynamic_shared_size_bytes#

函数可使用的最大动态分配共享内存大小（字节）。可设置。

max_threads_per_block#

可以在设备上成功启动函数的每个块的最大线程数。

name#
num_regs#

函数使用的寄存器数量。

options#
preferred_shared_memory_carveout#

在具有统一 L1 缓存和共享内存的设备上，指示用作共享内存占总内存的百分比。如果该比例不完全等于支持的共享内存容量，则使用下一个更大的受支持容量。可设置。

ptx_version#

编译期间使用的 PTX 虚拟架构版本，格式为：10*主版本号 + 次版本号。

shared_size_bytes#

函数使用的静态分配共享内存的大小（字节）。这与任何动态分配的共享内存是分开的，后者必须在调用函数时指定。