1. 内存拷贝基础

各位同学，咱们今天聊点实在的。内存拷贝，听起来简单吧？就是把数据从CPU搬到GPU，或者搬回来。但我跟你说，这里面门道可多了。我刚开始做CUDA那会儿，就因为这个吃了不少亏。

先看一张图，帮你快速建立整体认知：

1.1 CPU与GPU内存架构差异

CPU和GPU的内存设计，出发点完全不同。CPU追求低延迟，GPU追求高吞吐。说白了，CPU是「反应快」，GPU是「力气大」。

CPU内存架构：CPU的内存是统一寻址的，所有核心共享一个内存空间。它的缓存层级很复杂——L1、L2、L3，目的就是让数据离核心更近。我做过一个测试，CPU从L1缓存读数据只要1ns，但从主存读要100ns。差了整整两个数量级。

GPU内存架构：GPU的显存（Global Memory）是所有线程块共享的，但延迟很高。好在GPU有大量线程并行执行，可以通过「线程切换」来隐藏延迟。嗯，这里要注意：GPU还有共享内存（Shared Memory），它就在SM内部，延迟低得多，但容量很小，通常只有几十KB。

1.2 PCIe传输原理

PCIe是连接CPU和GPU的「高速公路」。我当年第一次看PCIe的spec，差点睡着——太枯燥了。但后来调性能时才发现，不理解PCIe真的不行。

PCIe的核心是「点对点串行传输」。什么意思？就是数据像流水一样，一位一位地传。它有几个关键参数：

PCIe 版本	单通道带宽 (x1)	x16 带宽	编码方式
Gen3	~1 GB/s	~16 GB/s	128b/130b
Gen4	~2 GB/s	~32 GB/s	128b/130b
Gen5	~4 GB/s	~64 GB/s	128b/130b

你想想看，PCIe Gen3 x16的理论带宽是16 GB/s，但实际能跑到多少？我实测过，大概12-13 GB/s。为什么？因为还有协议开销、数据包头的损耗。所以写代码时别按理论值算，留点余量。

1.3 cudaMemcpy 详解

cudaMemcpy是CUDA里最常用的内存拷贝函数。它的原型长这样：

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

参数不多，但每个都有讲究：

• dst：目标地址。可以是主机内存，也可以是设备显存。
• src：源地址。同上。
• count：拷贝的字节数。注意别越界。
• kind：拷贝方向。有四种：

枚举值	方向	说明
cudaMemcpyHostToDevice	主机 → 设备	最常用，把数据从CPU内存搬到GPU显存
cudaMemcpyDeviceToHost	设备 → 主机	把结果从GPU搬回CPU
cudaMemcpyDeviceToDevice	设备 → 设备	显存内部拷贝，不经过PCIe
cudaMemcpyHostToHost	主机 → 主机	基本不用，用memcpy就行

我个人的习惯是，每次调用cudaMemcpy后都检查返回值。别嫌麻烦，出错了能立刻知道。比如：

cudaError_t err = cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
if (err != cudaSuccess) {
    printf("cudaMemcpy failed: %s\n", cudaGetErrorString(err));
    exit(-1);
}

1.4 实战经验：如何避免「拷贝瓶颈」

我做过一个图像处理项目，数据量很大，每帧要传几十MB。一开始直接cudaMemcpy，结果发现大部分时间都花在拷贝上了。后来怎么优化的？

1. 减少拷贝次数：能一次传完就别分多次。合并小数据成大块。
2. 使用固定内存：用cudaHostAlloc分配主机内存，可以提升HostToDevice的带宽。我实测过，能快20%-30%。
3. 异步传输：用cudaMemcpyAsync配合流（Stream），让拷贝和计算重叠。

嗯，这里要特别说一下固定内存。普通malloc分配的内存是分页的，GPU访问时可能触发缺页中断。而cudaHostAlloc分配的是「钉住」的内存，不会被换出，所以PCIe传输时可以直接DMA访问，效率高很多。

好了，这一章就到这里。记住我上面说的那些坑，写代码时多留个心眼。