CUDA 内存模型 Q&A

Q1：CUDA 中有哪几种内存类型？各自的特点、带宽、延迟是什么？

GPU 的存储层次和 CPU 类似，都是"越快越小、越慢越大"，但 GPU 拥有更丰富的片上可编程内存，这是高性能 CUDA 编程的核心所在。

CUDA 设备上的内存类型汇总如下：

内存类型 物理位置 访问范围 生命周期 速度 大小
寄存器（Register） SM 片上 单线程私有 Kernel 执行期间 最快 极小（每线程约 255 个 32 位寄存器）
共享内存（Shared Memory） SM 片上 线程块内所有线程 Kernel 执行期间 极快（低延迟） 每 SM 最大约 100~228 KB（因架构而异）
L1 Cache SM 片上（与共享内存共用物理空间） 自动缓存全局内存访问 自动管理 极快 与共享内存共享配额
L2 Cache 设备片上，全局共享 所有 SM 共享 自动管理 较快 数 MB 至数十 MB
全局内存（Global Memory） 设备 DRAM（显存） Grid 内所有线程 应用程序生命周期 高带宽但高延迟 数 GB（即整块显存）
常量内存（Constant Memory） 设备 DRAM + 专用缓存 Grid 内所有线程（只读） 应用程序生命周期 广播命中时快，否则慢 固定 64 KB
本地内存（Local Memory） 设备 DRAM（逻辑私有） 单线程私有 Kernel 执行期间 与全局内存相同（慢） 每线程最大 512 KB
纹理内存（Texture Memory） 设备 DRAM + 专用缓存 Grid 内所有线程（只读） 应用程序生命周期 空间局部性好时有缓存加速 —

性能感知： 寄存器和共享内存是 GPU 最快的存储，访问延迟只有几个时钟周期；全局内存延迟高达数百个周期，但带宽极高（如 H100 的 HBM3 带宽超过 3 TB/s），合理利用这一带宽是高性能 kernel 设计的关键。

Q2：各内存类型的用途详解

全局内存（Global Memory）

全局内存是 GPU 的主存储区域，相当于 CPU 系统中的 RAM。所有 kernel 线程都可以读写，容量等同于整块显存（数 GB）。

  • 持久性强：cudaMalloc 分配的全局内存在调用 cudaFree 或程序结束前一直存在。
  • 是 CPU 和 GPU 之间数据传输的中转区。
  • 所有线程可访问，但访问延迟高，必须通过合并访问（Coalescing）优化。
  • Kernel 的计算结果唯一的返回途径：kernel 没有返回值，计算结果只能写入全局内存，再由 CPU 读取。
1
2
3
4
5
// 在全局内存上分配和使用
float* devPtr;
cudaMalloc(&devPtr, N * sizeof(float));
// kernel 中直接通过指针访问
__global__ void kernel(float* data) { data[threadIdx.x] = 1.0f; }

共享内存（Shared Memory）

共享内存是 SM 上的片上可编程缓存，速度远高于全局内存，是同一线程块内线程协作的核心机制。

  • 物理上与 L1 缓存共享同一块硅片资源（统一数据缓存），两者的配比可以通过 API 调节。
  • 作用域仅限于单个线程块，线程块结束后数据消失。
  • 用途：线程间数据共享、缓存频繁访问的全局内存数据、用作用户管理的 scratchpad（便签本）。
  • 访问需要用 __syncthreads() 防止数据竞争。
1
2
3
4
5
6
__global__ void kernel(float* input, float* output) {
    __shared__ float tile[256];           // 静态分配
    tile[threadIdx.x] = input[threadIdx.x];
    __syncthreads();                       // 等所有线程写完再读
    output[threadIdx.x] = tile[255 - threadIdx.x];
}

动态分配方式：

1
2
3
4
// 启动时指定字节数
kernel<<<grid, block, sharedMemBytes>>>(args);
// kernel 内部
extern __shared__ float dynamicTile[];

寄存器（Register）

寄存器是最快的存储，每个线程独占，由编译器自动分配管理。

  • 存储 kernel 中的局部变量、循环计数器、临时计算中间值。
  • 无需程序员显式声明，编译器自动决定哪些变量放入寄存器。
  • 每个 SM 的寄存器总数固定（通常 64K 个 32 位寄存器），所有线程共享，用量直接影响 Occupancy。
  • 超出可用寄存器数量时会发生"寄存器溢出"（Register Spilling），详见 Q5。

