CUDA内存空间全景图

CUDA 的程序员之所以与 CPU 程序员有着截然不同的思维模式，根源就在内存模型上。CPU 程序面对的是一个统一的大内存池，而 CUDA 程序面对的则是一个由多种内存空间构成的异构体系——它们散布在芯片的不同物理位置，拥有迥异的容量、速度和生命周期。在写下第一行 cudaMemcpy 之前，必须先建立这张全景图，否则后续的每一次优化都将无从谈起。

容量与速度：一场金字塔式的博弈

CUDA 内存空间并非平权设计。从硬件物理位置与访问延迟来看，它们天然地呈现为一座金字塔：

• 寄存器（Register）：位于 SM 内部的超高速存储单元，容量极小（每 SM 约 256KB，均分给线程后每线程通常只有几十到上百个 32 位寄存器），但访问延迟为零（与指令执行同步），生命周期与线程一致。变量声明后未加任何修饰符的局部变量，若编译器判定可放入寄存器则放入寄存器。
• 共享内存（Shared Memory）：同样位于片内，是 SM 上的一块可编程管理的高速存储，每 SM 约 48KB 到 228KB（取决于架构和显式配置）。访问延迟约为全局内存的 1/20 到 1/50，带宽极高。它属于线程块私有，块内所有线程可读写同一份数据——这正是线程协作的基础设施。
• L1 缓存（L1 Cache）：位于片内，由硬件自动管理，用于缓存全局内存和局部内存的访问。共享内存在物理上常与 L1 缓存共享同一块 SRAM，这也是为什么两者容量此消彼长的原因。
• 全局内存（Global Memory）：即显存（GPU 板载 DRAM），容量最大（数 GB 到数十 GB），是所有线程（包括不同的 kernel 启动）都能访问的唯一空间。但访问延迟高达 400-800 个时钟周期，是所有内存空间中速度最慢的常规空间。所有从主机传输到设备的数据都暂时栖身于此。
• 本地内存（Local Memory）：一个容易误导的名字。它逻辑上是线程私有的，但物理上存放在全局内存（DRAM）中。当寄存器被耗尽（溢出）、或访问动态索引的局部数组时，编译器会将数据放在本地内存。代价是访问延迟与全局内存相当。
• 常量内存（Constant Memory）：容量固定 64KB，位于片外 DRAM，但拥有专属的缓存（常量缓存）。当所有线程访问同一地址时，缓存广播机制让读取速度逼近寄存器；若各线程访问不同地址，则退化为普通全局内存读取。数据在 kernel 运行期间只读。
• 纹理内存（Texture Memory）：同样是只读的片外空间，专为图形学设计的缓存和硬件插值功能，适合二维空间局部性强的访问模式（如图像处理）。对通用计算而言，其地位已被现代架构的只读缓存部分替代，但在特定场景仍有价值。

作用域：谁看得见谁

这八种内存空间的可见性呈现严格的层级关系，几乎所有编程错误都源于对作用域的误判：

内存空间	作用域	可见性
寄存器	单个线程	仅该线程
本地内存	单个线程	仅该线程（物理上存储在全局内存，逻辑上私有）
共享内存	线程块	块内所有线程
全局内存	网格（整个设备）	所有线程 + 主机
常量内存	网格（整个设备）	所有线程可读（只读）
纹理内存	网格（整个设备）	所有线程可读（只读）
固定内存	主机	主机 + 设备（通过零拷贝或 DMA）
统一内存	主机 + 设备	主机与设备（由运行时自动管理迁移）

生命周期：谁创建，谁销毁

每种内存空间的生与死遵循着不同的规则，这决定了数据在不同 kernel 启动之间能否存活。快速一览：

内存空间	生命周期	分配/释放方式
寄存器、本地内存	随线程诞生与消亡	编译器自动管理
共享内存	随线程块创建/销毁	静态声明或动态配置
全局、常量、纹理内存	显式管理，跨越内核启动	cudaMalloc / cudaFree
固定内存	主机 API 控制	cudaMallocHost / cudaFreeHost
统一内存	运行时托管	cudaMallocManaged / cudaFree

几个值得注意的细节：

