内存管理

CUDA 内置目标弃用通知

Numba 内置的 CUDA 目标已弃用，后续开发已移至 NVIDIA numba-cuda 包。请参阅内置 CUDA 目标弃用和维护状态。

数据传输

尽管 Numba 可以自动将 NumPy 数组传输到设备，但它只能保守地进行，即在内核完成时总是将设备内存传输回主机。为避免只读数组的不必要传输，您可以使用以下 API 手动控制传输：

numba.cuda.device_array(shape, dtype=np.float64, strides=None, order='C', stream=0): 分配一个空的设备 ndarray。类似于 numpy.empty()。

numba.cuda.device_array_like(ary, stream=0): 使用数组信息调用 device_array()。

numba.cuda.to_device(obj, stream=0, copy=True, to=None)

分配并将 numpy ndarray 或结构化标量传输到设备。

从主机到设备复制一个 numpy 数组

ary = np.arange(10)
d_ary = cuda.to_device(ary)

将传输排队到流

stream = cuda.stream()
d_ary = cuda.to_device(ary, stream=stream)

结果的 d_ary 是一个 DeviceNDArray。

从设备到主机复制

hary = d_ary.copy_to_host()

将设备复制到主机到一个现有数组

ary = np.empty(shape=d_ary.shape, dtype=d_ary.dtype)
d_ary.copy_to_host(ary)

将传输排队到流

hary = d_ary.copy_to_host(stream=stream)

除了设备数组之外，Numba 还可以处理任何实现 CUDA 数组接口的对象。这些对象也可以通过使用以下 API 创建 GPU 缓冲区的视图，手动转换为 Numba 设备数组：

numba.cuda.as_cuda_array(obj, sync=True)

从任何实现 CUDA 数组接口的对象创建 DeviceNDArray。

创建底层 GPU 缓冲区的视图。不进行数据复制。结果的 DeviceNDArray 将从 obj 获取一个引用。

如果 sync 为 True，则导入的流（如果存在）将被同步。

numba.cuda.is_cuda_array(obj)

测试对象是否定义了 __cuda_array_interface__ 属性。

不验证接口的有效性。

设备数组

设备数组引用具有以下方法。这些方法应在主机代码中调用，而不是在 CUDA-jit 编译的函数中调用。

class numba.cuda.cudadrv.devicearray.DeviceNDArray(shape, strides, dtype, stream=0, gpu_data=None)

一种 GPU 上的数组类型

copy_to_host(ary=None, stream=0)

将 self 复制到 ary，如果 ary 为 None，则创建一个新的 Numpy ndarray。

如果给定 CUDA stream，则传输将作为给定流的一部分异步进行。否则，传输是同步的：函数在复制完成后返回。

总是返回主机数组。

示例

import numpy as np
from numba import cuda

arr = np.arange(1000)
d_arr = cuda.to_device(arr)

my_kernel[100, 100](d_arr)

result_array = d_arr.copy_to_host()

is_c_contiguous(): 如果数组是 C 连续的，则返回 true。

is_f_contiguous(): 如果数组是 Fortran 连续的，则返回 true。

ravel(order='C', stream=0): 展平一个连续数组而不改变其内容，类似于 numpy.ndarray.ravel()。如果数组不连续，则引发异常。

reshape(*newshape, **kws)

重塑数组而不改变其内容，类似于 numpy.ndarray.reshape()。示例：

d_arr = d_arr.reshape(20, 50, order='F')

注意

DeviceNDArray 定义了 CUDA 数组接口。

固定内存

numba.cuda.pinned(*arylist): 一个用于临时固定一系列主机 ndarray 的上下文管理器。

numba.cuda.pinned_array(shape, dtype=np.float64, strides=None, order='C'): 分配一个带有固定（页锁定）缓冲区的 ndarray。类似于 np.empty()。

numba.cuda.pinned_array_like(ary): 使用数组信息调用 pinned_array()。

映射内存

numba.cuda.mapped(*arylist, **kws): 一个用于临时映射一系列主机 ndarray 的上下文管理器。

numba.cuda.mapped_array(shape, dtype=np.float64, strides=None, order='C', stream=0, portable=False, wc=False)

分配一个映射 ndarray，其缓冲区固定并映射到设备。类似于 np.empty()

参数

portable – 一个布尔标志，用于允许分配的设备内存可在多个设备中使用。
wc – 一个布尔标志，用于启用写组合（writecombined）分配，这种分配主机写入速度更快，设备读取速度更快，但主机写入和设备写入速度较慢。

numba.cuda.mapped_array_like(ary, stream=0, portable=False, wc=False): 使用数组信息调用 mapped_array()。