常量内存（Constant Memory）

常量内存是只读的全局内存区域，配有专用的 SM 级常量缓存。

  • 容量固定为 64 KB，在 host 代码中初始化，kernel 中只能读取。
  • 当 Warp 内所有线程访问同一地址时，硬件以**广播（Broadcast）**机制一次性分发，效率极高。
  • 适合存储所有线程都会用到的只读参数、系数、查找表等。
  • 若不同线程访问不同地址（非广播），则串行化，性能退化为与全局内存相当。
1
2
3
4
5
6
7
8
// host 端声明并初始化
__constant__ float coeffs[4];
cudaMemcpyToSymbol(coeffs, h_coeffs, sizeof(h_coeffs));

// kernel 中直接使用
__global__ void compute(float* out) {
    out[threadIdx.x] = coeffs[0] * threadIdx.x + coeffs[1];
}

纹理内存（Texture Memory）

纹理内存在老旧 GPU 上因其空间局部性缓存机制而有性能优势。但在所有当前支持的 NVIDIA GPU 上，直接的 load/store 指令已经能达到同等效果，纹理内存不再具有性能优势。

如果你是新项目开发者，可以直接忽略纹理内存 API。如果维护老代码时遇到，只需知道它提供二维空间局部性缓存即可。

本地内存（Local Memory）

本地内存名字中有"local"，但物理上位于全局 DRAM，访问延迟与全局内存完全相同。"本地"只是指其逻辑作用域是线程私有的。

它是寄存器的溢出区，详见 Q5。

分布式共享内存（Distributed Shared Memory，CC ≥ 9.0）

从 Hopper 架构（CC 9.0）起，同一 Thread Block Cluster 内的所有线程块可以互相访问对方的共享内存，这个跨块可访问的共享内存空间称为分布式共享内存。其总容量等于 Cluster 内各线程块共享内存之和，为需要比单个 SM 更大协作范围的算法（如大规模直方图）提供了新的优化手段。

Q3：什么是内存合并访问（Memory Coalescing）？为什么重要？如何实现？

为什么重要

全局内存的访问以 32 字节为最小事务单位进行。当一个 Warp（32 个线程）同时请求全局内存时，硬件会将这 32 个线程的请求合并（coalesce）成尽量少的内存事务来完成。合并访问与非合并访问的性能差距可以高达 8 倍。

合并 vs 非合并对比

理想情况（完全合并）： 32 个连续线程访问连续的 4 字节数据，共需 128 字节，通过 4 次 32 字节事务完成，内存带宽利用率 100%。

最坏情况（完全不合并）： 32 个线程访问的地址彼此间隔 32 字节以上，每个线程触发一次独立的 32 字节事务，共需传输 1024 字节，但实际只用到 128 字节，带宽利用率仅 12.5%。

如何实现合并访问

核心原则：让连续线程 ID 的线程访问连续的内存地址。

1
2
3
4
5
6
7
8
9
10
11
// 合并访问：thread i 访问 data[i]
__global__ void good(float* data) {
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    data[i] = data[i] * 2.0f;   // 连续线程访问连续地址 ✓
}

// 非合并访问：每个线程跳跃访问，步长过大
__global__ void bad(float* data, int stride) {
    int i = (threadIdx.x + blockDim.x * blockIdx.x) * stride;
    data[i] = data[i] * 2.0f;   // stride 大时严重不合并 ✗
}

矩阵转置的典型案例： 朴素的转置 kernel 读操作合并但写操作不合并（列方向写入）。解决方案是用共享内存做"中转"：先按行（合并）读入共享内存，经 __syncthreads() 后，再按行（合并）写出到转置后的位置：

1
2
3
4
5
6
7
8
__global__ void smem_transpose(int m, float* a, float* c) {
    __shared__ float tile[32][32];
    // 合并读取：连续线程读连续列
    tile[threadIdx.x][threadIdx.y] = a[(tileRow + threadIdx.y)*m + (tileCol + threadIdx.x)];
    __syncthreads();
    // 合并写出：连续线程写连续行
    c[(tileCol + threadIdx.y)*m + (tileRow + threadIdx.x)] = tile[threadIdx.y][threadIdx.x];
}

合并访问的本质衡量标准： 最大化"实际使用的字节数 / 实际传输的字节数"之比。只要 Warp 内所有线程访问的地址落在尽量少的 32 字节段内，就能得到高利用率。

