企业官网建设流程全解析

• 编译器
• 高性能计算

本篇指南以 Numba 官方 CUDA 文档 docs/source/cuda/memory.rst 为骨架，深入讲解 Numba CUDA 编程中的内存管理全貌：如何在主机与设备之间手动控制数据传输、使用页锁定（pinned）内存、映射内存与托管（managed）内存加速异步拷贝，如何通过流（stream）组织异步执行，以及在设备端如何使用共享内存（shared memory）、本地内存（local memory）与常量内存（constant memory）提升内核性能，最后剖析 Numba 的延迟释放（deferred deallocation）机制。读完本文，你将能够为数据密集型 CUDA 内核设计一套高效、可预测的内存策略，并理解 Numba 在底层究竟如何管理设备资源。

为什么需要手动管理数据传输

Numba 在调用 CUDA 内核时，可以自动把 NumPy 数组传输到设备端。但这种自动传输是保守的：它无法判断某个数组在内核中是否只读，因此总是会在内核结束后把设备内存无条件回传到主机。对于只读的输入数组，这种隐式回传是纯粹的浪费。

为了避免不必要的传输，Numba 提供了手动控制数据传输的 API，让你精确决定"何时传上去、何时传回来、要不要传回来"。这些 API 全部定义在 numba/cuda/api.py 中，是宿主端（host）代码的核心工具。

数据传输 API

device_array 与 device_array_like：在设备上分配空数组

cuda.device_array(shape, dtype=np.float64, strides=None, order='C', stream=0)在设备端分配一块未初始化的数组，语义上等价于numpy.empty（源码见 api.py）：

from numba import cuda import numpy as np d_ary = cuda.device_array(shape=(100, 100), dtype=np.float64) # 等价于 np.empty((100, 100), dtype=np.float64)，但内存位于 GPU 上

cuda.device_array_like(ary, stream=0)则根据已有数组ary的 shape、dtype 和内存布局（C 连续或 F 连续）创建一块同构的设备数组（见 api.py）：

host_ary = np.arange(1000).reshape(10, 100) # C 连续 d_ary = cuda.device_array_like(host_ary) # 设备端同样是 C 连续的 10x100 数组

to_device：主机到设备的核心入口

cuda.to_device(obj, stream=0, copy=True, to=None)负责分配设备内存并把 NumPy 数组或结构化标量拷贝上去（见 api.py）：

ary = np.arange(10) d_ary = cuda.to_device(ary) # 默认在流 0（同步）上执行 host->device 拷贝

把传输挂到某个流上，实现异步拷贝：

stream = cuda.stream() d_ary = cuda.to_device(ary, stream=stream) # 异步入队

回传设备数据到主机有两种方式：

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)

从源码看，to_device实际调用的是devicearray.auto_device(obj, stream=stream, copy=copy, user_explicit=True)；当传入to参数时，则直接把数据拷贝到已有的设备数组上（to.copy_to_device(obj, stream=stream)），避免重复分配。

设备数组：DeviceNDArray 及其方法

手动分配得到的是numba.cuda.cudadrv.devicearray.DeviceNDArray，它定义在 numba/cuda/cudadrv/devicearray.py 中，只能在宿主代码中调用，不能用于 CUDA 设备函数内部。其关键方法如下（对应文档中列出的成员）：

需要注意，DeviceNDArray还实现了 CUDA Array Interface（__cuda_array_interface__，见 devicearray.py），这意味着它可以与其他支持该接口的库无缝互操作，详见 cuda_array_interface.rst。

as_cuda_array 与 is_cuda_array：消费外部 GPU 缓冲区

除了 Numba 自己分配的设备数组，Numba 还可以消费任何实现了 CUDA Array Interface 的对象，通过创建 GPU 缓冲区的视图（不拷贝数据）将其包装为DeviceNDArray：

def is_cuda_array(obj): return hasattr(obj, '__cuda_array_interface__')

is_cuda_array仅检查对象是否定义了__cuda_array_interface__属性，不验证接口的合法性；as_cuda_array(obj, sync=True)则会真正创建视图，并对导入的流（如果有）执行同步（见 api.py）：

if not is_cuda_array(obj): raise TypeError("*obj* doesn't implement the cuda array interface.")

as_cuda_array返回的数组会持有obj的引用，保证底层 GPU 缓冲区的生命周期；sync=True时还会同步导入的流，确保其他库在该流上排队的工作完成后再使用该视图。

页锁定（Pinned）内存

