内存与数据移动#

GPU 有自己的内存，与主机分离。将数据传输到设备和从设备传回—— 以及在数据到达后选择正确的内存类型——是每个 CUDA 程序的基础。 本章涵盖 cuda-oxide 的内存抽象，从主机/设备传输到共享内存和 kernel ABI。

参见

CUDA 编程指南 -- 设备内存 关于 CUDA 内存层次结构和访问模式的权威参考。

CUDA 内存层次结构#

NVIDIA GPU 暴露多个内存层级，每层有不同的容量、延迟和作用范围：

内存

作用范围

典型大小

延迟

cuda-oxide API

寄存器

每线程

~255 × 32-bit

0 周期

局部变量

共享内存

每个 block

48--228 KB（取决于架构）

~5 周期

SharedArray、DynamicSharedArray

L1 缓存

每个 SM

与共享内存共享

硬件管理

自动

L2 缓存

芯片范围

最多 50 MB（Hopper）

~30 周期

自动

全局内存（DRAM）

所有线程

16--80 GB（HBM）

~400 周期

DeviceBuffer、DeviceBox

指导原则：将频繁访问的数据移到更快、更近的内存中。 寄存器最快但是每线程的；共享内存快速且对整个 block 可见； 全局内存容量大但速度慢。

gpu-programming/images/memory-hierarchy.svg

CUDA 内存层次结构，从最快（寄存器，每线程）到最大（全局 DRAM，所有线程）。 每一层以容量换取延迟。右侧面板显示了作用范围和每个级别的 cuda-oxide API。#

Context 和 Stream#

在深入内存 API 之前，需要介绍两个在每个代码示例中出现的 主机端概念：context 和 stream。

CUDA context（CudaContext）将主机线程绑定到特定的 GPU。 它拥有该设备上的所有资源——模块、stream、分配。通常在程序开始时创建一个：

use cuda_core::CudaContext;

let ctx = CudaContext::new(0).unwrap();   // 绑定到 GPU 0

CUDA stream（CudaStream）是 GPU 操作的有序队列。 在同一个 stream 上排队的操作按 FIFO 顺序执行——每个操作 看到前面所有操作的副作用。在不同 stream 上的操作可以 重叠并并发运行，它们之间没有顺序保证。

let stream = ctx.default_stream();     // 隐式的、始终可用的 stream
let work_stream = stream.fork()?;      // 一个新 stream，连接到父 stream 的当前进度位置

每次内存传输和 kernel 启动都需要一个 stream。对于单 stream 程序 （涵盖本书中的大部分示例），默认 stream就是你所需要的全部—— 一切按顺序执行，构造即正确。多 stream 流水线可以解锁计算和数据移动 之间的重叠，但需要通过事件或 join 进行显式同步：

同一 stream：      [kernel_A] → [memcpy_B] → [kernel_C]     （自动排序）
不同 stream：      [kernel_A on stream 1] | [memcpy_B on stream 2]  （并发，需要事件）
gpu-programming/images/cuda-streams.svg

上图：单 stream 执行，操作在 FIFO 中自动排序。 下图：多 stream 执行，stream A 和 B 并发运行，通过事件 建立 kernel_write 和 kernel_read 之间的数据依赖关系。#

参见

启动 Kernel 章节涵盖了启动宏中的 stream 使用， 异步 GPU 编程 部分涵盖了自动管理 stream 的 DeviceOperation。

DeviceBuffer -- 主机/设备传输#

cuda_core 中的 DeviceBuffer<T> 是分配设备内存和在主机与 GPU 之间 移动数据的主要方式：

use cuda_core::{CudaContext, DeviceBuffer};

let ctx = CudaContext::new(0).unwrap();
let stream = ctx.default_stream();

// 主机 → 设备：将主机切片复制到 GPU 内存
let a_dev = DeviceBuffer::from_host(&stream, &host_data).unwrap();

// 分配零初始化的设备内存
let mut c_dev = DeviceBuffer::<f32>::zeroed(&stream, 1024).unwrap();

// 设备 → 主机：读回结果
let results = c_dev.to_host_vec(&stream).unwrap();

关键方法#

方法

方向

描述

from_host(&stream, &[T])

主机 → 设备

分配 + 异步复制

zeroed(&stream, len)

--

分配 + 零填充