Q4：什么是 Shared Memory 的 Bank 冲突？如何避免？

共享内存的 Bank 结构

共享内存被划分为 32 个 Bank，连续的 32 位字依次映射到 Bank 0、Bank 1、…、Bank 31，然后循环。每个 Bank 每个时钟周期可以独立服务一次 32 位访问。

什么是 Bank 冲突

当同一个 Warp 内的多个线程在同一个时钟周期访问同一个 Bank 中的不同地址时，就发生了 Bank 冲突（Bank Conflict）。发生冲突时，这些访问必须被串行化执行，几路冲突就需要几倍的时间。

不会发生冲突的两种情况：

  1. 多个线程访问不同的 Bank（无论地址是否相邻）
  2. 多个线程访问同一 Bank 的同一地址——硬件以广播方式一次性分发

图解三种访问模式

  • 步长为 1（左图）：连续线程访问连续地址，每个线程命中不同 Bank，无冲突。
  • 步长为 2（中图）：每隔一个 Bank 访问，32 个线程中有 16 对线程命中相同 Bank，发生 2 路冲突，性能减半。
  • 步长为 3（右图）：3 与 32 互质，所有线程仍然命中不同 Bank，无冲突。
  • 左图：随机排列但每个线程命中不同 Bank，无冲突。
  • 中图：多个线程访问 Bank 5 的同一地址，触发广播，无冲突。
  • 右图：所有线程访问同一地址，广播，无冲突。

矩阵转置中的 Bank 冲突与规避

在上面 Q3 的共享内存转置代码中，声明 __shared__ float tile[32][32]，写操作 tile[threadIdx.x][threadIdx.y] 中，同一 Warp 内连续线程对应相同的 threadIdx.x，不同的 threadIdx.y，访问同一列——这 32 个元素在内存中每隔 32 个字，全部落在同一个 Bank，造成 32 路 Bank 冲突。

经典解决方法：在共享内存数组中加一列 padding（填充）：

1
__shared__ float tile[32][33];  // 加 1 列 padding，打破 Bank 对齐

加了一列 padding 后，同一列的元素不再对齐到相同 Bank，Bank 冲突消除，吞吐量大幅提升。

实践建议

  • 声明共享内存数组时，如果第一维大小是 2 的幂次（16、32、64 等），优先考虑在行末加 1 个 padding 元素。
  • 用 CUDA profiler（Nsight Compute）可以直接测量 shared memory bank conflict 次数，是优化的有力工具。

Q5：寄存器溢出（Register Spilling）是什么？有何影响？

什么是寄存器溢出

每个 SM 的寄存器文件大小有限（通常 64K 个 32 位寄存器），由当前驻留的所有线程共享。如果一个 kernel 中单个线程需要的寄存器数量超过了分配给它的上限，编译器会把多余的变量"溢出"到**本地内存（Local Memory）**中存放，这就是寄存器溢出（Register Spilling）。

本地内存在逻辑上是线程私有的，但物理上位于设备 DRAM（全局内存区域）。因此，一旦发生溢出，原本只需寄存器访问（接近零延迟）的操作，变成了访问全局内存（数百周期延迟），性能严重下降。

哪些情况会导致溢出

  • kernel 中声明了大量局部变量，超出寄存器预算
  • 使用了不能在编译期确定下标的局部数组（编译器无法将其放入寄存器）
  • 大型结构体或数组占用了过多寄存器空间
  • 使用 --maxrregcount 强制限制寄存器数量后，超出部分自动溢出

溢出的影响

  1. 直接影响： 访问溢出变量的延迟从接近 0 变为数百个时钟周期（等同于全局内存访问），严重拖慢 kernel 性能。
  2. 间接影响： 由于溢出到本地内存，额外占用显存带宽，可能影响其他访问的带宽。
  3. 正面效果（有时）： 限制寄存器用量可以让更多线程块同时驻留在 SM 上，提高 Occupancy。若 kernel 本身是延迟密集型（而非计算密集型），更高的 Occupancy 可能带来净性能提升，即使有溢出开销也值得。

如何诊断和优化

1
2
# 编译时查看寄存器用量
nvcc --resource-usage my_kernel.cu

编译器输出示例：

1
2
ptxas info: Used 48 registers, 0 bytes smem, 0 bytes lmem
ptxas info: Used 72 registers, 256 bytes lmem  ← lmem > 0 说明有溢出