• 寄存器与本地内存：随线程的诞生而诞生，随线程的消亡而消亡。线程执行完毕，这些数据便灰飞烟灭。
• 共享内存：生命周期与线程块绑定。块创建时分配，块销毁时回收。注意：块内的共享内存在同一块内的不同 kernel 启动之间不保留（除非使用 CUDA 12 的动态共享内存持久化特性，那是后话）。
• 全局内存、常量内存、纹理内存：生命周期由 cudaMalloc / cudaFree（或 cudaMemcpyToSymbol 等）显式管理。它们跨越 kernel 启动而存在，直到主机显式释放，或程序退出。这是数据在不同 kernel 之间传递的通道。

CPU 直觉的陷阱：CPU 程序员常默认"分配的内存只要没释放就一直有效"，在 CUDA 中这条直觉在共享内存和寄存器上完全失效——它们是硬件资源，不是软件管理的堆。

---

这张全景图回答了一个核心问题：选择哪种内存，本质上是在容量、速度、作用域和生命周期之间做权衡。全局内存在容量上占优但速度垫底，共享内存在速度上接近寄存器但容量受限，寄存器最快却无处藏身。理解了这些边界条件，接下来就可以开始动手实践——通过 cudaMalloc 与 cudaMemcpy 打通主机与设备之间的数据传输通道，这是所有 CUDA 程序的第一个性能瓶颈，也是第一个优化突破口。

显存分配与拷贝：cudaMalloc/cudaMemcpy

全景图已经告诉我们，全局内存是主机与设备之间交换数据的"码头"。那么，如何在这座码头上装卸货物？答案就是这一节的主角：cudaMalloc 与 cudaMemcpy。这对 API 构成了所有 CUDA 程序中最基础的数据通路，几乎所有需要将数据从 CPU 端送入 GPU 的场景，都从这两个函数开始。

显存分配：cudaMalloc

CPU 端用 malloc 分配堆内存，GPU 端则用 cudaMalloc 分配全局内存。两者的接口形式极为相似：

cudaError_t cudaMalloc(void** devPtr, size_t size);

第一个参数是指向指针的指针，因为 cudaMalloc 需要修改调用者传入的指针变量，使其指向新分配的显存地址；第二个参数指定分配的字节数。一个值得注意的细节是：cudaMalloc 分配的是设备端全局内存，这片内存在主机端代码中无法直接解引用——你拿到的只是一个"远程地址"。

来看一个完整的分配与释放流程：

float *d_data = nullptr;      // 设备端指针，初始化为空
size_t bytes = 1024 * sizeof(float);  // 1024 个 float

// 分配显存；返回 cudaError_t 用于错误检查
cudaError_t err = cudaMalloc(&d_data, bytes);
if (err != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed: %s\n", cudaGetErrorString(err));
    exit(EXIT_FAILURE);
}

// ... 使用 d_data（需要通过拷贝或内核写入） ...

cudaFree(d_data);   // 释放显存，与 cudaMalloc 配对使用

每个 cudaMalloc 都必须对应一个 cudaFree，否则会造成显存泄漏——这一点与 malloc/free 的纪律完全一致。不过在 CUDA 程序中，显存泄漏往往更难排查，因为 nvcc 编译期不会给出任何警告，而运行时错误又可能在数小时之后才浮出水面。

双向拷贝：cudaMemcpy 与方向参数

有了分配好的显存，下一步就是数据的搬运。cudaMemcpy 是 CUDA 中最核心的数据传输函数，它的签名如下：

cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, 
                       cudaMemcpyKind kind);

前两个参数分别指定目标地址和源地址，第三个参数是拷贝的字节数，第四个参数 kind 决定了拷贝的方向。cudaMemcpyKind 共有四种取值：

枚举值	方向	典型场景
cudaMemcpyHostToDevice	主机 → 设备	将输入数据从 CPU 内存送入显存
cudaMemcpyDeviceToHost	设备 → 主机	将计算结果从显存取回 CPU 内存
cudaMemcpyDeviceToDevice	设备 → 设备	显存内部搬运（如交换两块缓冲区）
cudaMemcpyHostToHost	主机 → 主机	等价于 memcpy，几乎不用

最容易犯的错误是搞反 dst 和 src 的顺序。一个实用的记忆技巧：cudaMemcpy 的参数顺序与 memcpy 完全一致，永远是"目标在前，源头在后"。方向参数 kind 则是告诉 CUDA 运行时，两个地址分别位于哪一侧，以便正确配置 DMA 引擎。

阻塞式拷贝的同步语义

cudaMemcpy 是一个同步（阻塞）函数——这是理解其行为的关键。当主机线程调用 cudaMemcpy 时，它会阻塞等待拷贝完成才返回。具体来说：

