3.3。内存管理

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

3.3.1。数据传输

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

  1. numba.cuda.device_array(shape, dtype=np.float, strides=None, order='C', stream=0)

分配一个空设备 ndarray。与numpy.empty()类似。

  1. numba.cuda.device_array_like(ary, stream=0)

使用数组中的信息调用 cuda.devicearray（）。

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

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

要将 host-＆gt;设备复制为 numpy 数组：

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

要将传输排入队列：

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

得到的d_ary是DeviceNDArray。

要复制 device-＆gt;主机：

  1. hary = d_ary.copy_to_host()

要将 device-＆gt;主机复制到现有数组：

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

要将传输排入队列：

  1. hary = d_ary.copy_to_host(stream=stream)

除了设备阵列，Numba 还可以使用任何实现 cuda 阵列接口的对象。通过使用以下 API 创建 GPU 缓冲区视图，还可以将这些对象手动转换为 Numba 设备阵列：

  1. numba.cuda.as_cuda_array(obj)

从任何实现 cuda-array-interface 的对象创建 DeviceNDArray。

创建基础 GPU 缓冲区的视图。没有复制数据。生成的 DeviceNDArray 将从 <cite>obj</cite> 获取引用。

  1. numba.cuda.is_cuda_array(obj)

测试对象是否已定义 <cite>cuda_array_interface</cite> 。

不验证接口的有效性。

3.3.1.1。设备阵列

设备阵列引用具有以下方法。这些方法将在主机代码中调用，而不是在 CUDA-jitted 函数中调用。

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

GPU 上阵列类型

  1. copy_to_host(ary=None, stream=0)

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

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

始终返回主机阵列。

例：

  1. import numpy as np
  2. from numba import cuda
  4. arr = np.arange(1000)
  5. d_arr = cuda.to_device(arr)
  7. my_kernel[100, 100](d_arr)
  9. result_array = d_arr.copy_to_host()
  1. is_c_contiguous()

如果数组是 C-contiguous，则返回 true。

  1. is_f_contiguous()

如果数组是 Fortran-contiguous，则返回 true。

  1. ravel(order='C', stream=0)

在不改变其内容的情况下展平阵列，类似于 numpy.ndarray.ravel() 。

  1. reshape(*newshape, **kws)

与 numpy.ndarray.reshape() 类似，重塑阵列而不改变其内容。例：

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

注意

DeviceNDArray 定义 cuda 阵列接口。

3.3.2。固定内存

  1. numba.cuda.pinned(*arylist)

用于临时固定主机 ndarray 序列的上下文管理器。

  1. numba.cuda.pinned_array(shape, dtype=np.float, strides=None, order='C')

使用固定（页面锁定）的缓冲区分配 np.ndarray。与 np.empty（）类似。

3.3.3。流

  1. numba.cuda.stream()

创建表示设备命令队列的 CUDA 流。

CUDA 流有以下方法：

  1. class numba.cuda.cudadrv.driver.Stream(context, handle, finalizer)
  1. auto_synchronize()

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

  1. synchronize()

等待此流中的所有命令执行。这将提交任何挂起的内存传输。

3.3.4。共享内存和线程同步

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

与传统的动态内存管理不同，内存在内核持续时间内分配一次。

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

在设备上分配给定 形状 和 类型 的共享数组。必须在设备上调用此函数（即，从内核或设备函数）。 shape 是一个整数或整数元组，表示数组的维度，必须是一个简单的常量表达式。 类型 是需要存储在数组中的元素的 Numba 类型。

可以像任何普通设备阵列一样读取和写入返回的类似阵列的对象（例如通过索引）。

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

  1. numba.cuda.syncthreads()

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

也可以看看

矩阵乘法示例。

3.3.5。本地记忆

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

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

在设备上分配给定 形状 和 类型 的局部数组。 shape 是一个整数或整数元组，表示数组的维度，必须是一个简单的常量表达式。 类型 是需要存储在数组中的元素的 Numba 类型。该数组对当前线程是私有的。返回类似于数组的对象，可以像任何标准数组一样读取和写入（例如通过索引）。

3.3.6。恒定记忆

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

  1. numba.cuda.const.array_like(arr)