优化方法：

  • 拆分大 kernel，减少单个 kernel 内的变量数量
  • 减少不必要的局部数组，改用全局内存或共享内存
  • 使用 __launch_bounds__(maxThreadsPerBlock, minBlocksPerSM) 给编译器提供提示，让其优化寄存器分配
  • 慎用 --maxrregcount，确保限制后的性能确实提升（通过 profiler 验证）

本地内存的访问模式有一个好消息：连续线程 ID 对本地内存中相同偏移量的访问是天然合并的，因此溢出的性能损失比随机全局内存访问要小，但仍远不如寄存器。

Q6：L1/L2 Cache 在 GPU 中如何工作？

整体结构

GPU 采用两级缓存结构：

  • L1 Cache（每 SM 独立）：物理上与共享内存共用同一块片上资源（称为统一数据缓存，Unified Data Cache）。两者的配比可以由程序员配置。
  • L2 Cache（全局共享）：位于设备上，所有 SM 共享，是 L1 和全局 DRAM 之间的缓冲层。容量通常为数 MB 至数十 MB，可通过 cudaDeviceProp.l2CacheSize 查询。

每个 SM 还有独立的常量缓存（Constant Cache），专用于缓存常量内存中的只读数据。

L1/共享内存比例配置

由于 L1 和共享内存物理共用，可以按需调整两者的比例：

1
2
3
4
5
6
7
8
// 推荐方式（hint，驱动可忽略）
cudaFuncSetAttribute(myKernel,
    cudaFuncAttributePreferredSharedMemoryCarveout, 50);  // 50% 给共享内存

// 旧方式（硬性要求，可能导致 kernel 串行化）
cudaFuncSetCacheConfig(myKernel, cudaFuncCachePreferL1);
// 或
cudaFuncSetCacheConfig(myKernel, cudaFuncCachePreferShared);
  • 如果 kernel 不使用共享内存，整个统一数据缓存都作为 L1 使用，自动获得最大 L1。
  • 如果 kernel 大量使用共享内存，L1 配额相应减少。

最佳实践： 优先使用 cudaFuncSetAttribute 设置 carveout 百分比，而不是 cudaFuncSetCacheConfig，前者是 hint，驱动可灵活处理，不会导致 kernel 串行；后者是硬性要求，相邻不同配置的 kernel 交替时会触发重配置开销。

L2 缓存的持久化控制（CC ≥ 8.0）

从 Ampere 架构（CC 8.0）起，可以将 L2 缓存的一部分划出来，专门用于持久化（persisting）访问——即某段数据在 L2 中享有高优先级驻留权，不会轻易被其他访问驱逐：

1
2
3
4
5
// 划出 75% 的 L2 用于持久化访问
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, device_id);
size_t size = min((size_t)(prop.l2CacheSize * 0.75), prop.persistingL2CacheMaxSize);
cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, size);

对于反复访问同一数据集（如大型权重矩阵、查找表）的 kernel，持久化 L2 可以显著降低全局内存带宽压力，提升性能。

GPU Cache 与 CPU Cache 的关键区别

GPU L1 缓存对于全局内存写操作**默认不是写回（write-back）**的——全局内存写通常直接穿透到 L2 或 DRAM。而 CPU 的 L1 通常是写回缓存。这是 GPU 内存一致性模型的一个重要差异，在多 kernel 访问共享数据时需要注意内存一致性。

Q7：什么是 Unified Memory？有何优缺点？

Unified Memory 的定义

统一内存（Unified Memory）是 CUDA 的一项内存管理特性，允许创建一块 CPU 和 GPU 都能直接访问的内存区域，程序员无需手动在两端之间拷贝数据——CUDA 运行时或底层硬件会自动在需要时将数据迁移到正确的位置。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
float* data;
cudaMallocManaged(&data, N * sizeof(float));  // 分配统一内存

// CPU 端初始化
for (int i = 0; i < N; i++) data[i] = i;

// GPU kernel 直接使用，无需 cudaMemcpy
kernel<<<grid, block>>>(data, N);
cudaDeviceSynchronize();

// CPU 端直接读取结果
printf("%f\n", data[0]);

cudaFree(data);

三种统一内存范式

根据系统硬件和操作系统不同，统一内存的能力分三个级别：

1. 受限统一内存（Limited Unified Memory）
主要出现在 Windows、WSL 和部分 Tegra 设备上：

  • 托管内存首先分配在 CPU 物理内存
  • GPU 开始执行时，整块托管内存迁移到 GPU
  • GPU 执行期间 CPU 不能访问托管内存
  • GPU 同步后数据迁移回 CPU
  • 不支持 GPU 显存超额分配

