4.2。 CUDA 内核 API

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

4.2.1。内核声明

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

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 - 限制内核每个线程最多使用这个数量的寄存器。有助于增加入住率。

| | --- | --- |

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

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

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

extensions

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

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

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

inspect_asm(signature=None, compute_capability=None)

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

inspect_llvm(signature=None, compute_capability=None)

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

inspect_types(file=None)

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

specialize(*args)

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

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

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

bind()

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

device

获取当前活动上下文

inspect_asm()

返回此内核的 PTX 代码。

inspect_llvm()

返回此内核的 LLVM IR。

inspect_types(file=None)

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

ptx

该内核的 PTX 代码。

4.2.2。内在属性和函数

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

4.2.2.1。线程索引

numba.cuda.threadIdx

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

numba.cuda.blockIdx

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

numba.cuda.blockDim

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

numba.cuda.gridDim

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

numba.cuda.laneid

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

numba.cuda.warpsize

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

numba.cuda.grid(ndim)

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

第一个整数的计算如下：

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

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

numba.cuda.gridsize(ndim)

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

第一个整数的计算如下：

cuda.blockDim.x * cuda.gridDim.x

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

4.2.2.2。内存管理

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

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

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

注意

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

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

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

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

注意

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

numba.cuda.const.array_like(ary)

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

返回类似ary参数的数组。

注意

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

4.2.2.3。同步和原子操作

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

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

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

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

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

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

numba.cuda.syncthreads()

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

numba.cuda.syncthreads_count(predicate)

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

numba.cuda.syncthreads_and(predicate)

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

numba.cuda.syncthreads_or(predicate)

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

警告

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

4.2.2.4。记忆栅栏

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

警告

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

numba.cuda.threadfence()

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

numba.cuda.threadfence_block()

线程块级别的内存栅栏。

numba.cuda.threadfence_system()

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

4.2.2.5。 Warp Intrinsics

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

numba.cuda.syncwarp(membermask)

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

numba.cuda.all_sync(membermask, predicate)

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

numba.cuda.any_sync(membermask, predicate)

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

numba.cuda.eq_sync(membermask, predicate)

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

numba.cuda.ballot_sync(membermask, predicate)

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

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

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

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

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

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

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

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

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

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

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

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 工具包文档。

numba.cuda.popc()

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

numba.cuda.brev()

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

numba.cuda.clz()

计算值中前导零的数量。

numba.cuda.ffs()

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

4.2.2.7。浮点内在函数

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

numba.cuda.fma()

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

4.2.2.8。控制流程说明

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

numba.cuda.selp()

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

我们一直在努力

apachecn/AiLearning