基于类似数组 arr ，在常量内存中分配并使数组可访问。

3.3.7。 SmartArrays（实验性）

Numba 提供类似数组的数据类型，可自动管理与设备之间的数据移动。在大多数情况下，它可以作为 <cite>numpy.ndarray</cite> 的直接替换，并且由 Numba 的 JIT 编译器支持’host’和’cuda’目标。

  1. class numba.SmartArray(obj=None, copy=True, shape=None, dtype=None, order=None, where='host')

一种支持主机和 GPU 存储的阵列类型。

  1. __init__(obj=None, copy=True, shape=None, dtype=None, order=None, where='host')

在’where’定义的内存空间中构造一个 SmartArray。有效的调用：

  • SmartArray（obj =＆lt; array-like object＆gt;，copy =＆lt; optional-true-or-false＆gt;）：

    从现有的类似数组的对象创建 SmartArray。 ‘copy’参数指定是采用还是复制它。

  • SmartArray（shape =＆lt; shape＆gt;，dtype =＆lt; dtype＆gt ;, order =＆lt; order＆gt;）

    在给定典型的 NumPy 数组属性的情况下，从头开始创建新的 SmartArray。

（可选的’where’参数指定最初分配数组的位置。（默认值：’host’）

  1. get(where='host')

返回给定内存空间中“self”的表示形式。

  1. mark_changed(where='host')

将给定位置标记为已更改，如果需要，则广播更新。

因此， <cite>SmartArray</cite> 对象可以作为函数参数传递给 jit 编译函数。每当执行 cuda.jit 编译的函数时，它将触发向 GPU 的数据传输（除非数据已存在）。但是，在函数完成后，不是将数据传回主机，而是将数据保留在设备上，如果有任何外部引用，则只更新主机端。因此，如果下一个操作是对 cuda.jit 编译函数的另一个调用，则不需要再次传输数据，从而使复合操作更有效（并且即使对于较小的数据大小也使得 GPU 的使用有利）。

3.3.8。解除分配行为

根据每个上下文跟踪所有 CUDA 资源的重新分配。当删除对设备存储器的最后一次引用时，将调度基础存储器以解除分配。释放不会立即发生。它被添加到待处理的解除分配队列中。这种设计有两个好处：

  1. 资源释放 API 可能导致设备同步;因此，打破任何异步执行。延迟重新分配可以避免性能关键代码部分的延迟。
  2. 一些释放错误可能导致所有剩余的解除分配失败。持续的释放错误可能会导致 CUDA 驱动程序级别的严重错误。在某些情况下，这可能意味着 CUDA 驱动程序中的分段错误。在最坏的情况下，这可能导致系统 GUI 冻结，并且只能通过系统重置进行恢复。在释放期间发生错误时，将取消剩余的挂起解除分配。将报告任何释放错误。当进程终止时，CUDA 驱动程序能够通过已终止的进程释放所有已分配的资源。

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

  • 由于内存不足错误导致分配失败。在清除所有解除分配后重试分配。
  • 释放队列已达到其最大大小，默认为 10.用户可以通过设置环境变量 <cite>NUMBA_CUDA_MAX_PENDING_DEALLOCS_COUNT</cite> 来覆盖。例如， <cite>NUMBA_CUDA_MAX_PENDING_DEALLOCS_COUNT = 20</cite> ，将限制增加到 20。
  • 达到挂起释放的资源的最大累积字节大小。这默认为设备内存容量的 20％。用户可以通过设置环境变量 <cite>NUMBA_CUDA_MAX_PENDING_DEALLOCS_RATIO</cite> 来覆盖。例如， <cite>NUMBA_CUDA_MAX_PENDING_DEALLOCS_RATIO = 0.5</cite> 将限制设置为容量的 50％。

有时，需要推迟资源释放，直到代码段结束。大多数情况下，用户希望避免因释放而导致的任何隐式同步。这可以通过使用以下上下文管理器来完成：

  1. numba.cuda.defer_cleanup()

暂时禁用内存释放。使用它来防止资源释放打破异步执行。

例如：

  1. with defer_cleanup():
  2.     # all cleanup is deferred in here
  3.     do_speed_critical_code()
  4. # cleanup can occur here

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