2. 完全统一内存（Full Unified Memory，软件一致性）
大多数 Linux 系统（PCIe 连接 GPU + HMM 支持）：

  • 支持 CPU 和 GPU 并发访问
  • 按页（page）粒度按需迁移
  • 支持 GPU 显存超额分配（oversubscription）
  • 可以使用 cudaMemAdvise 和 cudaMemPrefetchAsync 优化迁移行为

3. 完全统一内存（硬件一致性）
Grace Hopper、Grace Blackwell 等 NVLink C2C 连接架构（ATS 支持）：

  • 硬件级别维护一致性，无需软件页面迁移开销
  • GPU 可直接访问所有 CPU 内存，CPU 也可直接访问 GPU 内存
  • 性能最优，是真正的"零拷贝"架构

优缺点

优点：

  • 代码简洁，无需手写 cudaMemcpy，快速原型开发
  • 自动处理数据位置，降低 CPU/GPU 编程复杂度
  • 支持显存超额分配（Full 模式），可运行超过 GPU 内存大小的数据集
  • 可以通过 cudaMemAdvise/cudaMemPrefetchAsync 进一步优化

缺点：

  • 按需迁移存在运行时开销（缺页中断、数据迁移），首次访问延迟高
  • 性能上通常不如手动显式内存管理（optimal 时两者可以接近）
  • Windows/WSL 上限制较多，CPU 不能在 GPU 执行期间访问数据
  • 对迁移行为的控制粒度不如显式管理精细

性能建议： 即使使用统一内存，也应尽量使用 cudaMemPrefetchAsync 提前将数据迁移到目标设备，避免 kernel 执行时触发大量缺页中断：

1
2
3
// 将数据预取到 GPU，在 kernel 启动前完成迁移
cudaMemPrefetchAsync(data, N * sizeof(float), deviceId);
kernel<<<grid, block>>>(data, N);

Q8：cudaMalloc、cudaMallocManaged、cudaMallocHost 的区别？

这三个 API 是 CUDA 内存分配的核心，分别在不同的内存空间分配，适用于不同场景：

API 分配位置 谁能访问 用途
cudaMalloc GPU 显存（全局内存） 仅 GPU kernel 标准 GPU 计算缓冲区
cudaMallocManaged 统一内存（自动迁移） CPU + GPU 均可 快速开发、超额分配
cudaMallocHost CPU 物理内存（页锁定） 仅 CPU（可用于 DMA 传输） 高速 CPU↔GPU 数据传输缓冲区

cudaMalloc

在 GPU 显存上分配内存，是最基础的 GPU 内存分配方式。

1
2
3
4
5
6
7
8
9
10
float* devPtr;
cudaMalloc(&devPtr, N * sizeof(float));

// CPU 端无法直接访问 devPtr，必须通过 cudaMemcpy 传输数据
cudaMemcpy(devPtr, hostPtr, N * sizeof(float), cudaMemcpyHostToDevice);

// kernel 中可以直接使用
kernel<<<grid, block>>>(devPtr);

cudaFree(devPtr);  // 释放

特点：性能最可预期，完全在显存中，无迁移开销，是生产代码的首选。

cudaMallocManaged

分配统一内存，CPU 和 GPU 均可直接访问，CUDA 运行时自动管理数据迁移。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
float* ptr;
cudaMallocManaged(&ptr, N * sizeof(float));

// CPU 端直接读写
for (int i = 0; i < N; i++) ptr[i] = i;

// GPU kernel 直接使用
kernel<<<grid, block>>>(ptr);
cudaDeviceSynchronize();

// CPU 端直接读结果
printf("%f\n", ptr[0]);

cudaFree(ptr);

特点：开发最简便，适合算法原型验证；生产环境需配合 prefetch/advise 优化性能。

cudaMallocHost

在 CPU 上分配页锁定内存（Pinned Memory），是 CPU↔GPU 高速数据传输的标准方式。

1
2
3
4
5
6
7
8
9
10
float* hostPtr;
cudaMallocHost(&hostPtr, N * sizeof(float));  // 分配页锁定 CPU 内存

// 初始化数据
for (int i = 0; i < N; i++) hostPtr[i] = i;

