4.2。 CUDA 内核 API

原文： http://numba.pydata.org/numba-doc/latest/cuda-reference/kernel.html

4.2.1。内核声明

@cuda.jit装饰器用于创建 CUDA 内核：

  1. numba.cuda.jit(func_or_sig=None, argtypes=None, device=False, inline=False, bind=True, link=[], debug=None, **kws)

JIT 编译符合 CUDA Python 规范的 python 函数。如果提供了签名，则返回一个函数进行编译的函数。如果

| 参数： |

  • func_or_sig (function or numba.typing.Signature) –

    JIT 编译的函数，或要编译的函数的签名。如果提供了函数，则返回AutoJitCUDAKernel。如果提供了签名，则返回一个函数，该函数接受函数编译并返回AutoJitCUDAKernel。

    注意

    内核不能有任何返回值。

  • 设备（ bool） - 表示这是否是设备功能。

  • 结合（ bool） - 立即强制结合 CUDA 环境
  • 链接（ 列表 ） - 包含 PTX 源的文件列表，用于链接功能
  • debug - 如果为 True，检查执行内核时抛出的异常。由于这会降低性能，因此应仅用于调试目的。默认为 False。 （可以通过设置环境变量NUMBA_CUDA_DEBUGINFO=1来覆盖默认值。）
  • fastmath - 如果为 true，则启用 flush-to-zero 和 fusion-multiply-add，禁用精确除法和平方根。此参数对设备功能没有影响，其 fastmath 设置取决于调用它们的内核函数。

  • max_registers - 限制内核每个线程最多使用这个数量的寄存器。有助于增加入住率。

    | | —- | —- |

  1. class numba.cuda.compiler.AutoJitCUDAKernel(func, bind, targetoptions)

CUDA 内核对象。调用时，内核对象将专门为给定的参数（如果没有合适的专用版本已经存在）＆amp;计算功能，并在与当前上下文关联的设备上启动。

内核对象不是由用户构造的，而是使用 numba.cuda.jit() 装饰器创建的。

  1. extensions

必须具有 <cite>prepare_args</cite> 函数的对象列表。当调用专用内核时，每个参数将传递给 <cite>prepare_args</cite> （从此列表中的最后一个对象到第一个对象）。 <cite>prepare_args</cite> 的参数是：

  • <cite>ty</cite> numba 类型的参数
  • <cite>val</cite> 参数值本身
  • <cite>stream</cite> 用于当前调用内核的 CUDA 流
  • <cite>retr</cite> 一个零 arg 函数列表，你可能想要将调用后的清理工作附加到。

<cite>prepare_args</cite> 函数必须返回一个元组<cite>（ty，val）</cite>，它将依次传递给下一个最右侧<cite>扩展名</cite>。在调用所有扩展之后，生成的<cite>（ty，val）</cite>将被传递到 Numba 的默认参数编组逻辑中。

  1. inspect_asm(signature=None, compute_capability=None)

返回到目前为止遇到的所有签名的生成的汇编代码，或者返回 LLVM IR 以获取特定签名和 compute_capability（如果给定）。

  1. inspect_llvm(signature=None, compute_capability=None)

返回到目前为止遇到的所有签名的 LLVM IR，或者给出特定签名和 compute_capability 的 LLVM IR。

  1. inspect_types(file=None)

生成此函数的 Python 源代码的转储，并使用相应的 Numba IR 和类型信息进行注释。如果 文件 为 无 ，转储将写入 文件 或 sys.stdout 。

  1. specialize(*args)

编译并绑定当前上下文专用于给定 args 的此内核版本。

各个专用内核是 numba.cuda.compiler.CUDAKernel 的实例：

  1. class numba.cuda.compiler.CUDAKernel(llvm_module, name, pretty_name, argtypes, call_helper, link=(), debug=False, fastmath=False, type_annotation=None, extensions=[], max_registers=None)

CUDA 内核专门用于一组给定的参数类型。调用时，此对象将验证参数类型是否与其专用的参数类型匹配，然后在设备上启动内核。

  1. bind()

强制绑定到当前的 CUDA 上下文

  1. device

获取当前活动上下文

  1. inspect_asm()

返回此内核的 PTX 代码。

  1. inspect_llvm()

返回此内核的 LLVM IR。

  1. inspect_types(file=None)

