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)[源代码]#

用户定义的自定义内核。

此类可用于使用原始 CUDA 源代码定义自定义内核。

内核在调用 __call__() 方法时编译，每个设备都会缓存该编译结果。编译后的二进制文件也会缓存到 $HOME/.cupy/kernel_cache/ 目录下的一个带有哈希文件名的文件中。缓存的二进制文件可以被其他进程重用。

参数:

  • code (str) – CUDA 源代码。

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

  • options (tuple of str) – 传递给后端（NVRTC 或 NVCC）的编译器选项。详情请参阅 https://docs.nvidia.com/cuda/nvrtc/index.html#group__options 或 https://docs.nvidia.com/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源代码中启用协作组。如果设置为``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 (tuple) – 网格的块大小。

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

  • args (tuple) – 内核的参数。

  • 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#

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