to_host_vec(&stream)

设备 → 主机

异步复制 + 返回 Vec<T>

copy_to_host(&stream, &mut [T])

设备 → 主机

复制到现有切片

cu_deviceptr()

--

原始 CUdeviceptr 用于 FFI

所有权和 drop#

DeviceBuffer 在 drop 时通过 cuMemFree 同步释放其分配。 这是一个阻塞的驱动调用——它内部会同步整个设备以确保没有正在执行的 kernel 仍在使用该内存。实际上，这意味着：

  • 在 kernel 运行时释放会阻塞主机线程，直到 GPU 完全空闲，然后释放内存。

  • 在同步后释放（例如在 to_host_vec 或 stream.synchronize() 之后） 没有额外开销，因为设备已经空闲。

对于单 stream 工作负载这是没问题的——一切都按 FIFO 顺序执行，所以当你 读回结果时，所有 kernel 都已完成，释放操作是瞬间的。在多 stream 场景中， 当你希望计算与内存操作重叠时，这一成本会显现出来；在一个 stream 上的 同步释放可能会阻塞所有其他 stream 上的工作。

DeviceBox -- 异步友好的设备内存#

cuda_async 中的 DeviceBox<T> 解决了同步释放的问题。在 drop 时， 它通过 cuMemFreeAsync 在专用的释放器 stream 上释放内存。 这是一个 stream 排序操作——释放在释放器 stream 上排队，仅在 该 stream 上的所有先前工作完成后才执行。关键的是，它不会同步设备：

use cuda_async::device_box::DeviceBox;
use cuda_async::device_context::init_device_contexts;

init_device_contexts(0, 1)?;  // 初始化设备 context 映射（默认设备 0）

// DeviceBox 包装一个设备指针；在 drop 时异步释放
let dev_ptr: DeviceBox<f32> = /* 由 DeviceOperation 链分配 */;
// 当 dev_ptr 被 drop 时，cuMemFreeAsync 会在释放器 stream 上被调用。
// 其他 stream 继续运行而不会阻塞。

在两者之间选择#

DeviceBuffer

DeviceBox

Crate

cuda_core

cuda_async

Drop 时释放

cuMemFree（同步 -- 阻塞设备）

cuMemFreeAsync（异步 -- 不阻塞）

与...一起使用

类型化同步启动

类型化异步启动

主机读回

to_host_vec()

通过显式 memcpy 操作

最适合

单 stream、阻塞式工作负载

多 stream、流水线工作负载

小技巧

对于多 stream 流水线中的延迟敏感型清理，优先使用 DeviceBox。 对于简单的单 stream 示例，DeviceBuffer 更简单，同步释放几乎零成本。

参数标量化#

当你编写一个接受 &[f32] 的 kernel 时，主机和设备对于 如何在内存中表示 Rust 切片并不一致——结构体布局在主机的 x86 ABI 和 NVPTX ABI 之间可能存在差异。cuda-oxide 通过在 kernel 边界处 标量化聚合类型来解决这一问题：将其分解为基础值，使 双方以相同方式解释这些值。

Kernel 参数类型

主机实际传递的内容

&[T]

ptr: *const T + len: u64

DisjointSlice<T>

ptr: *mut T + len: u64

T（标量）

直接传递 T

结构体 { a: u32, b: f32 }

一个 byval 值（整个结构体）

闭包（带有 N 个捕获）

一个 byval 值（整个结构体）

零大小类型

完全剥离

这就是为什么类型化 #[cuda_module] 方法接受 &DeviceBuffer<T> 作为 &[T] 参数，接受 &mut DeviceBuffer<T> 作为可写类切片参数。生成的 方法为你提取指针和长度。在 kernel 内部，编译器 从标量参数中重新构造切片结构体，因此你的 kernel 代码 看到的是正常的 &[T] 类型。

gpu-programming/images/scalarization-abi.svg

参数标量化：主机通过 ABI 边界将 Rust 切片作为 (ptr, len) 对传递。 设备 kernel 接收平坦的标量参数，编译器在 kernel 内部重新构造 原始的 Rust 类型。#

小技巧

标量化在正常 kernel 代码中完全不可见。你在签名中写 &[f32] 并像普通切片一样使用它。生成的启动方法和编译器处理其他一切。