常规主机内存是可分页的，CUDA 在 host->device 拷贝时需要通过驱动临时把页面锁定，这一过程会带来额外开销。页锁定（pinned / page-locked）内存从分配之初就锁定在物理内存中，驱动可以直接使用 DMA 进行传输，显著提升拷贝带宽，也是异步传输（overlapped transfer）的前提。

Numba 提供三个相关 API（见 api.py）：

• cuda.pinned_array(shape, dtype=np.float64, strides=None, order='C')：分配一块页锁定的 NumPy ndarray，语义类似np.empty。底层调用current_context().memhostalloc(bytesize)，在 CUDA 驱动层面对应cuMemHostAlloc。
• cuda.pinned_array_like(ary)：按ary的 shape/dtype/布局创建页锁定数组（见 api.py）。
• cuda.pinned(*arylist)：一个上下文管理器，用于临时把一批已经存在的普通主机数组页锁定，退出with块后自动解锁（见 api.py）：

import numpy as np from numba import cuda host_ary = np.arange(1000) with cuda.pinned(host_ary): d_ary = cuda.to_device(host_ary) # 拷贝期间内存已被页锁定

典型用法是：先pinned_array分配主机缓冲区，然后用cuda.to_device(..., stream=stream)发起异步传输，同时 CPU 可以继续做其他计算，实现拷贝与计算的重叠。

映射（Mapped）内存

映射内存是页锁定内存的进阶形态：缓冲区同时映射到主机地址空间与设备地址空间，主机和设备可以共享同一块物理内存，无需显式拷贝。设备通过 PCIe 直接访问这块内存（零拷贝），代价是每次访问都经过 PCIe，适合小规模、访问稀疏的数据。

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

其中portable=True允许该内存被多个设备使用；wc=True启用 write-combined 分配——主机写入更快、设备读取更快，但主机读取和设备写入更慢（见 api.py）。返回的是MappedNDArray，可通过cuda.mapped_array_like(ary, stream=0, portable=False, wc=False)按已有数组创建。

cuda.mapped(*arylist, stream=0)则是上下文管理器形式，临时把一批主机数组映射到设备，退出时自动释放映射（见 api.py）：

with cuda.mapped(ary1, ary2) as mapped_arys: # mapped_arys 为设备视图列表，可直接作为内核参数 ...

托管（Managed）内存

托管内存基于 CUDA 的Unified Memory（统一寻址）：一块内存同时可被 CPU 与 GPU 访问，CUDA 驱动自动在两者之间按需迁移页面，程序员无需关心拷贝。这在内存访问模式复杂或无法预先规划传输时非常方便。

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

attach_global=True表示全局附加，内存可被任意设备上的任意流访问；attach_global=False则退化为仅主机附加（host attachment），只有计算能力（Compute Capability）6.0 及以上的设备才能访问（见 api.py）。返回类型为ManagedNDArray。

需要说明的是，托管内存的支持存在平台差异：在 Linux/x86 与 PowerPC 上完全支持，在 Windows 与 Linux/AArch64 上仍视为实验特性（见managed_array文档字符串）。

流（Streams）：组织异步执行

CUDA 流是一个命令队列：放入同一流中的操作（内核启动、内存拷贝）按顺序执行，不同流之间的操作可以并行。Numba 中流可以传给接受流的函数（如 host<->device 拷贝），也可以写进内核启动配置，从而实现异步执行。

创建与获取流

• cuda.stream()：创建一条新流，等价于驱动的cuStreamCreate（见 api.py 与 driver.py）。
• cuda.default_stream()：获取默认流。CUDA 语义中默认流既可能是 legacy 默认流也可能是 per-thread 默认流，取决于使用哪套 CUDA API；Numba 目前总是使用 legacy 默认流对应的 API，但保留未来切换的选项（见 api.py）。
• cuda.legacy_default_stream()：获取 legacy 默认流（对应驱动常量CU_STREAM_LEGACY，见 driver.py）。
• cuda.per_thread_default_stream()：获取 per-thread 默认流（对应CU_STREAM_PER_THREAD，见 driver.py）。
• cuda.external_stream(ptr)：用外部分配（如通过其他 CUDA 库、C/C++ 代码创建的）的流指针，包装成一个 Numba 流对象；ptr必须是int类型（见 api.py 与 driver.py）。

stream = cuda.stream() d_ary = cuda.to_device(host_ary, stream=stream) # 异步拷贝入队 kernelgrid, block, stream # 内核在同一个流中排队 result = d_ary.copy_to_host(stream=stream) # 异步回传

流对象的方法

numba.cuda.cudadrv.driver.Stream（见 driver.py）提供两个文档明确列出的方法：