• 对于 HostToDevice 方向的拷贝，函数返回时，数据已经安全地驻留在显存中，主机端可以立即释放源缓冲区或复用它来准备下一批数据；
• 对于 DeviceToHost 方向的拷贝，函数返回时，数据已经完整地拷贝到主机内存中，设备端的计算（如果之前有启动过内核）也已经完成。

这意味着 cudaMemcpy 天然地充当了同步屏障：它隐式地等待在此之前启动的所有内核执行完毕，并确保后续操作看到最新数据。这种语义简化了编程模型——你不需要手动管理同步，代价则是在数据搬运期间 GPU 计算单元可能处于空闲状态（因为等待数据传输完成）。数据传输与计算无法重叠，这正是后续使用 cudaMemcpyAsync + 流（Stream）要解决的痛点。此外，与 cudaMalloc 一样，cudaMemcpy 的返回值也应当检查，这里出于简洁省略，但生产代码中务必处理。

完整示例：向量加法

将以上的知识串起来，来看一个完整的向量加法程序。它演示了从分配、拷贝、内核执行到结果回传的完整数据流：

#include <stdio.h>
#include <cuda_runtime.h>
#include <stdlib.h>

// 简单的向量加法内核：每个线程负责计算一个元素
__global__ void vecAdd(const float *a, const float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];   // 从全局内存读取、写入
    }
}

int main() {
    // ========== 1. 主机端准备数据 ==========
    const int N = 1024;
    size_t bytes = N * sizeof(float);
    float *h_a = (float*)malloc(bytes);
    float *h_b = (float*)malloc(bytes);
    float *h_c = (float*)malloc(bytes);   // 存放结果

    for (int i = 0; i < N; i++) {
        h_a[i] = i * 1.0f;
        h_b[i] = 1.0f;
    }

    // ========== 2. 设备端分配显存 ==========
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // ========== 3. 主机 → 设备：拷贝输入数据 ==========
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // ========== 4. 启动内核：每个 block 256 线程 ==========
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;  // 向上取整
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, N);

    // ========== 5. 设备 → 主机：取回结果 ==========
    // cudaMemcpy 阻塞等待内核完成，此时结果已就绪
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    // ========== 6. 验证与清理 ==========
    for (int i = 0; i < N; i++) {
        if (h_c[i] != h_a[i] + h_b[i]) {
            printf("Mismatch at index %d\n", i);
            exit(EXIT_FAILURE);
        }
    }
    printf("All %d elements matched!\n", N);

    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    free(h_a); free(h_b); free(h_c);
    return 0;
}

编译命令为 nvcc -o vecAdd vecAdd.cu，运行输出应为 All 1024 elements matched!。这个程序中，步骤 3 的两次拷贝和步骤 5 的一次回传都是阻塞式的：步骤 3 完成后才启动内核，内核完成后步骤 5 才返回。整条数据通路是严格串行的。

数据通路的边界

至此，一条基本的数据通路已经打通：malloc 分配主机内存 → cudaMalloc 分配显存 → cudaMemcpy 双向搬运 → cudaFree 回收。这套流程是 CUDA 编程的"基本功"，适用面最广，也是绝大多数入门教程采用的模式。但它有两个天然局限：拷贝的同步阻塞（计算与传输无法重叠）和主机的页式内存（传输走 DMA 时需要额外的隐式拷贝）。前者可以通过流与异步拷贝解决，后者则引入了固定内存（Pinned Memory）——这是下一节的主角。

固定内存与零拷贝

前两节建立的数据通路——cudaMalloc 分配显存、cudaMemcpy 搬运数据——已经能支撑起绝大多数 CUDA 程序的正确运行。但细心的读者可能会问：cudaMemcpy 是一个同步函数，它阻塞等待拷贝完成。那么在这段时间里，数据究竟经历了怎样的旅程？理解了这个问题，就能明白为什么有一种内存分配方式能让拷贝速度提升近一倍。

分页内存与固定内存：数据传输的隐藏瓶颈

默认情况下，malloc 分配的主机内存是分页内存（Pageable Memory）。操作系统的虚拟内存机制允许这些页面在物理内存与磁盘之间换入换出。当 cudaMemcpy 需要将数据从分页内存拷贝到显存时，CUDA 驱动无法直接通过 DMA（Direct Memory Access） 访问这些页面——因为它们在物理内存中的位置可能随时发生变化。

