编译器高性能计算【免费下载链接】numbaNumPy aware dynamic Python compiler using LLVM项目地址https://gitcode.com/gh_mirrors/nu/numba点击查看免费下载本篇指南以 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代码的核心工具。数据传输 APIdevice_array 与 device_array_like在设备上分配空数组cuda.device_array(shape, dtypenp.float64, stridesNone, orderC, stream0)在设备端分配一块未初始化的数组语义上等价于numpy.empty源码见 api.pyfrom numba import cuda import numpy as np d_ary cuda.device_array(shape(100, 100), dtypenp.float64) # 等价于 np.empty((100, 100), dtypenp.float64)但内存位于 GPU 上cuda.device_array_like(ary, stream0)则根据已有数组ary的 shape、dtype 和内存布局C 连续或 F 连续创建一块同构的设备数组见 api.pyhost_ary np.arange(1000).reshape(10, 100) # C 连续 d_ary cuda.device_array_like(host_ary) # 设备端同样是 C 连续的 10x100 数组to_device主机到设备的核心入口cuda.to_device(obj, stream0, copyTrue, toNone)负责分配设备内存并把 NumPy 数组或结构化标量拷贝上去见 api.pyary np.arange(10) d_ary cuda.to_device(ary) # 默认在流 0同步上执行 host-device 拷贝把传输挂到某个流上实现异步拷贝stream cuda.stream() d_ary cuda.to_device(ary, streamstream) # 异步入队回传设备数据到主机有两种方式hary d_ary.copy_to_host() # 创建新数组并回传 # 或者回传到已有数组 ary np.empty(shaped_ary.shape, dtyped_ary.dtype) d_ary.copy_to_host(ary) # 或者挂到流上异步回传 hary d_ary.copy_to_host(streamstream)从源码看to_device实际调用的是devicearray.auto_device(obj, streamstream, copycopy, user_explicitTrue)当传入to参数时则直接把数据拷贝到已有的设备数组上to.copy_to_device(obj, streamstream)避免重复分配。设备数组DeviceNDArray 及其方法手动分配得到的是numba.cuda.cudadrv.devicearray.DeviceNDArray它定义在 numba/cuda/cudadrv/devicearray.py 中只能在宿主代码中调用不能用于 CUDA 设备函数内部。其关键方法如下对应文档中列出的成员方法作用copy_to_host(aryNone, stream0)拷贝到主机ary为None时新建 NumPy 数组返回指定stream时异步执行否则同步返回见 devicearray.pyis_c_contiguous()判断数组是否 C 连续见 devicearray.pyis_f_contiguous()判断数组是否 Fortran 连续见 devicearray.pyravel(orderC, stream0)展平连续数组若数组不连续则抛出异常见 devicearray.pyreshape(*newshape, **kws)不改变数据地重塑形状与numpy.ndarray.reshape类似若 reshape 需要拷贝则抛出NotImplementedError见 devicearray.py需要注意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 缓冲区的视图不拷贝数据将其包装为DeviceNDArraydef is_cuda_array(obj): return hasattr(obj, __cuda_array_interface__)is_cuda_array仅检查对象是否定义了__cuda_array_interface__属性不验证接口的合法性as_cuda_array(obj, syncTrue)则会真正创建视图并对导入的流如果有执行同步见 api.pyif not is_cuda_array(obj): raise TypeError(*obj* doesnt implement the cuda array interface.)as_cuda_array返回的数组会持有obj的引用保证底层 GPU 缓冲区的生命周期syncTrue时还会同步导入的流确保其他库在该流上排队的工作完成后再使用该视图。页锁定Pinned内存常规主机内存是可分页的CUDA 在 host-device 拷贝时需要通过驱动临时把页面锁定这一过程会带来额外开销。页锁定pinned / page-locked内存从分配之初就锁定在物理内存中驱动可以直接使用 DMA 进行传输显著提升拷贝带宽也是异步传输overlapped transfer的前提。Numba 提供三个相关 API见 api.pycuda.pinned_array(shape, dtypenp.float64, stridesNone, orderC)分配一块页锁定的 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.pyimport 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(..., streamstream)发起异步传输同时 CPU 可以继续做其他计算实现拷贝与计算的重叠。映射Mapped内存映射内存是页锁定内存的进阶形态缓冲区同时映射到主机地址空间与设备地址空间主机和设备可以共享同一块物理内存无需显式拷贝。设备通过 PCIe 直接访问这块内存零拷贝代价是每次访问都经过 PCIe适合小规模、访问稀疏的数据。cuda.mapped_array(shape, dtypenp.float64, stridesNone, orderC, stream0, portableFalse, wcFalse)其中portableTrue允许该内存被多个设备使用wcTrue启用 write-combined 分配——主机写入更快、设备读取更快但主机读取和设备写入更慢见 api.py。返回的是MappedNDArray可通过cuda.mapped_array_like(ary, stream0, portableFalse, wcFalse)按已有数组创建。cuda.mapped(*arylist, stream0)则是上下文管理器形式临时把一批主机数组映射到设备退出时自动释放映射见 api.pywith cuda.mapped(ary1, ary2) as mapped_arys: # mapped_arys 为设备视图列表可直接作为内核参数 ...托管Managed内存托管内存基于 CUDA 的Unified Memory统一寻址一块内存同时可被 CPU 与 GPU 访问CUDA 驱动自动在两者之间按需迁移页面程序员无需关心拷贝。这在内存访问模式复杂或无法预先规划传输时非常方便。cuda.managed_array(shape, dtypenp.float64, stridesNone, orderC, stream0, attach_globalTrue)attach_globalTrue表示全局附加内存可被任意设备上的任意流访问attach_globalFalse则退化为仅主机附加host attachment只有计算能力Compute Capability6.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 APINumba 目前总是使用 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, streamstream) # 异步拷贝入队 kernelgrid, block, stream # 内核在同一个流中排队 result d_ary.copy_to_host(streamstream) # 异步回传流对象的方法numba.cuda.cudadrv.driver.Stream见 driver.py提供两个文档明确列出的方法synchronize()阻塞等待流中所有命令执行完毕同时提交所有挂起的异步内存传输底层调用cuStreamSynchronize见 driver.py。auto_synchronize()上下文管理器进入with块后执行流中的操作退出时自动调用synchronize()见 driver.pywith stream.auto_synchronize(): kernelgrid, block, stream d_ary.copy_to_host(host_ary, streamstream) # 此处保证流内所有操作已完成此外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, dtypecuda.float32) tid cuda.threadIdx.x sdata[tid] x[tid] cuda.syncthreads() # 等待所有线程写入完成 x[tid] sdata[(tid 1) % 32] # 读取邻居线程的数据shape 必须是一个简单常量表达式文档明确规定了三种合法形式字面量如10局部变量其右侧是字面量或简单常量表达式如先定义shape 10再使用shape编译时已定义在 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, dtypenp.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, dtypenp.float32) i32_arr cuda.shared.array(0, dtypenp.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, dtypenp.float32) i32_arr cuda.shared.array(0, dtypenp.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同样必须是简单常量表达式规则与共享内存相同字面量、右侧为常量的局部变量、编译时可见的全局变量且结果必须是 Pythoninttype是元素的 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], dtypenp.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 BehaviorNumba 内部的内存管理采用**延迟释放deferred deallocation**策略。如果使用了外部内存管理插件EMM Plugin见 external-memory-management.rst释放行为可能不同应参考对应插件的文档。按上下文追踪的延迟释放所有 CUDA 资源的释放都按 CUDA 上下文context追踪。当某个设备内存的最后一个引用被丢弃时底层内存并不会立即释放而是被放入待释放队列pending deallocations queue。这种设计有两个好处避免隐式同步打断异步执行资源释放 API 可能导致设备同步从而打断性能关键路径上的异步执行推迟释放可以避免这部分延迟。降低释放错误的风险某些释放错误可能导致剩余释放全部失败持续的释放错误可能引发 CUDA 驱动层面的严重问题——严重时可能导致 CUDA 驱动段错误segmentation fault最坏情况下甚至冻结系统 GUI只能通过系统重启恢复。因此当释放过程中出现错误时其余待释放项会被取消且所有释放错误都会被上报。进程终止时CUDA 驱动能够回收该进程持有的全部资源。队列自动刷新的三个触发条件待释放队列在以下事件发生时自动刷新flush分配因内存不足OOM失败时先刷新所有待释放项再重试分配队列达到最大条目数时默认为 10可通过环境变量NUMBA_CUDA_MAX_PENDING_DEALLOCS_COUNT覆盖。例如NUMBA_CUDA_MAX_PENDING_DEALLOCS_COUNT20将上限提高到 20待释放资源累计字节数达到上限时默认为设备内存容量的 20%可通过环境变量NUMBA_CUDA_MAX_PENDING_DEALLOCS_RATIO覆盖。例如NUMBA_CUDA_MAX_PENDING_DEALLOCS_RATIO0.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.pywith cuda.defer_cleanup(): # 此区间内的所有清理操作都被推迟 do_speed_critical_code() # 离开 with 块后清理可以正常发生该上下文管理器可以嵌套使用。从源码看它分别调用了内存管理器的defer_cleanup()和待释放队列的disable()driver.py从而在区间内完全挂起释放动作。选择合适内存策略的实践建议场景推荐方案数据在内核结束后仍需回读且只读输入频繁用to_devicecopy_to_host手动控制避免自动回传追求最大 host-device 拷贝带宽、需要异步传输pinned_array 流stream实现拷贝与计算重叠数据小、访问稀疏、希望免拷贝共享mapped_array零拷贝映射访问模式复杂、难以预先规划传输managed_arrayUnified Memory 自动迁移块内协作归约、缓存频繁复用的数据shared.arraysyncthreads每个线程的私有临时工作区local.array全线程广播式只读数据系数表等const.array_like关键代码段想避免释放带来的同步defer_cleanup()上下文管理器在动手优化前建议先确认默认的保守自动传输确实成为瓶颈对只读数组使用手动to_device对回传需求使用显式copy_to_host再结合流与页锁定内存往往就能获得显著的带宽收益而共享内存与常量内存的优化则更依赖具体内核的访问模式。延伸阅读CUDA Array Interface 详解理解as_cuda_array依赖的跨库互操作协议内核启动与编译指南内核启动配置grid/block/stream的完整语法矩阵乘法示例共享内存 同步的经典实战外部内存管理插件EMM Plugin提案自定义内存分配/释放策略的扩展机制相关测试numba/cuda/tests/cudadrv/test_deallocations.py 验证了待释放队列的条数上限、字节上限与 OOM 刷新行为赞分享编译器高性能计算【免费下载链接】numbaNumPy aware dynamic Python compiler using LLVM项目地址https://gitcode.com/gh_mirrors/nu/numba点击查看免费下载相关推荐Numba CUDA 进程间内存共享CUDA IPC完全指南跨进程共享设备数组Numba CUDA 进程间内存共享CUDA IPC完全指南跨进程共享设备数组 CUDA IPCInter Process Communication编译器高性能计算Triton内存管理完全解析共享内存与缓存策略Triton内存管理完全解析共享内存与缓存策略 Triton语言和编译器作为深度学习计算的关键基础设施其内存管理机制直接影响着GPU计算的性能表现。本文将深编译器编程语言人工智能深度学习高性能计算终极指南ROCm GPU内存管理完全解析 - 高效内存分配与数据传输策略终极指南ROCm GPU内存管理完全解析 高效内存分配与数据传输策略 ROCm GPU内存管理是AMD开源计算平台中最重要的核心技术之一能够帮助开发者充分利开发工具高性能计算文档创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 Numba CUDA 内存管理完整指南：从数据传输到共享内存与释放策略
编译器高性能计算【免费下载链接】numbaNumPy aware dynamic Python compiler using LLVM项目地址https://gitcode.com/gh_mirrors/nu/numba点击查看免费下载本篇指南以 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代码的核心工具。数据传输 APIdevice_array 与 device_array_like在设备上分配空数组cuda.device_array(shape, dtypenp.float64, stridesNone, orderC, stream0)在设备端分配一块未初始化的数组语义上等价于numpy.empty源码见 api.pyfrom numba import cuda import numpy as np d_ary cuda.device_array(shape(100, 100), dtypenp.float64) # 等价于 np.empty((100, 100), dtypenp.float64)但内存位于 GPU 上cuda.device_array_like(ary, stream0)则根据已有数组ary的 shape、dtype 和内存布局C 连续或 F 连续创建一块同构的设备数组见 api.pyhost_ary np.arange(1000).reshape(10, 100) # C 连续 d_ary cuda.device_array_like(host_ary) # 设备端同样是 C 连续的 10x100 数组to_device主机到设备的核心入口cuda.to_device(obj, stream0, copyTrue, toNone)负责分配设备内存并把 NumPy 数组或结构化标量拷贝上去见 api.pyary np.arange(10) d_ary cuda.to_device(ary) # 默认在流 0同步上执行 host-device 拷贝把传输挂到某个流上实现异步拷贝stream cuda.stream() d_ary cuda.to_device(ary, streamstream) # 异步入队回传设备数据到主机有两种方式hary d_ary.copy_to_host() # 创建新数组并回传 # 或者回传到已有数组 ary np.empty(shaped_ary.shape, dtyped_ary.dtype) d_ary.copy_to_host(ary) # 或者挂到流上异步回传 hary d_ary.copy_to_host(streamstream)从源码看to_device实际调用的是devicearray.auto_device(obj, streamstream, copycopy, user_explicitTrue)当传入to参数时则直接把数据拷贝到已有的设备数组上to.copy_to_device(obj, streamstream)避免重复分配。设备数组DeviceNDArray 及其方法手动分配得到的是numba.cuda.cudadrv.devicearray.DeviceNDArray它定义在 numba/cuda/cudadrv/devicearray.py 中只能在宿主代码中调用不能用于 CUDA 设备函数内部。其关键方法如下对应文档中列出的成员方法作用copy_to_host(aryNone, stream0)拷贝到主机ary为None时新建 NumPy 数组返回指定stream时异步执行否则同步返回见 devicearray.pyis_c_contiguous()判断数组是否 C 连续见 devicearray.pyis_f_contiguous()判断数组是否 Fortran 连续见 devicearray.pyravel(orderC, stream0)展平连续数组若数组不连续则抛出异常见 devicearray.pyreshape(*newshape, **kws)不改变数据地重塑形状与numpy.ndarray.reshape类似若 reshape 需要拷贝则抛出NotImplementedError见 devicearray.py需要注意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 缓冲区的视图不拷贝数据将其包装为DeviceNDArraydef is_cuda_array(obj): return hasattr(obj, __cuda_array_interface__)is_cuda_array仅检查对象是否定义了__cuda_array_interface__属性不验证接口的合法性as_cuda_array(obj, syncTrue)则会真正创建视图并对导入的流如果有执行同步见 api.pyif not is_cuda_array(obj): raise TypeError(*obj* doesnt implement the cuda array interface.)as_cuda_array返回的数组会持有obj的引用保证底层 GPU 缓冲区的生命周期syncTrue时还会同步导入的流确保其他库在该流上排队的工作完成后再使用该视图。页锁定Pinned内存常规主机内存是可分页的CUDA 在 host-device 拷贝时需要通过驱动临时把页面锁定这一过程会带来额外开销。页锁定pinned / page-locked内存从分配之初就锁定在物理内存中驱动可以直接使用 DMA 进行传输显著提升拷贝带宽也是异步传输overlapped transfer的前提。Numba 提供三个相关 API见 api.pycuda.pinned_array(shape, dtypenp.float64, stridesNone, orderC)分配一块页锁定的 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.pyimport 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(..., streamstream)发起异步传输同时 CPU 可以继续做其他计算实现拷贝与计算的重叠。映射Mapped内存映射内存是页锁定内存的进阶形态缓冲区同时映射到主机地址空间与设备地址空间主机和设备可以共享同一块物理内存无需显式拷贝。设备通过 PCIe 直接访问这块内存零拷贝代价是每次访问都经过 PCIe适合小规模、访问稀疏的数据。cuda.mapped_array(shape, dtypenp.float64, stridesNone, orderC, stream0, portableFalse, wcFalse)其中portableTrue允许该内存被多个设备使用wcTrue启用 write-combined 分配——主机写入更快、设备读取更快但主机读取和设备写入更慢见 api.py。返回的是MappedNDArray可通过cuda.mapped_array_like(ary, stream0, portableFalse, wcFalse)按已有数组创建。cuda.mapped(*arylist, stream0)则是上下文管理器形式临时把一批主机数组映射到设备退出时自动释放映射见 api.pywith cuda.mapped(ary1, ary2) as mapped_arys: # mapped_arys 为设备视图列表可直接作为内核参数 ...托管Managed内存托管内存基于 CUDA 的Unified Memory统一寻址一块内存同时可被 CPU 与 GPU 访问CUDA 驱动自动在两者之间按需迁移页面程序员无需关心拷贝。这在内存访问模式复杂或无法预先规划传输时非常方便。cuda.managed_array(shape, dtypenp.float64, stridesNone, orderC, stream0, attach_globalTrue)attach_globalTrue表示全局附加内存可被任意设备上的任意流访问attach_globalFalse则退化为仅主机附加host attachment只有计算能力Compute Capability6.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 APINumba 目前总是使用 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, streamstream) # 异步拷贝入队 kernelgrid, block, stream # 内核在同一个流中排队 result d_ary.copy_to_host(streamstream) # 异步回传流对象的方法numba.cuda.cudadrv.driver.Stream见 driver.py提供两个文档明确列出的方法synchronize()阻塞等待流中所有命令执行完毕同时提交所有挂起的异步内存传输底层调用cuStreamSynchronize见 driver.py。auto_synchronize()上下文管理器进入with块后执行流中的操作退出时自动调用synchronize()见 driver.pywith stream.auto_synchronize(): kernelgrid, block, stream d_ary.copy_to_host(host_ary, streamstream) # 此处保证流内所有操作已完成此外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, dtypecuda.float32) tid cuda.threadIdx.x sdata[tid] x[tid] cuda.syncthreads() # 等待所有线程写入完成 x[tid] sdata[(tid 1) % 32] # 读取邻居线程的数据shape 必须是一个简单常量表达式文档明确规定了三种合法形式字面量如10局部变量其右侧是字面量或简单常量表达式如先定义shape 10再使用shape编译时已定义在 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, dtypenp.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, dtypenp.float32) i32_arr cuda.shared.array(0, dtypenp.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, dtypenp.float32) i32_arr cuda.shared.array(0, dtypenp.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同样必须是简单常量表达式规则与共享内存相同字面量、右侧为常量的局部变量、编译时可见的全局变量且结果必须是 Pythoninttype是元素的 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], dtypenp.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 BehaviorNumba 内部的内存管理采用**延迟释放deferred deallocation**策略。如果使用了外部内存管理插件EMM Plugin见 external-memory-management.rst释放行为可能不同应参考对应插件的文档。按上下文追踪的延迟释放所有 CUDA 资源的释放都按 CUDA 上下文context追踪。当某个设备内存的最后一个引用被丢弃时底层内存并不会立即释放而是被放入待释放队列pending deallocations queue。这种设计有两个好处避免隐式同步打断异步执行资源释放 API 可能导致设备同步从而打断性能关键路径上的异步执行推迟释放可以避免这部分延迟。降低释放错误的风险某些释放错误可能导致剩余释放全部失败持续的释放错误可能引发 CUDA 驱动层面的严重问题——严重时可能导致 CUDA 驱动段错误segmentation fault最坏情况下甚至冻结系统 GUI只能通过系统重启恢复。因此当释放过程中出现错误时其余待释放项会被取消且所有释放错误都会被上报。进程终止时CUDA 驱动能够回收该进程持有的全部资源。队列自动刷新的三个触发条件待释放队列在以下事件发生时自动刷新flush分配因内存不足OOM失败时先刷新所有待释放项再重试分配队列达到最大条目数时默认为 10可通过环境变量NUMBA_CUDA_MAX_PENDING_DEALLOCS_COUNT覆盖。例如NUMBA_CUDA_MAX_PENDING_DEALLOCS_COUNT20将上限提高到 20待释放资源累计字节数达到上限时默认为设备内存容量的 20%可通过环境变量NUMBA_CUDA_MAX_PENDING_DEALLOCS_RATIO覆盖。例如NUMBA_CUDA_MAX_PENDING_DEALLOCS_RATIO0.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.pywith cuda.defer_cleanup(): # 此区间内的所有清理操作都被推迟 do_speed_critical_code() # 离开 with 块后清理可以正常发生该上下文管理器可以嵌套使用。从源码看它分别调用了内存管理器的defer_cleanup()和待释放队列的disable()driver.py从而在区间内完全挂起释放动作。选择合适内存策略的实践建议场景推荐方案数据在内核结束后仍需回读且只读输入频繁用to_devicecopy_to_host手动控制避免自动回传追求最大 host-device 拷贝带宽、需要异步传输pinned_array 流stream实现拷贝与计算重叠数据小、访问稀疏、希望免拷贝共享mapped_array零拷贝映射访问模式复杂、难以预先规划传输managed_arrayUnified Memory 自动迁移块内协作归约、缓存频繁复用的数据shared.arraysyncthreads每个线程的私有临时工作区local.array全线程广播式只读数据系数表等const.array_like关键代码段想避免释放带来的同步defer_cleanup()上下文管理器在动手优化前建议先确认默认的保守自动传输确实成为瓶颈对只读数组使用手动to_device对回传需求使用显式copy_to_host再结合流与页锁定内存往往就能获得显著的带宽收益而共享内存与常量内存的优化则更依赖具体内核的访问模式。延伸阅读CUDA Array Interface 详解理解as_cuda_array依赖的跨库互操作协议内核启动与编译指南内核启动配置grid/block/stream的完整语法矩阵乘法示例共享内存 同步的经典实战外部内存管理插件EMM Plugin提案自定义内存分配/释放策略的扩展机制相关测试numba/cuda/tests/cudadrv/test_deallocations.py 验证了待释放队列的条数上限、字节上限与 OOM 刷新行为赞分享编译器高性能计算【免费下载链接】numbaNumPy aware dynamic Python compiler using LLVM项目地址https://gitcode.com/gh_mirrors/nu/numba点击查看免费下载相关推荐Numba CUDA 进程间内存共享CUDA IPC完全指南跨进程共享设备数组Numba CUDA 进程间内存共享CUDA IPC完全指南跨进程共享设备数组 CUDA IPCInter Process Communication编译器高性能计算Triton内存管理完全解析共享内存与缓存策略Triton内存管理完全解析共享内存与缓存策略 Triton语言和编译器作为深度学习计算的关键基础设施其内存管理机制直接影响着GPU计算的性能表现。本文将深编译器编程语言人工智能深度学习高性能计算终极指南ROCm GPU内存管理完全解析 - 高效内存分配与数据传输策略终极指南ROCm GPU内存管理完全解析 高效内存分配与数据传输策略 ROCm GPU内存管理是AMD开源计算平台中最重要的核心技术之一能够帮助开发者充分利开发工具高性能计算文档创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考