• synchronize()：阻塞等待流中所有命令执行完毕，同时提交所有挂起的（异步）内存传输，底层调用cuStreamSynchronize（见 driver.py）。
• auto_synchronize()：上下文管理器，进入with块后执行流中的操作，退出时自动调用synchronize()（见 driver.py）：

with stream.auto_synchronize(): kernelgrid, block, stream d_ary.copy_to_host(host_ary, stream=stream) # 此处保证流内所有操作已完成

此外，Stream还提供add_callback（向流添加回调）、async_done（返回一个asyncio.Future，流内操作全部完成后 resolve）等进阶能力。

共享内存与线程同步

共享内存：块内协作的手动缓存

**共享内存（shared memory）**是设备上一块容量有限的片上内存：同一线程块（block）内的所有线程均可读写它，访问速度远快于普通设备内存（DRAM），因此既可用于块内线程协作，也可当作手动管理的缓存。

• 内存只在内核运行期间分配一次，与传统运行时动态内存管理不同；
• 容量有限，需要根据设备规格合理规划。

在设备端（内核或设备函数中）通过cuda.shared.array(shape, type)分配：

from numba import cuda @cuda.jit def kernel(x): # 每个线程块分配一个 32 个 float32 的共享数组 sdata = cuda.shared.array(32, dtype=cuda.float32) tid = cuda.threadIdx.x sdata[tid] = x[tid] cuda.syncthreads() # 等待所有线程写入完成 x[tid] = sdata[(tid + 1) % 32] # 读取邻居线程的数据

shape 必须是一个简单常量表达式，文档明确规定了三种合法形式：

1. 字面量，如10；
2. 局部变量，其右侧是字面量或简单常量表达式，如先定义shape = 10再使用shape；
3. 编译时已定义在 jitted 函数全局命名空间中的全局变量。

并且该表达式的求值结果必须是 Python 的int（不能是 NumPy 标量或其他整数类标量类型）。type是元素类型对应的 Numba 类型，返回的数组对象可像普通设备数组一样通过索引读写。

从源码看，cuda.shared.array的整数维度与元组维度分别由 cudaimpl.py 中的 lowering 函数处理，它们把数组分配到 NVVM 的共享地址空间（nvvm.ADDRSPACE_SHARED），并通过_get_unique_smem_id为共享内存符号生成唯一名称，以避免 NVVM 在 PTX 输出中错误内部化共享内存导致的缺陷。

syncthreads：块内屏障

cuda.syncthreads()实现与多线程编程中**屏障（barrier）**相同的语义：它等待线程块内的所有线程都调用它，然后一起返回。典型模式是"每个线程写入共享数组的一个元素，然后syncthreads等待全部写完，再安全地读取别人的数据"。文档同时指出，矩阵乘法示例（cuda-matmul）是共享内存+同步的经典参考实现。

动态共享内存

共享内存既可以是静态的（编译期确定大小），也可以是动态的（启动内核时指定字节数）。要使用动态共享内存，先在内核中声明大小为 0 的共享数组：

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

然后在内核启动配置的第四个参数中指定动态共享内存的字节数：

kernel_funcgrid_dim, block_dim, stream, dyn_shared_mem_size

即启动配置kernel_func[grid_dim, block_dim, stream, dyn_shared_mem_size]中的 4 个参数依次为：网格维度、块维度、流、动态共享内存字节数。

关键陷阱：所有动态共享内存数组彼此别名（alias）。因为它们都指向同一块动态分配的共享内存。例如：

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解释的结果。

如果希望多个动态共享数组互不干扰，需要取互不相交（disjoint）的视图：

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（通过[1:]偏移一个元素）。两个值互不覆盖，输出：

3.140000 1

本地内存（Local memory）

本地内存是每个线程私有的内存区域，在标量局部变量不够用时充当线程的临时工作区（scratchpad）。与共享内存类似，它在内核运行期间只分配一次，而非传统意义上的运行时动态分配。本地内存在物理上通常位于设备内存（DRAM）中，访问延迟高于寄存器，因此应谨慎使用。

cuda.local.array(shape, type)

shape同样必须是简单常量表达式（规则与共享内存相同：字面量、右侧为常量的局部变量、编译时可见的全局变量），且结果必须是 Pythonint；type是元素的 Numba 类型。返回的数组仅当前线程可见，可像标准数组一样索引读写。相关 lowering 实现见 cudaimpl.py，将数组分配到本地地址空间。

常量内存（Constant memory）

常量内存是一块只读、有缓存、位于片外（off-chip）的内存，所有线程都可以访问，且由宿主端分配。由于有专用缓存，当块内所有线程读取同一地址时（广播式访问），常量内存访问效率极高，适合存储系数表、查找表等只读数据。