为了解决这个问题，CUDA 驱动必须执行一次额外的内部拷贝：先将分页内存中的数据复制到一块临时的固定内存（Pinned Memory）（也称为页锁定内存）中，再由 DMA 引擎将数据从这块固定内存传输到设备端。整个流程如下图所示：

这额外的第一步带来了两个代价：其一是时间开销，数据多了一次拷贝，整体传输时间几乎翻倍；其二是 CPU 参与度，第一步需要 CPU 介入，使得整个过程中 CPU 无法完全腾出来做其他事情。

解法也很直接——既然驱动总要找一块固定内存来中转，不如让用户直接绕过 malloc，用 cudaMallocHost 来分配内存：

float *h_data;
cudaMallocHost(&h_data, nbytes);  // 分配固定内存
// ... 使用 h_data 填充数据 ...
cudaMemcpy(d_data, h_data, nbytes, cudaMemcpyHostToDevice);
cudaFreeHost(h_data);  // 注意用 cudaFreeHost 而非 free

cudaMallocHost 分配的页面被锁定在物理内存中，操作系统不会将其换出。CUDA 驱动可以直接向 DMA 引擎提供这些页面的物理地址，省去了中转拷贝。在实测中，对于大规模数据传输（数百 MB 以上），使用固定内存通常能获得 20% 到 100% 的带宽提升，具体幅度取决于 PCIe 版本、数据规模和系统负载。

需要强调的是，固定内存的分配成本远高于普通内存：调用 cudaMallocHost 时，驱动需要锁定物理页并建立映射关系。正确做法是一次性分配、反复使用，而非在循环内部反复调用。

cudaHostAlloc：更精细的控制

cudaHostAlloc 是 cudaMallocHost 的超集，额外的标志位提供了更精细的控制：

标志	作用	适用场景
cudaHostAllocDefault	行为等价于 cudaMallocHost	一般用途
cudaHostAllocWriteCombined	分配**写合并（Write-Combined）**内存	仅用于设备读取的主机内存
cudaHostAllocMapped	将固定内存映射到设备地址空间	零拷贝内存
cudaHostAllocPortable	允许该内存被所有 CUDA 上下文访问	多 GPU 或多线程场景

其中 cudaHostAllocWriteCombined 值得留意。写合并内存使用了一种特殊的缓存策略，专为 PCIe 总线上的批量写入优化。它有两个特点：其一，写入操作不会被 CPU 缓存，因此 CPU 端读取这种内存的性能极差（可能比普通内存慢一个数量级）；其二，设备端读取这种内存时不需要维护缓存一致性，带宽更高。因此它只适合"CPU 写、GPU 读"的单向传输场景。

零拷贝内存：取消显式的拷贝

将 cudaHostAlloc 的 cudaHostAllocMapped 标志与固定内存结合，就得到了零拷贝内存（Zero-Copy Memory）。所谓零拷贝，并非数据传输不需要时间，而是不需要程序员显式调用 cudaMemcpy——设备端通过映射直接读取主机内存。

使用方式如下：

float *h_data, *d_ptr;
cudaHostAlloc(&h_data, nbytes, cudaHostAllocMapped);  // 分配映射的固定内存
cudaHostGetDevicePointer(&d_ptr, h_data, 0);          // 获取设备端指针
// 在 kernel 中直接使用 d_ptr 读写，无需 cudaMemcpy
kernel<<<grid, block>>>(d_ptr);
cudaFreeHost(h_data);

设备端访问 d_ptr 时，每次读写都会通过 PCIe 总线实时访问主机内存。这意味着：

• 零拷贝消除了显式的 cudaMemcpy 调用，代码更简洁
• 零拷贝并没有消除数据传输——每次内核访问都会触发一次 PCIe 事务
• 对于小数据量、多次访问的场合（如参数传递），零拷贝的延迟远低于"拷贝一次 + 访问多次"
• 对于大数据量、高频访问的场合（如大规模矩阵运算），零拷贝的性能会急剧恶化——每次访问都是一次 PCIe 往返

另一个关键约束是：零拷贝内存要求主机端必须是固定内存，否则无法建立设备映射。

一张表总结：何时使用哪种方案