// 传输到 GPU（速度比普通 malloc 更快，支持异步传输）
cudaMemcpy(devPtr, hostPtr, N * sizeof(float), cudaMemcpyHostToDevice);

cudaFreeHost(hostPtr);  // 释放（注意用 cudaFreeHost）

特点：是所有需要频繁 CPU↔GPU 数据传输的场景的最佳实践，详见 Q9。

选择建议

1
2
3
4
5
6
7
8
需要频繁 CPU↔GPU 数据传输？
  → CPU 端用 cudaMallocHost，GPU 端用 cudaMalloc + cudaMemcpy

快速原型开发，或数据超过 GPU 内存？
  → cudaMallocManaged

生产级 GPU 计算，数据主要在 GPU 端处理？
  → cudaMalloc

Q9：Pinned Memory（页锁定内存）是什么？为何能加速数据传输？

什么是 Pinned Memory

普通 CPU 内存（通过 malloc/new/mmap 分配）是可分页内存（Pageable Memory）——操作系统可以随时将其换出到磁盘（swap），也可以在物理内存中重新定位（内存整理）。

Pinned Memory（页锁定内存，也称固定内存）是用 cudaMallocHost 分配的 CPU 内存，操作系统保证永远不会将其换出到磁盘，也不会改变其物理地址。

为什么能加速数据传输

原因一：DMA 直接传输

GPU 通过 DMA（Direct Memory Access）控制器在 PCIe/NVLink 上传输数据，DMA 控制器只能访问物理地址固定的内存。

  • 普通可分页内存：在传输之前，CUDA 运行时必须先把数据临时拷贝到一块内部的 Pinned 缓冲区，再由 DMA 传输到 GPU。这是两段传输，有额外开销。
  • Pinned 内存：物理地址固定，DMA 可以直接访问，省去了中间拷贝，传输速度更快。

原因二：支持异步传输（cudaMemcpyAsync）

异步内存传输要求 CPU 端内存在传输期间保持有效且地址固定——这是 Pinned Memory 的天然属性。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
// 异步传输 + kernel 重叠执行（流水线）
cudaStream_t stream;
cudaStreamCreate(&stream);

float* pinnedSrc;
cudaMallocHost(&pinnedSrc, N * sizeof(float));

// 在 stream 中异步发起传输，CPU 立即继续执行
cudaMemcpyAsync(devPtr, pinnedSrc, N * sizeof(float),
                cudaMemcpyHostToDevice, stream);

// 与传输同时，做其他 CPU 工作或启动另一个 kernel
doOtherWork();

cudaStreamSynchronize(stream);  // 等传输完成

普通可分页内存无法用于 cudaMemcpyAsync，因为操作系统随时可能移动数据的物理位置，DMA 传输中途就会读到错误数据。

原因三：可以被映射供 GPU 直接访问

Pinned Memory 可以通过 cudaHostAlloc 的 cudaHostAllocMapped 标志被映射到 GPU 的地址空间，kernel 可以直接通过指针读写 CPU 端的 Pinned Memory（数据走 PCIe/NVLink）。虽然带宽低于显存访问，但在数据量小、只需访问一次时可以省去显式传输步骤。

API 汇总

1
2
3
4
5
6
7
8
9
10
11
12
13
// 分配页锁定内存（最常用）
cudaMallocHost(&ptr, size);

// 带标志的分配（更灵活）
cudaHostAlloc(&ptr, size, cudaHostAllocDefault);     // 同 cudaMallocHost
cudaHostAlloc(&ptr, size, cudaHostAllocMapped);      // 映射到 GPU 地址空间
cudaHostAlloc(&ptr, size, cudaHostAllocWriteCombined); // 写合并，提升 CPU→GPU 传输速度，但 CPU 读慢

// 对已有内存进行页锁定（适合第三方库分配的内存）
cudaHostRegister(existingPtr, size, cudaHostRegisterMapped);

// 释放
cudaFreeHost(ptr);

注意事项

  • 不要过量分配 Pinned Memory。页锁定的内存不能被操作系统换出，会长期占用物理内存，分配过多会导致系统整体内存压力增大，影响其他进程。
  • 最佳实践：只对用于 CPU↔GPU 传输的缓冲区使用 Pinned Memory，其他 CPU 端数据使用普通 malloc 即可。
  • 在配备 HMM 或 ATS 的系统上，所有 CPU 内存对 GPU 直接可见，Pinned Memory 的传输优势会有所改变，具体以 NVIDIA 官方文档为准。