cuda.const.array_like(arr)

cuda.const.array_like(arr)基于类数组对象arr在常量内存中分配并填充一个数组。从源码看，其 lowering 是一个空操作（cudaimpl.py）：实际数组在CUDATargetContext.make_constant_array阶段就已经被创建为常量数组，运行时无需再做任何处理。

import numpy as np from numba import cuda coefficients = np.array([1.0, 2.0, 3.0], dtype=np.float32) @cuda.jit def kernel(x): i = cuda.grid(1) if i < x.size: x[i] = x[i] * cuda.const.array_like(coefficients)[i % 3]

释放行为（Deallocation Behavior）

Numba 内部的内存管理采用**延迟释放（deferred deallocation）**策略。如果使用了外部内存管理插件（EMM Plugin，见 external-memory-management.rst），释放行为可能不同，应参考对应插件的文档。

按上下文追踪的延迟释放

所有 CUDA 资源的释放都按 CUDA 上下文（context）追踪。当某个设备内存的最后一个引用被丢弃时，底层内存并不会立即释放，而是被放入待释放队列（pending deallocations queue）。这种设计有两个好处：

1. 避免隐式同步打断异步执行：资源释放 API 可能导致设备同步，从而打断性能关键路径上的异步执行；推迟释放可以避免这部分延迟。
2. 降低释放错误的风险：某些释放错误可能导致剩余释放全部失败；持续的释放错误可能引发 CUDA 驱动层面的严重问题——严重时可能导致 CUDA 驱动段错误（segmentation fault），最坏情况下甚至冻结系统 GUI，只能通过系统重启恢复。因此当释放过程中出现错误时，其余待释放项会被取消，且所有释放错误都会被上报。进程终止时，CUDA 驱动能够回收该进程持有的全部资源。

队列自动刷新的三个触发条件

待释放队列在以下事件发生时自动刷新（flush）：

1. 分配因内存不足（OOM）失败时：先刷新所有待释放项，再重试分配；
2. 队列达到最大条目数时：默认为 10，可通过环境变量NUMBA_CUDA_MAX_PENDING_DEALLOCS_COUNT覆盖。例如NUMBA_CUDA_MAX_PENDING_DEALLOCS_COUNT=20将上限提高到 20；
3. 待释放资源累计字节数达到上限时：默认为设备内存容量的 20%，可通过环境变量NUMBA_CUDA_MAX_PENDING_DEALLOCS_RATIO覆盖。例如NUMBA_CUDA_MAX_PENDING_DEALLOCS_RATIO=0.5将上限设为容量的 50%。

这两个环境变量在 numba/core/config.py 中定义，源码注释明确给出默认值：CUDA_DEALLOCS_COUNT默认 10、CUDA_DEALLOCS_RATIO默认 0.2。底层实现为_PendingDeallocs类（driver.py）：它用双端队列维护待释放项，add_item入队并累计字节数，一旦条目数超过CUDA_DEALLOCS_COUNT或累计字节超过memory_capacity * CUDA_DEALLOCS_RATIO（_max_pending_bytes）就立即clear()刷新；disable()/is_disabled则支持暂时挂起刷新。

defer_cleanup：手动推迟清理

有时我们希望把资源释放推迟到某段代码结束，最常见的动机是避免释放带来的隐式同步打断关键区间的异步执行。cuda.defer_cleanup()正是为此设计的上下文管理器（见 api.py）：

with cuda.defer_cleanup(): # 此区间内的所有清理操作都被推迟 do_speed_critical_code() # 离开 with 块后，清理可以正常发生

该上下文管理器可以嵌套使用。从源码看，它分别调用了内存管理器的defer_cleanup()和待释放队列的disable()（driver.py），从而在区间内完全挂起释放动作。

选择合适内存策略的实践建议

在动手优化前，建议先确认默认的保守自动传输确实成为瓶颈：对只读数组使用手动to_device，对回传需求使用显式copy_to_host，再结合流与页锁定内存，往往就能获得显著的带宽收益；而共享内存与常量内存的优化则更依赖具体内核的访问模式。

延伸阅读

• CUDA Array Interface 详解：理解as_cuda_array依赖的跨库互操作协议
• 内核启动与编译指南：内核启动配置grid/block/stream的完整语法
• 矩阵乘法示例：共享内存 + 同步的经典实战
• 外部内存管理插件（EMM Plugin）提案：自定义内存分配/释放策略的扩展机制
• 相关测试：numba/cuda/tests/cudadrv/test_deallocations.py 验证了待释放队列的条数上限、字节上限与 OOM 刷新行为

• 编译器
• 高性能计算