场景	推荐方案	理由
大块数据、一次拷贝后被反复使用	cudaMalloc + cudaMemcpy	传输成本摊薄到每次内核访问中
大块数据、需要最高传输带宽	cudaMallocHost（固定内存）	省去中转拷贝，带宽最高
小块数据、频繁与设备交互	cudaHostAllocMapped（零拷贝）	省去显式拷贝，延迟最低
CPU 只写、GPU 只读的大数据流	cudaHostAllocWriteCombined	写合并优化单向带宽

这四种方案构成了主机与设备之间数据传输的完整工具箱。但无论是固定内存还是零拷贝，它们优化的都是"数据进出的速度"——真正决定内核性能的，是数据进入芯片之后如何被高效地利用。接下来的视角将转向统一内存——一种试图彻底消除手动搬运的编程模型。

统一内存与按需迁移

固定内存的确显著提升了 cudaMemcpy 的带宽，但仍要求程序员手工管理两块内存的布局、分配时机与拷贝方向。当数据规模增长、代码路径分支增多时，这种显式搬运的负担会迅速膨胀。统一内存（Unified Memory）正是为解决这一痛点而生的编程模型——它试图让 GPU 程序员重新体验 CPU 编程中"一个指针走天下"的便利。

简化编程：一个指针访问两种内存

cudaMallocManaged 是进入统一内存世界的大门。它分配的内存同时映射到主机与设备地址空间，返回的指针既可以在 CPU 函数中解引用，也可以在核函数中直接访问：

float* data;
cudaMallocManaged(&data, N * sizeof(float));

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

// GPU 端直接读取
kernel<<<blocks, threads>>>(data);
cudaDeviceSynchronize();

// 回到 CPU 端直接读取结果
float sum = 0;
for (int i = 0; i < N; i++) sum += data[i];

从编程体验来看，这与纯 CPU 代码几乎无异——没有 cudaMemcpy，没有 dst 与 src 的区分，更不存在搞反方向的低级错误。对于原型验证、数据结构复杂的应用（如链表、树）以及难以预判数据访问模式的算法，统一内存能显著降低心智负担。

但"方便"并非免费。要理解其代价，需要先了解它背后的机制。

页面错误与按需迁移：懒加载的代价

统一内存的核心机制是按需迁移。当 CPU 访问 cudaMallocManaged 分配的内存时，系统将该页面驻留在主机侧；当 GPU 某个线程块首次访问该页面时，硬件触发页面错误（Page Fault），驱动程序将对应页面从主机内存迁移至设备显存，随后核函数继续执行。这个过程完全由运行时自动完成，对程序员透明：

这种机制在 Pascal 架构（Compute Capability 6.0）之后成为硬件特性，也带来了实质性的性能代价。每一次页面错误的处理需要经历：陷入驱动 → 定位页面 → 传输数据 → 更新页表 → 恢复执行，延迟通常在数十微秒量级。相比之下，一次常规的 cudaMemcpy 虽然也要传输数据，但它是批量、连续的，且没有页面错误处理的开销。因此，如果核函数反复触发页面迁移，统一内存的性能可能比显式拷贝差数倍。

并发核函数访问：数据局部性的新问题

统一内存在多核函数并发场景下还有一个容易被忽视的细节。考虑两个并发执行的核函数，分别访问同一块统一内存的不同区域，硬件会尝试将这些页面分布到两个流各自的物理位置。若访问模式不幸重叠，页面会在两个 GPU 引擎之间来回迁移，形成抖动（Thrashing）。一个典型的反面模式是：

cudaStream_t s1, s2;
cudaStreamCreate(&s1); cudaStreamCreate(&s2);
kernelA<<<blocks, threads, 0, s1>>>(data, offsetA);
kernelB<<<blocks, threads, 0, s2>>>(data, offsetB);

当 offsetA 与 offsetB 指向的区间在物理页面上交错时，两个流会争夺同一批页面的所有权。正确做法是让每个流访问尽量互不重叠的大块连续区域，或使用 cudaMemPrefetchAsync 主动将页面预迁到目标设备：

// 在核函数启动前，明确告诉驱动：这些页面将主要被 GPU 使用
cudaMemPrefetchAsync(data, N * sizeof(float), deviceId, stream);

这一 API 相当于手动的"页面预热"，能大幅减少运行时页面错误的次数。

性能注意事项：何时选择统一内存

综合以上讨论，可以给出统一内存的适用性判断：

场景	显式拷贝	统一内存
数据结构简单、访问模式固定	推荐	可用
指针复杂（链表、树）	难以实现	推荐
数据量巨大且只访问一次	推荐	不推荐
原型开发/教学示例	可用	推荐
多流并发且页面交错	推荐	需配合 cudaMemPrefetchAsync