DisjointSlice -- 安全的并行写入#

在 CUDA C++ 中，并行输出的标准模式是原始的 __global__ 指针， 每个线程对其进行索引。这本质上是不安全的——没有任何东西能阻止 两个线程写入同一位置。

cuda-oxide 提供了 DisjointSlice<T, IndexSpace> 作为安全替代。它 包装了一个可变切片，只允许通过其 IndexSpace 匹配的 ThreadIndex 进行写入，确保每个线程访问唯一的元素：

use cuda_device::{kernel, DisjointSlice};

#[kernel]
pub fn double(input: &[f32], mut out: DisjointSlice<f32>) {
    if let Some((out_elem, idx)) = out.get_mut_indexed() {
        *out_elem = input[idx.get()] * 2.0;
    }
}

  • get_mut_indexed() 是一次调用形式：它在单次调用中创建每线程见证并 将其解析为 &mut T。None 涵盖越界线程（例如对于 2D 场景 col >= ROW_STRIDE） 和越界索引两种情形。

  • 显式的两步形式 let idx = thread::index_1d(); out.get_mut(idx) 同样可用，适用于需要跨多个切片进行并行算术运算时。

  • 对于类似归约的模式，多个线程有意写入同一位置时， get_unchecked_mut（unsafe）提供了逃逸出口。

为什么 ThreadIndex 使这成为安全的#

DisjointSlice 安全性的关键是 ThreadIndex<'kernel, IndexSpace>—— 一个不透明的见证，没有公共构造函数。获取它的唯一方式 是通过受信任的索引函数，这些函数从硬件内置变量 （threadIdx、blockIdx、blockDim）中派生值：

let idx = thread::index_1d();          // ThreadIndex<'_, Index1D> -- 可以
let bad = ThreadIndex::new(42);        // 不存在 -- 私有构造函数

这之所以有效是因为 CUDA 的线程索引是硬件提供的统一值： block 中的每个线程从 GPU 的 warp 调度器接收到唯一的 threadIdx。 对于 1D grid 启动（仅使用 x 维度），从 blockIdx.x * blockDim.x + threadIdx.x 派生的全局索引在整个 grid 中 是唯一的。

该见证还是 !Send + !Sync + !Copy + !Clone 的，其 'kernel 生命周期借用自宏注入的栈局部作用域——因此线程不能将其 ThreadIndex 存放在共享内存中供邻居稍后取用，见证也不能在 kernel 体之外存活。 再加上 IndexSpace 参数（Index1D、Index2D<S>、Runtime2DIndex）， 类型系统在编译时也会拒绝不匹配的 2D 步幅——数据竞争风险变成一个类型错误。

共享内存#

共享内存是快速的片上内存，对 block 内的所有线程可见。 它是 block 内线程间通信和数据复用的主要工具，在速度和作用范围上 介于寄存器（每线程）和全局内存（所有线程）之间。

静态共享内存 -- SharedArray#

当大小在编译时已知时，使用 SharedArray<T, N>：

use cuda_device::{kernel, thread, SharedArray, DisjointSlice};

static mut TILE: SharedArray<f32, 256> = SharedArray::UNINIT;

#[kernel]
pub fn smem_example(input: &[f32], mut out: DisjointSlice<f32>) {
    let idx = thread::index_1d();
    let tid = thread::threadIdx_x() as usize;

    // 从全局内存加载到共享内存
    unsafe { TILE[tid] = input[idx.get()]; }
    thread::sync_threads();

    // 从共享内存读取邻居（比全局内存快得多）
    let neighbor = if tid > 0 {
        unsafe { TILE[tid - 1] }
    } else {
        0.0
    };
    thread::sync_threads();

    if let Some(out_elem) = out.get_mut(idx) {
        *out_elem = unsafe { TILE[tid] } + neighbor;
    }
}

每个 static mut SharedArray 映射到 PTX 中一个单独的 .shared 分配。 sync_threads() 屏障确保所有线程在完成写入之前，没有线程进行读取。

动态共享内存 -- DynamicSharedArray#

当大小取决于运行时参数时，使用 DynamicSharedArray<T> 并通过 LaunchConfig::shared_mem_bytes 指定分配大小：

use cuda_device::{kernel, thread, DynamicSharedArray, DisjointSlice};