管理内存

numba.cuda.managed_array(shape, dtype=np.float64, strides=None, order='C', stream=0, attach_global=True)

分配一个带有管理缓冲区的 np.ndarray。类似于 np.empty()。

管理内存支持 Linux / x86 和 PowerPC，在 Windows 和 Linux / AArch64 上被认为是实验性的。

参数: attach_global – 一个标志，指示是否全局附加。全局附加意味着内存可从任何设备上的任何流访问。如果为 False，则附加为主机，并且内存只能由计算能力 6.0 及更高版本的设备访问。

流

流可以传递给接受它们的函数（例如主机和设备之间的复制），并传递给内核启动配置，以便操作异步执行。

numba.cuda.stream(): 创建一个 CUDA 流，它表示设备的命令队列。

numba.cuda.default_stream(): 获取默认 CUDA 流。CUDA 语义通常是，默认流是旧版默认流或每线程默认流，具体取决于使用的 CUDA API。在 Numba 中，始终使用旧版默认流的 API，但将来可能会提供使用每线程默认流 API 的选项。

numba.cuda.legacy_default_stream(): 获取旧版默认 CUDA 流。

numba.cuda.per_thread_default_stream(): 获取每线程默认 CUDA 流。

numba.cuda.external_stream(ptr)

为在 Numba 外部分配的流创建一个 Numba 流对象。

参数: ptr (int) – 指向要在 Numba Stream 中封装的外部流的指针

CUDA 流具有以下方法：

class numba.cuda.cudadrv.driver.Stream(context, handle, finalizer, external=False)

auto_synchronize(): 一个上下文管理器，它等待此流中的所有命令执行完毕，并在退出上下文时提交任何待处理的内存传输。

synchronize(): 等待此流中的所有命令执行完毕。这将提交任何待处理的内存传输。

共享内存和线程同步

必要时，可以在设备上分配有限的共享内存以加快数据访问速度。该内存将在属于给定块的所有线程之间共享（即既可读又可写），并且比常规设备内存具有更快的访问时间。它还允许线程在给定解决方案上协作。您可以将其视为手动管理的数据缓存。

与传统的动态内存管理不同，内存只在内核执行期间分配一次。

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

在设备上分配给定 shape 和 type 的共享数组。此函数必须在设备上调用（即从内核或设备函数中）。shape 是一个整数或一个整数元组，表示数组的维度，并且必须是一个简单的常量表达式。“简单常量表达式”包括但不限于：

一个字面量（例如 10）

一个局部变量，其右侧是字面量或简单的常量表达式（例如 shape，其中 shape 在函数中早先定义为 shape = 10）

一个全局变量，在编译时在 jit 编译函数的全局变量中定义（例如 shape，其中 shape 使用任何全局范围的表达式定义）。

定义结果必须是 Python int 类型（即不是 NumPy 标量或其他标量/整数类类型）。type 是要存储在数组中的元素的 Numba 类型。返回的类数组对象可以像任何普通设备数组一样进行读写（例如通过索引）。

一个常见模式是让每个线程填充共享数组中的一个元素，然后使用 syncthreads() 等待所有线程完成。

numba.cuda.syncthreads(): 同步同一线程块中的所有线程。此函数实现了与传统多线程编程中的屏障相同的模式：此函数会等待块中的所有线程都调用它，然后将控制权返回给所有调用者。

另请参阅

矩阵乘法示例.

动态共享内存

为了在内核代码中使用动态共享内存，请声明一个大小为 0 的共享数组

@cuda.jit
def kernel_func(x):
   dyn_arr = cuda.shared.array(0, dtype=np.float32)
   ...

并在内核调用期间以字节为单位指定动态共享内存的大小

kernel_func[32, 32, 0, 128](x)

在上面的代码中，内核启动配置有 4 个参数

kernel_func[grid_dim, block_dim, stream, dyn_shared_mem_size]

注意：所有动态共享内存数组都存在别名现象，因此如果您想拥有多个动态共享数组，您需要获取数组的 不相交 视图。例如，考虑：

from numba import cuda
import numpy as np

@cuda.jit
def f():
   f32_arr = cuda.shared.array(0, dtype=np.float32)
   i32_arr = cuda.shared.array(0, dtype=np.int32)
   f32_arr[0] = 3.14
   print(f32_arr[0])
   print(i32_arr[0])

f[1, 1, 0, 4]()
cuda.synchronize()

这分配了 4 字节的共享内存（足以容纳一个 int32 或一个 float32），并声明了 int32 类型和 float32 类型的动态共享内存数组。当设置 f32_arr[0] 时，这也会设置 i32_arr[0] 的值，因为它们指向相同的内存。因此我们看到输出为：