核心结论是：统一内存降低了编程复杂度，但将性能优化的责任转交给了运行时与页式迁移机制。它在数据复用率高、访问模式难以预判的场景最能发挥价值；而在数据一次性流过、带宽敏感的生产级代码中，手工 cudaMemcpy 的确定性优势依然不可替代。

从全景图的视角看，统一内存是一条新的"河流"——它将主机与设备的内存空间连成了一片，但水流的方向与速度不再由程序员完全掌控。理解了它的运行机制，就能在便利与性能之间做出有依据的选择。接下来我们将从数据传输的宏观层面下沉到核函数内部，看看那个容量最小却速度最快的内存空间——寄存器——是如何在程序运行时被精细分配的。

数据传输优化策略

前三节构建的主机-设备数据通路——cudaMalloc 分配、cudaMemcpy 搬运、固定内存加速、统一内存简化——解决的都是"如何正确传输"的问题。但正如全景图所揭示的，PCIe 总线是 CPU 与 GPU 之间的唯一物理通道，它的带宽（PCIe 3.0 x16 约为 16 GB/s）相比设备端的显存带宽（数百 GB/s）低了一个数量级。这意味着，哪怕 cudaMemcpy 的参数写得再正确、固定内存用得再到位，如果传输策略本身是低效的，性能天花板依然牢牢锁死在 PCIe 的物理极限上。本节从三个层面讨论：如何压缩传输开销、如何减少传输次数、以及如何让传输与计算重叠。

传输开销：隐藏的固定成本

一次 cudaMemcpy 的开销远不止"按字节搬运"那么简单。它包含三个组成部分：

开销来源	说明	量级
调用开销	API 调用进入驱动、参数校验、上下文切换	微秒级（~5-10μs）
启动开销	DMA 引擎初始化、PCIe 链路握手	微秒级（~5-20μs）
传输本身	数据量 ÷ PCIe 有效带宽	与数据量成正比

前两项是固定成本——无论传输 1 字节还是 1 GB，这部分时间几乎不变。当传输数据量很小时，固定成本占比极高，有效带宽极低。一个具体的感受：传输 1 KB 数据，假设固定开销 10μs，PCIe 带宽 12 GB/s，则传输本身约 0.08μs，总耗时约 10.08μs，有效带宽仅约 100 MB/s——不到峰值带宽的 1%。

这一观察直接引出第一条优化原则：不要传输小数据块。如果程序需要分多次传输总量为 100 KB 的数据，与其发起 100 次 1 KB 的传输（总耗时约 1 ms），不如合并为一次 100 KB 的传输（总耗时约 18μs），差距超过 50 倍。做法是将分散在多个小数组中的数据在主机端打包进一个连续的缓冲区，用一次 cudaMemcpy 完成搬运。这就是批量传输（Batched Transfer）的核心思想。批量传输的另一层含义是：尽量复用已经分配好的显存。cudaMalloc 和 cudaFree 本身也有不小的开销（涉及驱动级的显存管理），频繁分配释放往往比多传几 MB 数据还要昂贵。正确的做法是一次分配、反复使用——这也是固定内存章节中提到的原则在全局内存上的延伸。

数据复用：从根源上减少传输

批量传输优化的是"单次传输的效率"，但真正治本的方向是减少传输的总字节数。这引出一个关键洞察：GPU 的瓶颈往往不是算力，而是数据供给。如果数据只在主机端被访问一次、在设备端也只用一次，那么传输成本就是不可避免的净开销。但如果数据在设备端被多次复用，每次复用的成本就被摊薄了。

数据复用的典型模式是计算密集型的多阶段流水线。例如矩阵乘法或卷积运算，输入数据一旦从主机传输到显存，就会被核函数反复读取。此时正确的策略是：尽可能在设备端完成所有中间步骤，只传输入和最终输出。每减少一次主机-设备间的中间结果往返，就节省了一次完整的 PCIe 往返延迟。

另一个常见的反模式是在循环体内传输数据：

// 反模式：每个循环迭代都发起一次传输
for (int i = 0; i < 1000; i++) {
    cudaMemcpy(d_input, h_input + i * N, N * sizeof(float), cudaMemcpyHostToDevice);
    kernel<<<blocks, threads>>>(d_input, d_output);
    cudaMemcpy(h_result + i * N, d_output, N * sizeof(float), cudaMemcpyDeviceToHost);
}

