# 4.2。 CUDA 内核 API
> 原文： [http://numba.pydata.org/numba-doc/latest/cuda-reference/kernel.html](http://numba.pydata.org/numba-doc/latest/cuda-reference/kernel.html)
## 4.2.1。内核声明
`@cuda.jit`装饰器用于创建 CUDA 内核：
```py
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_](https://docs.python.org/3/library/functions.html#bool "(in Python v3.7)")） - 表示这是否是设备功能。
* **结合**（ [_bool_](https://docs.python.org/3/library/functions.html#bool "(in Python v3.7)")） - 立即强制结合 CUDA 环境
* **链接**（ [_ 列表 _](https://docs.python.org/3/library/stdtypes.html#list "(in Python v3.7)") ） - 包含 PTX 源的文件列表，用于链接功能
* **debug** - 如果为 True，检查执行内核时抛出的异常。由于这会降低性能，因此应仅用于调试目的。默认为 False。 （可以通过设置环境变量`NUMBA_CUDA_DEBUGINFO=1`来覆盖默认值。）
* **fastmath** - 如果为 true，则启用 flush-to-zero 和 fusion-multiply-add，禁用精确除法和平方根。此参数对设备功能没有影响，其 fastmath 设置取决于调用它们的内核函数。
* **max_registers** - 限制内核每个线程最多使用这个数量的寄存器。有助于增加入住率。
|
| --- | --- |
```py
class numba.cuda.compiler.AutoJitCUDAKernel(func, bind, targetoptions)
```
CUDA 内核对象。调用时，内核对象将专门为给定的参数（如果没有合适的专用版本已经存在）＆amp;计算功能，并在与当前上下文关联的设备上启动。
内核对象不是由用户构造的，而是使用 [`numba.cuda.jit()`](#numba.cuda.jit "numba.cuda.jit") 装饰器创建的。
```py
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 的默认参数编组逻辑中。
```py
inspect_asm(signature=None, compute_capability=None)
```
返回到目前为止遇到的所有签名的生成的汇编代码，或者返回 LLVM IR 以获取特定签名和 compute_capability（如果给定）。
```py
inspect_llvm(signature=None, compute_capability=None)
```
返回到目前为止遇到的所有签名的 LLVM IR，或者给出特定签名和 compute_capability 的 LLVM IR。
```py
inspect_types(file=None)
```
生成此函数的 Python 源代码的转储，并使用相应的 Numba IR 和类型信息进行注释。如果 _ 文件 _ 为 _ 无 _，转储将写入 _ 文件 _ 或 _sys.stdout_ 。
```py
specialize(*args)
```
编译并绑定当前上下文专用于给定 _args_ 的此内核版本。
各个专用内核是 [`numba.cuda.compiler.CUDAKernel`](#numba.cuda.compiler.CUDAKernel "numba.cuda.compiler.CUDAKernel") 的实例：
```py
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 内核专门用于一组给定的参数类型。调用时，此对象将验证参数类型是否与其专用的参数类型匹配，然后在设备上启动内核。
```py
bind()
```
强制绑定到当前的 CUDA 上下文
```py
device
```
获取当前活动上下文
```py
inspect_asm()
```
返回此内核的 PTX 代码。
```py
inspect_llvm()
```
返回此内核的 LLVM IR。
```py
inspect_types(file=None)
```
生成此函数的 Python 源代码的转储，并使用相应的 Numba IR 和类型信息进行注释。如果 _ 文件 _ 为 _ 无 _，转储将写入 _ 文件 _ 或 _sys.stdout_ 。
```py
ptx
```
该内核的 PTX 代码。
## 4.2.2。内在属性和函数
本节中的其余属性和函数只能在 CUDA 内核中调用。
### 4.2.2.1。线程索引
```py
numba.cuda.threadIdx
```
当前线程块中的线程索引，通过属性`x`，`y`和`z`访问。每个索引是一个整数，范围从 0 到 0 到 [`numba.cuda.blockDim`](#numba.cuda.blockDim "numba.cuda.blockDim") 不包含的属性的相应值。
```py
numba.cuda.blockIdx
```
线程块网格中的块索引，通过属性`x`，`y`和`z`访问。每个索引是一个整数，范围从 0 到 0 到 [`numba.cuda.gridDim`](#numba.cuda.gridDim "numba.cuda.gridDim") 不包含的属性的相应值。
```py
numba.cuda.blockDim
```
线程块的形状，在实例化内核时声明。对于给定内核中的所有线程，该值是相同的，即使它们属于不同的块（即每个块都是“满”）。
```py
numba.cuda.gridDim
```
块网格的形状，通过属性`x`，`y`和`z`访问。
```py
numba.cuda.laneid
```
当前 warp 中的线程索引，作为一个整数，范围从 0 到包含 [`numba.cuda.warpsize`](#numba.cuda.warpsize "numba.cuda.warpsize") 不包括。
```py
numba.cuda.warpsize
```
GPU 上的 warp 线程的大小。目前这总是 32。
```py
numba.cuda.grid(ndim)
```
返回整个块网格中当前线程的绝对位置。 _ndim_ 应该对应于实例化内核时声明的维数。如果 _ndim_ 为 1，则返回单个整数。如果 _ndim_ 为 2 或 3，则返回给定数量的整数的元组。
第一个整数的计算如下：
```py
cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
```
并且与其他两个索引类似，但使用`y`和`z`属性。
```py
numba.cuda.gridsize(ndim)
```
返回整个块网格的线程中的绝对大小（或形状）。 _ndim_ 应该对应于实例化内核时声明的维数。
第一个整数的计算如下：
```py
cuda.blockDim.x * cuda.gridDim.x
```
并且与其他两个索引类似，但使用`y`和`z`属性。
### 4.2.2.2。内存管理
```py
numba.cuda.shared.array(shape, dtype)
```
使用给定的`shape`和`dtype`在 CUDA 内核的本地内存空间中创建一个数组。
返回其内容未初始化的数组。
注意
同一线程块中的所有线程都看到相同的数组。
```py
numba.cuda.local.array(shape, dtype)
```
使用给定的`shape`和`dtype`在 CUDA 内核的本地内存空间中创建一个数组。
返回其内容未初始化的数组。
注意
每个线程都看到一个唯一的数组
```py
numba.cuda.const.array_like(ary)
```
在编译时将`ary`复制到 CUDA 内核上的常量内存空间。
返回类似`ary`参数的数组。
注意
所有线程和块都看到相同的数组。
### 4.2.2.3。同步和原子操作
```py
numba.cuda.atomic.add(array, idx, value)
```
执行`array[idx] += value`。仅支持 int32，int64，float32 和 float64。 `idx`参数可以是整数或整数索引的元组，用于索引到多维数组。 `idx`中的元素数必须与`array`的维数相匹配。
在存储新值之前返回`array[idx]`的值。表现得像原子载荷。
```py
numba.cuda.atomic.max(array, idx, value)
```
执行`array[idx] = max(array[idx], value)`。仅支持 int32，int64，float32 和 float64。 `idx`参数可以是整数或整数索引的元组，用于索引到多维数组。 `idx`中的元素数必须与`array`的维数相匹配。
在存储新值之前返回`array[idx]`的值。表现得像原子载荷。
```py
numba.cuda.syncthreads()
```
同步同一线程块中的所有线程。此函数实现与传统多线程编程中的障碍相同的模式：此函数等待，直到块中的所有线程调用它，此时它将控制权返回给所有调用者。
```py
numba.cuda.syncthreads_count(predicate)
```
[`numba.cuda.syncthreads`](#numba.cuda.syncthreads "numba.cuda.syncthreads") 的扩展，其中返回值是`predicate`为真的线程数。
```py
numba.cuda.syncthreads_and(predicate)
```
[`numba.cuda.syncthreads`](#numba.cuda.syncthreads "numba.cuda.syncthreads") 的扩展，如果`predicate`对所有线程都为真，则返回 1，否则返回 0。
```py
numba.cuda.syncthreads_or(predicate)
```
[`numba.cuda.syncthreads`](#numba.cuda.syncthreads "numba.cuda.syncthreads") 的扩展，如果任何线程的`predicate`为真，则返回 1，否则返回 0。
警告
所有 syncthreads 函数必须由线程块中的每个线程调用。如果这样做可能会导致未定义的行为。
### 4.2.2.4。记忆栅栏
内存屏障用于保证内存操作的效果可由同一线程块内的其他线程，相同的 GPU 设备和相同的系统（跨全局内存的 GPU）看到。内存加载和存储保证不会通过优化传递在内存栅栏中移动。
警告
内存栅栏被认为是高级 API，大多数用户使用线程屏障（例如`syncthreads()`）。
```py
numba.cuda.threadfence()
```
设备级别的存储器围栏（在 GPU 内）。
```py
numba.cuda.threadfence_block()
```
线程块级别的内存栅栏。
```py
numba.cuda.threadfence_system()
```
系统级别的内存栅栏（跨 GPU）。
### 4.2.2.5。 Warp Intrinsics
所有 warp 级操作至少需要 CUDA 9.参数`membermask`是一个 32 位整数掩码，每个位对应于 warp 中的一个线程，1 表示该线程位于函数调用中的线程子集中。如果 GPU 计算能力低于 7.x，则`membermask`必须全为 1。
```py
numba.cuda.syncwarp(membermask)
```
在 warp 中同步屏蔽的线程子集。
```py
numba.cuda.all_sync(membermask, predicate)
```
如果`predicate`对于屏蔽 warp 中的所有线程都为 true，则返回非零值，否则返回 0。
```py
numba.cuda.any_sync(membermask, predicate)
```
如果`predicate`对于屏蔽 warp 中的任何线程为 true，则返回非零值，否则返回 0。
```py
numba.cuda.eq_sync(membermask, predicate)
```
如果 boolean `predicate`对于屏蔽 warp 中的所有线程都相同，则返回非零值，否则返回 0。
```py
numba.cuda.ballot_sync(membermask, predicate)
```
返回 warp 中`predicate`为 true 并且在给定掩码内的所有线程的掩码。
```py
numba.cuda.shfl_sync(membermask, value, src_lane)
```
在屏蔽的扭曲中随机播放`value`并从`src_lane`返回`value`。如果这在 warp 之外，则返回给定的`value`。
```py
numba.cuda.shfl_up_sync(membermask, value, delta)
```
在屏蔽的扭曲中随机播放`value`并从`laneid - delta`返回`value`。如果这在 warp 之外，则返回给定的`value`。
```py
numba.cuda.shfl_down_sync(membermask, value, delta)
```
在屏蔽的扭曲中随机播放`value`并从`laneid + delta`返回`value`。如果这在 warp 之外，则返回给定的`value`。
```py
numba.cuda.shfl_xor_sync(membermask, value, lane_mask)
```
在屏蔽的扭曲中随机播放`value`并从`laneid ^ lane_mask`返回`value`。
```py
numba.cuda.match_any_sync(membermask, value, lane_mask)
```
返回与掩码 warp 中给定`value`具有相同`value`的线程掩码。
```py
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 工具包文档](https://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__INTRINSIC__INT.html)。
```py
numba.cuda.popc()
```
返回给定值中的设置位数。
```py
numba.cuda.brev()
```
反转整数值的位模式，例如 0b10110110 变为 0b01101101。
```py
numba.cuda.clz()
```
计算值中前导零的数量。
```py
numba.cuda.ffs()
```
在整数中查找设置为 1 的最低有效位的位置。
### 4.2.2.7。浮点内在函数
可以使用 CUDA Math API 的浮点内部函数的子集。有关进一步的文档，包括语义，请参阅 CUDA Toolkit 文档的[单](https://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__SINGLE.html)和[双](https://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__DOUBLE.html)精度部分。
```py
numba.cuda.fma()
```
执行融合乘法 - 加法运算。以 C api 中的`fma`和`fmaf`命名，但映射到`fma.rn.f32`和`fma.rn.f64`（舍入到最近 - 偶数）PTX 指令。
### 4.2.2.8。控制流程说明
CUDA 控制流指令的子集可直接作为内在函数使用。避免分支是提高 CUDA 性能的关键方法，使用这些内在函数意味着您不必依赖`nvcc`优化器来识别和删除分支。有关进一步的文档，包括语义，请参阅[相关的 CUDA 工具包文档](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#comparison-and-selection-instructions)。
```py
numba.cuda.selp()
```
根据第一个参数的值，在两个表达式之间进行选择。与 LLVM 的`select`指令类似。
- 1. 用户手册
- 1.1。 Numba 的约 5 分钟指南
- 1.2。概述
- 1.3。安装
- 1.4。使用@jit 编译 Python 代码
- 1.5。使用@generated_jit 进行灵活的专业化
- 1.6。创建 Numpy 通用函数
- 1.7。用@jitclass 编译 python 类
- 1.8。使用@cfunc 创建 C 回调
- 1.9。提前编译代码
- 1.10。使用@jit 自动并行化
- 1.11。使用@stencil装饰器
- 1.12。从 JIT 代码 中回调到 Python 解释器
- 1.13。性能提示
- 1.14。线程层
- 1.15。故障排除和提示
- 1.16。常见问题
- 1.17。示例
- 1.18。会谈和教程
- 2. 参考手册
- 2.1。类型和签名
- 2.2。即时编译
- 2.3。提前编译
- 2.4。公用事业
- 2.5。环境变量
- 2.6。支持的 Python 功能
- 2.7。支持的 NumPy 功能
- 2.8。与 Python 语义的偏差
- 2.9。浮点陷阱
- 2.10。 Python 2.7 寿命终止计划
- 3. 用于 CUDA GPU 的 Numba
- 3.1。概述
- 3.2。编写 CUDA 内核
- 3.3。内存管理
- 3.4。编写设备功能
- 3.5。 CUDA Python 中支持的 Python 功能
- 3.6。支持的原子操作
- 3.7。随机数生成
- 3.8。设备管理
- 3.10。示例
- 3.11。使用 CUDA 模拟器 调试 CUDA Python
- 3.12。 GPU 减少
- 3.13。 CUDA Ufuncs 和广义 Ufuncs
- 3.14。共享 CUDA 内存
- 3.15。 CUDA 阵列接口
- 3.16。 CUDA 常见问题
- 4. CUDA Python 参考
- 4.1。 CUDA 主机 API
- 4.2。 CUDA 内核 API
- 4.3。内存管理
- 5. 用于 AMD ROC GPU 的 Numba
- 5.1。概述
- 5.2。编写 HSA 内核
- 5.3。内存管理
- 5.4。编写设备功能
- 5.5。支持的原子操作
- 5.6。代理商
- 5.7。 ROC Ufuncs 和广义 Ufuncs
- 5.8。示例
- 6. 扩展 Numba
- 6.1。高级扩展 API
- 6.2。低级扩展 API
- 6.3。示例：间隔类型
- 7. 开发者手册
- 7.1。贡献给 Numba
- 7.2。 Numba 建筑
- 7.3。多态调度
- 7.4。关于发电机的注意事项
- 7.5。关于 Numba Runtime 的注意事项
- 7.6。使用 Numba Rewrite Pass 获得乐趣和优化
- 7.7。实时变量分析
- 7.8。上市
- 7.9。模板注释
- 7.10。关于自定义管道的注意事项
- 7.11。环境对象
- 7.12。哈希 的注意事项
- 7.13。 Numba 项目路线图
- 8. Numba 增强建议
- 9. 术语表