生成此函数的 Python 源代码的转储，并使用相应的 Numba IR 和类型信息进行注释。如果 文件 为 无 ，转储将写入 文件 或 sys.stdout 。

  1. ptx

该内核的 PTX 代码。

4.2.2。内在属性和函数

本节中的其余属性和函数只能在 CUDA 内核中调用。

4.2.2.1。线程索引

  1. numba.cuda.threadIdx

当前线程块中的线程索引，通过属性x，y和z访问。每个索引是一个整数，范围从 0 到 0 到 numba.cuda.blockDim 不包含的属性的相应值。

  1. numba.cuda.blockIdx

线程块网格中的块索引，通过属性x，y和z访问。每个索引是一个整数，范围从 0 到 0 到 numba.cuda.gridDim 不包含的属性的相应值。

  1. numba.cuda.blockDim

线程块的形状，在实例化内核时声明。对于给定内核中的所有线程，该值是相同的，即使它们属于不同的块（即每个块都是“满”）。

  1. numba.cuda.gridDim

块网格的形状，通过属性x，y和z访问。

  1. numba.cuda.laneid

当前 warp 中的线程索引，作为一个整数，范围从 0 到包含 numba.cuda.warpsize 不包括。

  1. numba.cuda.warpsize

GPU 上的 warp 线程的大小。目前这总是 32。

  1. numba.cuda.grid(ndim)

返回整个块网格中当前线程的绝对位置。 ndim 应该对应于实例化内核时声明的维数。如果 ndim 为 1，则返回单个整数。如果 ndim 为 2 或 3，则返回给定数量的整数的元组。

第一个整数的计算如下：

  1. cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x

并且与其他两个索引类似，但使用y和z属性。

  1. numba.cuda.gridsize(ndim)

返回整个块网格的线程中的绝对大小（或形状）。 ndim 应该对应于实例化内核时声明的维数。

第一个整数的计算如下：

  1. cuda.blockDim.x * cuda.gridDim.x

并且与其他两个索引类似，但使用y和z属性。

4.2.2.2。内存管理

  1. numba.cuda.shared.array(shape, dtype)

使用给定的shape和dtype在 CUDA 内核的本地内存空间中创建一个数组。

返回其内容未初始化的数组。

注意

同一线程块中的所有线程都看到相同的数组。

  1. numba.cuda.local.array(shape, dtype)

使用给定的shape和dtype在 CUDA 内核的本地内存空间中创建一个数组。

返回其内容未初始化的数组。

注意

每个线程都看到一个唯一的数组

  1. numba.cuda.const.array_like(ary)

在编译时将ary复制到 CUDA 内核上的常量内存空间。

返回类似ary参数的数组。

注意

所有线程和块都看到相同的数组。

4.2.2.3。同步和原子操作

  1. numba.cuda.atomic.add(array, idx, value)

执行array[idx] += value。仅支持 int32，int64，float32 和 float64。 idx参数可以是整数或整数索引的元组，用于索引到多维数组。 idx中的元素数必须与array的维数相匹配。

在存储新值之前返回array[idx]的值。表现得像原子载荷。

  1. numba.cuda.atomic.max(array, idx, value)

执行array[idx] = max(array[idx], value)。仅支持 int32，int64，float32 和 float64。 idx参数可以是整数或整数索引的元组，用于索引到多维数组。 idx中的元素数必须与array的维数相匹配。

在存储新值之前返回array[idx]的值。表现得像原子载荷。

  1. numba.cuda.syncthreads()

同步同一线程块中的所有线程。此函数实现与传统多线程编程中的障碍相同的模式：此函数等待，直到块中的所有线程调用它，此时它将控制权返回给所有调用者。

  1. numba.cuda.syncthreads_count(predicate)

numba.cuda.syncthreads 的扩展，其中返回值是predicate为真的线程数。

  1. numba.cuda.syncthreads_and(predicate)

numba.cuda.syncthreads 的扩展，如果predicate对所有线程都为真，则返回 1，否则返回 0。

  1. numba.cuda.syncthreads_or(predicate)

numba.cuda.syncthreads 的扩展，如果任何线程的predicate为真，则返回 1，否则返回 0。

警告

所有 syncthreads 函数必须由线程块中的每个线程调用。如果这样做可能会导致未定义的行为。

4.2.2.4。记忆栅栏