这段代码有 2000 次 PCIe 传输，实际上完全可以先拼接所有输入为一个大数组，一次传入、一次传出。不仅仅是因为调用次数多，更因为每次传输之间都夹着一次核函数启动，传输的固定开销与核函数启动的延迟叠加，共同拖垮了吞吐。

需要澄清的是，数据复用不是指"多传几次"，而是指"让同一份数据在设备端发挥更大价值"。理解这一点的关键指标是计算传输比（Compute-to-Transfer Ratio）：设备端每接收 1 字节数据所执行的浮点运算次数。这个比值越高，程序对 PCIe 带宽的依赖就越低，越有可能逼近 GPU 的算力上限。一个极端的例子是图像卷积：每读入一个像素块，可能要进行上百次乘加运算；而一个单纯的数组求和，计算传输比接近 1:1，注定受限于 PCIe 带宽。

流：让传输与计算重叠

前两节优化的是"减少传输量"和"提高单次效率"，但传输期间 GPU 的 SM（流多处理器）是否无事可做？答案是肯定的。默认情况下，所有 CUDA 操作（传输、核函数）都在**默认流（Default Stream）**中顺序执行——核函数等待传输完成，传输等待上一核函数结束。这种串行依赖造成了严重的硬件闲置：DMA 引擎搬运数据时，SM 空闲；SM 计算时，PCIe 空闲。

流（Stream） 正是用来打破这种串行的机制。流代表一组按序执行的操作序列，但不同流之间的操作可以并发执行。将传输放入一个流、将核函数放入另一个流，只要它们之间没有数据依赖，驱动就会尝试让 cudaMemcpyAsync 与核函数同时进行。例如，在深度学习训练中，一个 mini-batch 的数据传输可以与上一个 mini-batch 的梯度计算重叠，从而掩盖 PCIe 延迟。

// 创建两个流，使数据预取与计算重叠
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 流1：传输第 i+1 批数据
cudaMemcpyAsync(d_input_next, h_input_next, bytes, cudaMemcpyHostToDevice, stream1);

// 流2：计算第 i 批数据（假设 d_input_cur 已就绪）
kernel<<<blocks, threads, 0, stream2>>>(d_input_cur, d_output_cur);

这里需要区分 cudaMemcpy 与 cudaMemcpyAsync：前者是同步的，即使指定了流也会阻塞主机；后者才是真正的异步版本，通过流参数将传输操作交给 DMA 引擎，主机线程立即返回。注意，若要实现传输与计算的真正重叠，主机端的数据缓冲区必须是固定内存——因为异步传输需要 DMA 直接从主机内存读取数据，而分页内存的页面在物理内存中可能被换出，DMA 无法处理这种情况。

流的引入，为数据传输打开了第三维度的优化空间：延迟隐藏（Latency Hiding）。前两节减小了传输的体积，本节让传输的时间不再暴露在关键路径上。这种"计算与传输管线化"的思维，在大型应用中往往能带来数倍的端到端吞吐提升，也是后续深入优化（如多流并行、CUDA 事件同步、cudaMemcpyAsync 与持久内核的结合）的基石。

综合策略与全景回顾

将本节与前三节的内容串联起来，可以看到一条清晰的性能优化路径：

层次	策略	手段	解决的问题
传输效率	固定内存	cudaMallocHost	消除驱动中转拷贝
传输体积	批量传输 + 数据复用	打包缓冲区 + 设备端流水线	摊薄固定开销，减少总字节数
传输时机	流与异步拷贝	cudaMemcpyAsync + Stream	让传输时间不暴露在关键路径

三层策略互相独立、可以叠加使用：固定内存让单次传输更快，批量传输让传输次数更少，流让传输时间被隐藏。三者结合，PCIe 的物理带宽仍然存在，但程序的有效吞吐可以无限逼近这条带宽的上限。

至此，主机与设备之间的数据通路已经完整呈现：从全景图的内存类型认知，到 cudaMalloc/cudaMemcpy 的基础操作，再到固定内存与零拷贝的精细化控制，再到统一内存的托管模型，最后到批量传输与流并发的优化策略。这条学习路径覆盖了 CUDA 编程中所有与数据传输相关的核心概念。接下来的文章将把视角转向芯片内部——全局内存的访问模式、共享内存的手动管理、以及寄存器的精细分配，这些才是决定单个内核性能的真正战场。