#[kernel]
pub fn dynamic_smem_example(input: &[f32], mut out: DisjointSlice<f32>) {
    let smem = DynamicSharedArray::<f32>::get();
    let tid = thread::threadIdx_x() as usize;

    unsafe { *smem.add(tid) = input[thread::index_1d().get()]; }
    thread::sync_threads();
    // ... 使用 smem ...
}

在主机上，在启动时设置大小：

let config = LaunchConfig {
    grid_dim: (num_blocks, 1, 1),
    block_dim: (256, 1, 1),
    shared_mem_bytes: 256 * std::mem::size_of::<f32>() as u32,
};

多个动态数组可以通过使用 DynamicSharedArray::offset(byte_offset) 来 共享同一块分配，将其分区使用。

对齐#

类型

默认对齐

说明

SharedArray<T, N>

align_of::<T>()

标准对齐

SharedArray<T, N, 128>

128 字节

TMA 操作所需

DynamicSharedArray<T>

16 字节

与 nvcc 兼容的默认值

DynamicSharedArray<T, 128>

128 字节

TMA 所需

常见陷阱#

  • 缺少 sync_threads()： 在共享内存写入和读取之间没有屏障， 线程可能读取到过期或未初始化的数据。

  • 超出 SM 限制： 请求过多的共享内存会导致 CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES。请查看你所使用架构的限制。

参见

CUDA 编程指南 -- 共享内存 关于 bank 冲突、广播和最佳访问模式的详细信息。

Unified Memory 和 HMM#

默认情况下，GPU 在与 CPU 分离的地址空间中运行。GPU 不能解引用普通的主机指针——该地址根本无法映射到 GPU 的页表中。 因此传统的 CUDA 工作流需要在设备内存中显式分配，然后进行显式复制：

┌──────────────────┐              ┌──────────────────┐
│   CPU 内存       │    PCIe /    │   GPU 内存       │
│   （主机 DRAM）  │◄────────────►│   （设备 HBM）    │
│                  │   NVLink     │                  │
└──────────────────┘    复制      └──────────────────┘
    分离的地址空间 -- GPU 无法解引用主机指针

CUDA 提供了一些机制来放松这种限制，让 GPU 能够透明地访问 主机内存，代价是首次访问时的缺页延迟。

内存访问模式一览#

模式

GPU 可以访问什么

需要分配

首次访问成本

硬件要求

显式复制

仅设备内存

DeviceBuffer

无（数据已预先复制）

任何 CUDA GPU

Pinned（映射）

特定的主机缓冲区

cudaHostAlloc

高（每次访问约 10--20 µs）

任何 CUDA GPU

Unified Memory

托管分配

cudaMallocManaged

中（页面迁移）

Kepler+（sm_30+）

HMM

任意主机内存

无

中（缺页 + 获取）

Turing+ on Linux

cuda-oxide 主要使用显式复制（DeviceBuffer、DeviceBox）处理 大块数据，使用 HMM 处理非 move 闭包捕获和小型配置数据。

Unified Memory#

Unified Memory 是 CUDA 的托管内存分配器（cudaMallocManaged）。 生成的指针对 CPU 和 GPU 都有效——CUDA 运行时跟踪哪个处理器"拥有"每个页面， 并按需迁移。当 GPU 访问当前驻留在主机 DRAM 中的页面时，运行时在 kernel 读取数据之前透明地将其复制到设备内存。这种迁移对你的代码不可见， 但并非免费：从"错误"一侧的首次访问会产生缺页和通过互连进行 DMA 传输。 对同一页面的后续访问命中 GPU 的本地缓存。

cuda-oxide 目前不直接包装 cudaMallocManaged。对于托管内存工作流， 你需要通过原始绑定使用 CUDA 驱动 API。实际上， DeviceBuffer::from_host（显式复制）覆盖了大多数用例，并提供可预测的性能。

HMM（异构内存管理）#

HMM 是一个 Linux 内核特性，将 Unified Memory 的按需分页模型扩展到 所有系统内存——堆分配、mmap 区域，甚至栈变量。 启用 HMM 后，GPU 可以解引用任何有效的主机指针，无需特殊的 CUDA 分配器：

let factor = 5i32;                           // 普通栈变量
let scale = |x: i32| x * factor;            // 捕获 &factor（非 move）
cuda_launch! { kernel: scale, args: [...] }  // GPU 通过 HMM 读取 &factor

