# 3.2。编写 CUDA 内核
> 原文： [http://numba.pydata.org/numba-doc/latest/cuda/kernels.html](http://numba.pydata.org/numba-doc/latest/cuda/kernels.html)
## 3.2.1。简介
与用于编程 CPU 的传统顺序模型不同，CUDA 具有执行模型。在 CUDA 中，您编写的代码将由多个线程同时执行（通常为数百或数千）。您的解决方案将通过定义 _grid_ ， _blocks_ 和 _threads_ 的线程层次结构来建模。
Numba 的 CUDA 支持公开了用于声明和管理这种线程层次结构的工具。这些设施与 NVidia 的 CUDA C 语言大致相似。
Numba 还暴露了三种 GPU 内存：全局[设备内存](memory.html#cuda-device-memory)（连接到 GPU 本身的大型，相对较慢的片外内存），片上[共享内存](memory.html#cuda-shared-memory)和 []本地记忆](memory.html#cuda-local-memory)。除了最简单的算法外，您必须仔细考虑如何使用和访问内存，以最大限度地减少带宽需求和争用。
## 3.2.2。内核声明
_ 内核函数 _ 是一个 GPU 函数，用于从 CPU 代码（*）调用。它赋予它两个基本特征：
* 内核无法显式返回值;所有结果数据必须写入传递给函数的数组（如果计算一个标量，你可能会传递一个单元素数组）;
* 内核在被调用时显式声明它们的线程层次结构：即线程块的数量和每个块的线程数（注意，当内核被编译一次时，可以使用不同的块大小或网格大小多次调用它）。
乍一看，使用 Numba 编写 CUDA 内核看起来非常像为 CPU 编写 [JIT 函数](../glossary.html#term-jit-function)：
```py
@cuda.jit
def increment_by_one(an_array):
"""
Increment all array elements by one.
"""
# code elided here; read further for different implementations
```
（*）注意：较新的 CUDA 设备支持设备端内核启动;此功能称为 _ 动态并行 _，但 Numba 目前不支持它
## 3.2.3。内核调用
内核通常以以下方式启动：
```py
threadsperblock = 32
blockspergrid = (an_array.size + (threadsperblock - 1)) // threadsperblock
increment_by_one[blockspergrid, threadsperblock](an_array)
```
我们在这里注意两个步骤：
* 通过指定多个块（或“每个网格的块”）以及每个块的多个线程来实例化内核。两者的乘积将给出启动的线程总数。内核实例化是通过编译内核函数（此处为`increment_by_one`）并使用整数元组对其进行索引来完成的。
* 运行内核，通过传递输入数组（以及任何必要的单独输出数组）。默认情况下，运行内核是同步的：当内核完成执行并且数据被同步回来时，函数返回。
### 3.2.3.1。选择块大小
在声明内核所需的线程数时，拥有两级层次结构似乎很奇怪。块大小（即每个块的线程数）通常至关重要：
* 在软件方面，块大小决定共享[共享内存](memory.html#cuda-shared-memory)的给定区域的线程数。
* 在硬件方面，块大小必须足够大才能完全占用执行单元;建议可在 [CUDA C 编程指南](http://docs.nvidia.com/cuda/cuda-c-programming-guide)中找到。
### 3.2.3.2。多维块和网格
为了帮助处理多维数组，CUDA 允许您指定多维块和网格。在上面的示例中，您可以使`blockspergrid`和`threadsperblock`元组为一个，两个或三个整数。与等效大小的 1D 声明相比，这不会改变生成代码的效率或行为，但可以帮助您以更自然的方式编写算法。
## 3.2.4。螺纹定位
运行内核时，每个线程执行一次内核函数的代码。因此，它必须知道它所在的线程，以便知道它负责哪个数组元素（复杂的算法可能定义更复杂的责任，但基本原理是相同的）。
一种方法是让线程定位它在网格和块中的位置，并手动计算在数组中对应的位置：
```py
@cuda.jit
def increment_by_one(an_array):
# Thread id in a 1D block
tx = cuda.threadIdx.x
# Block id in a 1D grid
ty = cuda.blockIdx.x
# Block width, i.e. number of threads per block
bw = cuda.blockDim.x
# Compute flattened index inside the array
pos = tx + ty * bw
if pos < an_array.size: # Check array boundaries
an_array[pos] += 1
```
注意
除非您确定块大小和网格大小是数组大小的除数，否则**必须**检查边界，如上所示。
[`threadIdx`](../cuda-reference/kernel.html#numba.cuda.threadIdx "numba.cuda.threadIdx") ， [`blockIdx`](../cuda-reference/kernel.html#numba.cuda.blockIdx "numba.cuda.blockIdx") ， [`blockDim`](../cuda-reference/kernel.html#numba.cuda.blockDim "numba.cuda.blockDim") 和 [`gridDim`](../cuda-reference/kernel.html#numba.cuda.gridDim "numba.cuda.gridDim") 是 CUDA 后端为鞋底提供的特殊对象了解线程层次结构的几何以及当前线程在该几何中的位置的目的。
这些对象可以是 1D，2D 或 3D，具体取决于内核[调用的方式](#cuda-kernel-invocation)。要访问每个维度的值，请分别使用这些对象的`x`，`y`和`z`属性。
```py
numba.cuda.threadIdx
```
当前线程块中的线程索引。对于 1D 块，索引（由`x`属性给出）是一个整数，范围从 0 到包括 [`numba.cuda.blockDim`](../cuda-reference/kernel.html#numba.cuda.blockDim "numba.cuda.blockDim") 不包括。当使用多个维度时，每个维度都存在类似的规则。
```py
numba.cuda.blockDim
```
线程块的形状，在实例化内核时声明。对于给定内核中的所有线程，该值是相同的，即使它们属于不同的块（即每个块都是“满”）。
```py
numba.cuda.blockIdx
```
线程网格中的块索引启动了一个内核。对于 1D 网格，索引（由`x`属性给出）是一个整数，范围从 0 到包括 [`numba.cuda.gridDim`](../cuda-reference/kernel.html#numba.cuda.gridDim "numba.cuda.gridDim") 不包括。当使用多个维度时，每个维度都存在类似的规则。
```py
numba.cuda.gridDim
```
块网格的形状，即由内核调用启动的块的总数，在实例化内核时声明。
### 3.2.4.1。绝对位置
简单的算法倾向于始终以与上面示例中所示相同的方式使用线程索引。 Numba 提供额外的设施来自动进行这样的计算：
```py
numba.cuda.grid(ndim)
```
返回整个块网格中当前线程的绝对位置。 _ndim_ 应该对应于实例化内核时声明的维数。如果 _ndim_ 为 1，则返回单个整数。如果 _ndim_ 为 2 或 3，则返回给定数量的整数的元组。
```py
numba.cuda.gridsize(ndim)
```
返回整个块网格的线程中的绝对大小（或形状）。 _ndim_ 具有与上述 [`grid()`](../cuda-reference/kernel.html#numba.cuda.grid "numba.cuda.grid") 相同的含义。
使用这些函数，增量示例可以变为：
```py
@cuda.jit
def increment_by_one(an_array):
pos = cuda.grid(1)
if pos < an_array.size:
an_array[pos] += 1
```
2D 阵列和线程网格的相同示例是：
```py
@cuda.jit
def increment_a_2D_array(an_array):
x, y = cuda.grid(2)
if x < an_array.shape[0] and y < an_array.shape[1]:
an_array[x, y] += 1
```
请注意，实例化内核时的网格计算仍必须手动完成，例如：
```py
from __future__ import division # for Python 2
threadsperblock = (16, 16)
blockspergrid_x = math.ceil(an_array.shape[0] / threadsperblock[0])
blockspergrid_y = math.ceil(an_array.shape[1] / threadsperblock[1])
blockspergrid = (blockspergrid_x, blockspergrid_y)
increment_a_2D_array[blockspergrid, threadsperblock](an_array)
```
### 3.2.4.2。进一步阅读
有关 CUDA 编程的详细讨论，请参见 [CUDA C 编程指南](http://docs.nvidia.com/cuda/cuda-c-programming-guide)。
- 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. 术语表