3.140000
1078523331

因为 1078523331 是 float32 值 3.14 的位表示的 int32。

如果我们获取动态共享内存的不相交视图

from numba import cuda
import numpy as np

@cuda.jit
def f_with_view():
   f32_arr = cuda.shared.array(0, dtype=np.float32)
   i32_arr = cuda.shared.array(0, dtype=np.int32)[1:] # 1 int32 = 4 bytes
   f32_arr[0] = 3.14
   i32_arr[0] = 1
   print(f32_arr[0])
   print(i32_arr[0])

f_with_view[1, 1, 0, 8]()
cuda.synchronize()

这次我们声明了 8 字节的动态共享内存，其中前 4 字节用于 float32 值，后 4 字节用于 int32 值。现在我们可以设置 int32 和 float32 值，而它们不会发生别名冲突。

3.140000
1

局部内存

局部内存是每个线程私有的内存区域。当标量局部变量不足时，使用局部内存有助于分配一些暂存区。与传统的动态内存管理不同，内存只在内核执行期间分配一次。

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

在设备上分配给定 shape 和 type 的局部数组。shape 是一个整数或一个整数元组，表示数组的维度，并且必须是一个简单的常量表达式。“简单常量表达式”包括但不限于：

一个字面量（例如 10）

一个局部变量，其右侧是字面量或简单的常量表达式（例如 shape，其中 shape 在函数中早先定义为 shape = 10）

一个全局变量，在编译时在 jit 编译函数的全局变量中定义（例如 shape，其中 shape 使用任何全局范围的表达式定义）。

定义结果必须是 Python int 类型（即不是 NumPy 标量或其他标量/整数类类型）。type 是要存储在数组中的元素的 Numba 类型。该数组对于当前线程是私有的。返回一个类数组对象，可以像任何标准数组一样进行读写（例如通过索引）。

另请参阅

CUDA 编程指南中设备内存访问的局部内存部分。

常量内存

常量内存是一个只读、缓存且片外（off-chip）的内存区域，所有线程都可以访问，并由主机分配。在常量内存中创建数组的一种方法是通过使用：

numba.cuda.const.array_like(arr): 根据类数组对象 arr 分配并在常量内存中使数组可访问。

释放行为

本节描述 Numba 内部内存管理的释放行为。如果正在使用外部内存管理插件（参见外部内存管理 (EMM) 插件接口），则释放行为可能有所不同；您可以查阅 EMM 插件的文档以了解其释放行为。

所有 CUDA 资源的释放都是按上下文跟踪的。当设备内存的最后一个引用被删除时，底层内存被调度以进行释放。释放不会立即发生。它被添加到待处理释放的队列中。这种设计有两个好处：

资源释放 API 可能会导致设备同步；从而破坏任何异步执行。延迟释放可以避免在性能关键的代码段中产生延迟。
一些释放错误可能会导致所有剩余的释放失败。持续的释放错误可能在 CUDA 驱动程序级别引起严重错误。在某些情况下，这可能意味着 CUDA 驱动程序中的段错误。在最坏的情况下，这可能导致系统 GUI 冻结，并且只能通过系统重置来恢复。当在释放过程中发生错误时，剩余的待处理释放将被取消。任何释放错误都将被报告。当进程终止时，CUDA 驱动程序能够释放终止进程分配的所有资源。

一旦发生以下事件，释放队列将自动刷新：

由于内存不足错误导致分配失败。在刷新所有释放后，重新尝试分配。
释放队列已达到其最大大小，默认为 10。用户可以通过设置环境变量 NUMBA_CUDA_MAX_PENDING_DEALLOCS_COUNT 来覆盖。例如，NUMBA_CUDA_MAX_PENDING_DEALLOCS_COUNT=20 将限制增加到 20。
待释放资源的累计最大字节大小已达到。这默认为设备内存容量的 20%。用户可以通过设置环境变量 NUMBA_CUDA_MAX_PENDING_DEALLOCS_RATIO 来覆盖。例如，NUMBA_CUDA_MAX_PENDING_DEALLOCS_RATIO=0.5 将限制设置为容量的 50%。

有时，希望将资源释放推迟到代码段结束。通常，用户希望避免由于释放而引起的任何隐式同步。这可以通过使用以下上下文管理器来完成：

numba.cuda.defer_cleanup()

临时禁用内存释放。使用此功能可防止资源释放中断异步执行。

例如

with defer_cleanup():
    # all cleanup is deferred in here
    do_speed_critical_code()
# cleanup can occur here

注意：此上下文管理器可以嵌套使用。