与 Unified Memory 不同，HMM 不需要特殊的分配 API——指针就是普通的主机地址。 当 ATS（地址转换服务）在 Grace Hopper 等硬件一致性平台上可用时， 它取代 HMM 并以缓存行粒度提供硬件一致性；HMM 会自动禁用。

缺页时会发生什么#

当 kernel 从一个页面不在设备内存中的地址加载数据时， 硬件和驱动程序协作将其获取：

  1. SM 对一个虚拟地址执行全局加载（ld.global）。

  2. GPU MMU 在 TLB 中查找该地址。如果未命中，遍历设备页表。

  3. 如果页表中没有映射，GPU 引发缺页。故障的 warp 暂停； 同一 SM 上的其他 warp 可以继续。

  4. CUDA 驱动故障处理程序确定页面的来源：

    • Unified Memory——CUDA 运行时识别托管分配并启动迁移。

    • HMM——Linux 内核的 HMM 层解析主机虚拟地址，固定主机页面， 然后迁移或创建远程映射。

  5. 通过 PCIe 或 NVLink 的 DMA 传输将页面从主机 DRAM 复制到设备 HBM。 GPU 内存控制器写入数据；主机内存控制器服务读取。

  6. GPU 页表被更新，TLB 被重新填充，warp 恢复执行。 页面现在是本地的并缓存在 L2 中；后续访问只需数百个周期。

步骤 5 的延迟取决于互连方式：

互连方式

带宽

缺页延迟

说明

PCIe 4.0 x16

~25 GB/s

~10--20 µs

大多数桌面/工作站 GPU

PCIe 5.0 x16

~50 GB/s

~5--15 µs

Ada Lovelace + 较新平台

NVLink 4.0

~900 GB/s

~1--5 µs

数据中心 GPU（H100、B100）

Grace Hopper C2C

~900 GB/s

<1 µs

硬件一致性 -- 使用 ATS，而非 HMM

由于故障以页面粒度（4 KB 或 2 MB）处理，单次故障可以满足许多线程。 Warp 级别的合并也有帮助：32 个线程读取连续的 4 字节元素， 最多触及一两个页面，而不是 32 个页面。在 PCIe 系统上， 单次故障的成本大致相当于一次小型的 cudaMemcpy——按需分页的 优势在于你只为你实际触及的页面付费。

cuda-oxide 如何使用 HMM#

cuda-oxide 以两种方式利用 HMM：

  1. 非 move 闭包捕获。 当非 move 闭包传递给 kernel 时， 捕获的变量保留在主机栈上，GPU 通过 HMM 指针访问它们。 这避免了复制 kernel 只读取一次或不常读取的数据。

  2. 动态布局的结构体 ABI。 cuda-oxide 在设备端匹配 Rust 实际的 结构体布局（包括 #[repr(Rust)] 字段重排），因此通过 HMM 访问的 主机结构体可以在没有 #[repr(C)] 或手动布局规范的情况下被正确读取。 编译器查询 rustc 获取字段偏移量并构建带有显式填充的匹配 LLVM 结构体类型。

HMM 系统要求#

要求

最低配置

GPU 架构

Turing（计算能力 7.5+）

Linux 内核

6.1.24+、6.2.11+ 或 6.3+

CUDA 驱动

535+ 配合 Open Kernel Modules

检查 HMM 是否在你的系统上激活：

nvidia-smi -q | grep Addressing
# Addressing Mode : HMM  ← HMM 已启用

何时使用 HMM vs 显式复制#

场景

推荐方法

由许多线程处理的大型数组

DeviceBuffer::from_host（显式复制）

小型只读配置数据

HMM（传递指针，让 GPU 缺页）

CPU 和 GPU 之间迭代共享的数据

双缓冲显式复制

原型/快速实验

HMM（最简单 -- 无需复制）

小技巧

HMM 是一种便利，而不是性能策略。对于带宽敏感的 kernel， 显式复制到设备内存总是更快，因为它们避免了缺页开销， 并使用完整的内存总线宽度。

参见

CUDA 编程指南 -- Unified Memory 和 NVIDIA 博客 -- 使用 HMM 简化 GPU 开发 关于页面迁移、预取和系统要求的完整详细信息。