内存管理

数据传输

尽管 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即时编译的函数内调用。

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连续的，则返回真。

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）

在编译时由 jitted 函数的全局变量定义的全局变量（例如 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）

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

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

参见

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

常量内存

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

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

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