内存屏障用于保证内存操作的效果可由同一线程块内的其他线程，相同的 GPU 设备和相同的系统（跨全局内存的 GPU）看到。内存加载和存储保证不会通过优化传递在内存栅栏中移动。

警告

内存栅栏被认为是高级 API，大多数用户使用线程屏障（例如syncthreads()）。

  1. numba.cuda.threadfence()

设备级别的存储器围栏（在 GPU 内）。

  1. numba.cuda.threadfence_block()

线程块级别的内存栅栏。

  1. numba.cuda.threadfence_system()

系统级别的内存栅栏（跨 GPU）。

4.2.2.5。 Warp Intrinsics

所有 warp 级操作至少需要 CUDA 9.参数membermask是一个 32 位整数掩码，每个位对应于 warp 中的一个线程，1 表示该线程位于函数调用中的线程子集中。如果 GPU 计算能力低于 7.x，则membermask必须全为 1。

  1. numba.cuda.syncwarp(membermask)

在 warp 中同步屏蔽的线程子集。

  1. numba.cuda.all_sync(membermask, predicate)

如果predicate对于屏蔽 warp 中的所有线程都为 true，则返回非零值，否则返回 0。

  1. numba.cuda.any_sync(membermask, predicate)

如果predicate对于屏蔽 warp 中的任何线程为 true，则返回非零值，否则返回 0。

  1. numba.cuda.eq_sync(membermask, predicate)

如果 boolean predicate对于屏蔽 warp 中的所有线程都相同，则返回非零值，否则返回 0。

  1. numba.cuda.ballot_sync(membermask, predicate)

返回 warp 中predicate为 true 并且在给定掩码内的所有线程的掩码。

  1. numba.cuda.shfl_sync(membermask, value, src_lane)

在屏蔽的扭曲中随机播放value并从src_lane返回value。如果这在 warp 之外，则返回给定的value。

  1. numba.cuda.shfl_up_sync(membermask, value, delta)

在屏蔽的扭曲中随机播放value并从laneid - delta返回value。如果这在 warp 之外，则返回给定的value。

  1. numba.cuda.shfl_down_sync(membermask, value, delta)

在屏蔽的扭曲中随机播放value并从laneid + delta返回value。如果这在 warp 之外，则返回给定的value。

  1. numba.cuda.shfl_xor_sync(membermask, value, lane_mask)

在屏蔽的扭曲中随机播放value并从laneid ^ lane_mask返回value。

  1. numba.cuda.match_any_sync(membermask, value, lane_mask)

返回与掩码 warp 中给定value具有相同value的线程掩码。

  1. numba.cuda.match_all_sync(membermask, value, lane_mask)

返回（mask，pred）的元组，其中 mask 是掩码 warp 中与给定value具有相同value的线程的掩码，如果它们都具有相同的值，则为 0。是掩码 warp 中的所有线程是否具有相同 warp 的布尔值。

4.2.2.6。整数内在函数

可以使用 CUDA Math API 的整数内在函数的子集。有关进一步的文档，包括语义，请参阅 CUDA 工具包文档。

  1. numba.cuda.popc()

返回给定值中的设置位数。

  1. numba.cuda.brev()

反转整数值的位模式，例如 0b10110110 变为 0b01101101。

  1. numba.cuda.clz()

计算值中前导零的数量。

  1. numba.cuda.ffs()

在整数中查找设置为 1 的最低有效位的位置。

4.2.2.7。浮点内在函数

可以使用 CUDA Math API 的浮点内部函数的子集。有关进一步的文档，包括语义，请参阅 CUDA Toolkit 文档的单和双精度部分。

  1. numba.cuda.fma()

执行融合乘法 - 加法运算。以 C api 中的fma和fmaf命名，但映射到fma.rn.f32和fma.rn.f64（舍入到最近 - 偶数）PTX 指令。

4.2.2.8。控制流程说明

CUDA 控制流指令的子集可直接作为内在函数使用。避免分支是提高 CUDA 性能的关键方法，使用这些内在函数意味着您不必依赖nvcc优化器来识别和删除分支。有关进一步的文档，包括语义，请参阅相关的 CUDA 工具包文档。

  1. numba.cuda.selp()

根据第一个参数的值，在两个表达式之间进行选择。与 LLVM 的select指令类似。