CUDA编程基础：从GPU并行架构到线程与内存模型

做GPU编程的同行，尤其是刚接触CUDA架构的朋友，大多有过这种体验：代码能跑，但心里没底。为什么数据要先拷到显存？为什么线程是这么编排的？为什么稍微改个block尺寸，性能差出好几倍？这些问题如果不从根上解决，后面看再多优化技巧也落不了地。我第一篇写的就是把这些基础概念掰开揉碎，把“大规模并发处理器程序设计”这门经典课程里最核心的地基打牢，适合正在学CUDA并行编程的在校学生，也适合刚想用GPU做加速的工程师。

这一期先不做花哨的优化，只解决三个问题：GPU凭什么快、CUDA的线程到底怎么组织、数据是怎么在显存里流动的。把这三个问题想清楚，再看kernel代码，基本就能做到“看到启动配置就猜到大概性能”。

1. CPU受限困境与GPU并行革命

1.1 性能瓶颈从哪来

先说一个反直觉的现象：现代CPU的主频已经很多年没有大幅提高了。十年前旗舰CPU能跑到3.5GHz，今天还是3.5GHz上下，甚至有往回缩的趋势。为什么？功耗墙。频率再往上拉，芯片发热量和供电压力都受不了。于是厂商换了一条路：加核。从双核到四核再到十六核、三十二核，但核心数的增长也遇到了天花板——内存带宽、缓存一致性、互连带宽都开始拖后腿。你现在拿一颗32核的服务器CPU去跑一个简单的大循环，会发现核心全开也未必能线性加速。

更麻烦的是功耗经济性。CPU为了兼顾延迟敏感型的任务（比如操作系统调度、数据库事务、网络栈处理），内部有大量复杂的控制逻辑、分支预测器、大容量缓存。这些硬件对单线程性能有帮助，但对于“纯粹的数据批量处理”来说，它们都在“空转”。打个比方，CPU像一个独立工作室：装修精美、工具齐全、什么活都能干，但单价高；GPU则更像一条流水线车间：每个工位都极简，但数量巨大，专攻那些能被拆成无数小工的重复劳动。

1.2 GPU的并行本色

GPU最初为图形渲染而生，图形渲染的本质就是对屏幕上几百万个像素做相同的计算：坐标变换、光照计算、颜色输出。这些像素之间互不依赖，天然适合并行。显卡厂商把大量晶体管投入到计算单元（ALU）而不是控制逻辑上，所以一块消费级显卡能有数千个计算核心，而一颗CPU通常只有十几二十个大核。

关键点是吞吐量（throughput）思维。CPU追求低延迟：一个请求进来，希望尽快算完。GPU追求高吞吐：一万个请求进来，希望在单位时间内处理掉最多。单个线程可能比CPU慢，但三万个线程并发执行，整体速度就反超了。这就是“大规模并发处理器”这个名字的含义——不是把单个任务跑得更快，而是把千万个任务的组织成本压下去。

举一个实际的例子：对两个长度为100万的数组做逐元素相加。单核CPU大概耗时几毫秒，32核CPU用OpenMP可能很快，但是同样的数据量在GPU上跑一个简单的kernel，往往只要几十微秒，而且数据规模越往上差距越明显。当然GPU也有短板：单个任务的依赖链长的话，或者频繁需要同步的话，优势就不明显。这也是为什么要先理解任务类型再选方案。

1.3 为什么选CUDA作为学习入口

市面上有OpenCL、DirectCompute、SYCL、HIP等多种异构编程框架，但选CUDA入门有几个很实际的理由。

第一是生态成熟。CUDA的官方文档、示例代码、工具链（nsight、nvcc、cuda-gdb）都是完整的，遇到问题搜到的答案也最多。第二是硬件占有率。不管你在学校里用的是Quadro还是GeForce，在公司用的是A100还是H100，只要跑深度学习或者科学计算，N卡占有率在各个榜单上都很靠前。第三是学习迁移性。CUDA学透了，再看HIP、SYCL基本是换个名字的事，因为它们很多概念同源。

还有一个隐性原因：CUDA的编程模型非常接近GPU硬件的真实行为。它不像某些框架那样给你一个“统一数组”的抽象，让你完全感觉不到内存拷贝。CUDA要求你显式地管理host端和设备端的数据，这种“麻烦”反而能逼着你理解硬件原理。基础阶段多点这种麻烦，是好事。

2. 把GPU组织起来：CUDA的硬件组成与线程模型

2.1 硬件视角：流处理器与流多处理器

在软件里写代码之前，先把GPU里的硬件单元认清楚。最底层的计算单元叫流处理器（Streaming Processor，SP）——俗称CUDA core。每个SP负责执行一个线程的算术运算。一个显卡有几千个SP，但它们不是独立的，而是以组为单位封装在流多处理器（Streaming Multiprocessor，SM）里。每个SM里有若干个SP、一块共享内存（Shared Memory）、一组寄存器文件（Register File）以及调度单元。整颗GPU由多个SM组成，具体数量取决于芯片规模。

拿一辆卡车做类比：SP就是轮子，SM就是车桥，GPU就是整车。轮子不单独转动，要跟着车桥走；车桥的数量决定了整车能装多少货。编程时你接触的是“车桥”的概念——CUDA的线程块最终会被分配到SM上运行，你不需要也不应该指定具体去哪个SM，那是硬件调度器的事。

2.2 软件视角：线程的网格与线程块

CUDA的线程组织是两级的：一个kernel（内核）启动时，创建一个网格（grid），网格由多个线程块（block）组成，每个线程块包含多个线程（thread）。这种三级结构是理解CUDA的核心，也是最容易被新手忽略的地方。

对应到索引上：

• gridDim.x：网格在x维度上有多少个线程块。
• blockIdx.x：当前线程块在网格中的编号。
• blockDim.x：每个线程块里的线程数量。
• threadIdx.x：当前线程在块内的编号。

一个经典的用法是把全局索引算出来：

int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;

这就是线程与数据之间的映射公式。你可以把网格当成一本作业本，线程块当成某一页，线程当成这个页面里的一行。你拿到第几页第几行，就能定位到作业本中的具体位置。真实场景里数据可能是二维的、三维的（比如图像），所以CUDA允许grid和block都定义到dim3类型，最多三个维度。初学阶段先吃透一维，二维三维只是索引公式多几项而已。

2.3 Warp调度与SIMT执行机制

每个名叫“线程”的执行单元，硬件上是按固定大小分组调度的：一组32个线程，叫做一个warp。这是GPU执行的最小硬件单位。也就是说，虽然你可以把block定义成128线程，但硬件实际运行时是4个warp逐个被调度到SM上的。

这背后叫SIMT（Single Instruction, Multiple Thread，单指令多线程）。同一时刻，一个warp里的32个线程在各自的SP上执行同一条指令，但操作的数据各不相同。这和CPU的SIMD（单指令多数据）不一样。CPU的SIMD需要你手动把数据打包到向量寄存器里；而SIMT完全不用管，你写的是普通标量指令，硬件自动按32个线程并行执行。

这个特性带来一个很现实的问题：warp divergence（线程分支发散）。如果代码里有if (condition) { ... } else { ... }，并且一个warp内32个线程里有的走if、有的走else，那么硬件得把两段代码都执行一遍，不满足条件的线程会被掩蔽掉。这等于浪费了一半执行资源。所以说，面向GPU编程要尽量避免warp内的分支不均。不要觉得这句话是理论书上的教条，实际写规约、排序、直方图的时候，它直接决定性能。

再补充一个容易被忽略的点：SM上可以同时驻留多个block，每个block又由多个warp组成。只要还有资源（寄存器、共享内存、线程槽位），SM就会尽量多塞block，这样当一个warp在等待内存数据时，调度器可以立刻切换到另一个warp去执行。这种切换几乎是零成本的，和CPU线程切换保存现场完全不同。这也是GPU能靠“人海战术”吃掉访存延迟的核心原因。

3. 第一个CUDA程序：从环境搭建到代码跑通

3.1 开发环境准备与nvcc

学习CUDA第一步是装工具链。去Nvidia官网下载对应显卡驱动版本的CUDA Toolkit即可。装完之后命令行里会有nvcc -V，这是CUDA的编译器驱动程序。注意：nvcc和其他编译器最大的不同在于，它要处理两种代码：

• host端代码，由NVCC传给宿主编译器（比如GCC/Clang/MSVC）编译，最终在CPU上运行。
• device端代码（也就是kernel），由NVCC编译成GPU能执行的二进制，最终在GPU上运行。

你的.cu文件里可以同时混着C++代码和CUDA kernel，nvcc会根据语法区分它们。建议开发环境用Visual Studio Code加CUDA插件，或者直接用Visual Studio也行。写代码的时候把硬件架构参数搞清楚，编译的时候指定架构可以避免运行时“不支持的二进制格式”这类问题。

如果你拿不准自己显卡的计算能力（Compute Capability，简称CC），运行nvidia-smi看到显卡型号，再去查对应CC版本。常见的比如RTX 3060是8.6，A100是8.0，RTX 4090是8.9。编译时用-arch参数来指定，比如：

nvcc -arch=sm_80 hello.cu -o hello

sm_80表示生成面向Ampere架构（CC 8.0）的代码。为了兼容性，也可以直接用-arch=compute_80,code=sm_80这种更明确的写法。新手阶段可以直接用-arch=native让编译器自动探测，但要注意它在跨机器部署时会有兼容性问题。

3.2 完整代码示例与逐行解读

下面是一个最简单的CUDA程序，它帮我们验证开发环境是否正常，顺便跑通“CPU端调用GPU端”的完整链路：

#include <cstdio> __global__ void helloKernel() { printf("Hello CUDA from thread %d in block %d\n", threadIdx.x, blockIdx.x); } int main() { helloKernel<<<2, 3>>>(); cudaDeviceSynchronize(); return 0; }

先说__global__关键字，它表示这个函数是一个kernel，运行在GPU上，但由CPU端发起调用。kernel的调用语法是function<<<gridSize, blockSize>>>，这里<<<2, 3>>>表示启动一个包含2个block的grid，每个block里有3个线程，合计6个线程并行执行。

打印语句里的threadIdx.x和blockIdx.x分别代表当前线程在block内的编号和block在grid内的编号。运行时你会看到六行输出，编号依次是block 0的0、1、2，以及block 1的0、1、2。输出顺序不一定，因为GPU执行是并行的。

为什么要有cudaDeviceSynchronize()？kernel的启动在CPU看来是异步的，也就是CPU把任务交给GPU后立刻返回，不等GPU执行完。如果没有这一行，程序可能直接return 0退出，GPU端还没开始做事，printf的结果自然也就看不到。这一行强制CPU端阻塞等待GPU端所有任务完成。后面写CUDA错误检查时会发现，这也是一个关键的调试点。

3.3 编译运行实操与常见坑

把上面的代码保存为hello.cu，打开命令行进入文件目录，编译：

nvcc -arch=sm_80 hello.cu -o hello

没有报错的话运行：

./hello

屏幕上会打印6行hello信息。如果这里踩坑了，最常见的有三种：

一是nvcc不是系统命令。原因通常是CUDA Toolkit的bin目录没加到系统PATH。装完后手动把/usr/local/cuda/bin加进去，或者直接使用绝对路径调用。

二是编译报错unsupported gpu architecture。这个通常是因为-arch参数填的架构比你的显卡新，或者填的数字格式写错。先查自己显卡的CC版本，填对再编。

三是运行时报runtime API error ？: no kernel image is available。这个问题经常出现在用过于老的编译选项跑新卡上面。解决方式很简单：指定匹配当前显卡的架构重新编译。例如显卡CC 8.6，那就用-arch=sm_86。

跑通hello之后，可以试试把grid和block的尺寸调成二维三维，观察索引变化。比如<<<dim3(2,2), dim3(2,2)>>>这种写法，你会看到blockIdx.y和threadIdx.y都会参与索引计算。亲手调一次，比看十遍文档都管用。

4. 内存模型与数据搬移：CUDA里的高速公路系统

4.1 六种内存空间速览

CUDA的内存模型是整个编程模型里最需要花工夫的地方。很多初学者一上来就写kernel，结果数据没拷对，算出来的全是垃圾值。先看一张总表：

你可以把全局内存理解为硬盘——容量大、慢、人人都能访问；共享内存反而有点像CPU的L1/L2缓存，但它是由开发者手动控制的，而且极快。每个SM里的共享内存通常只有几十KB到一百多KB，所以它是稀缺资源。寄存器则是每个线程独占的一小摞空间，存线程的局部变量，数量有限，用多了会溢出到局部内存，性能直接下降。

4.2 数据从哪来、到哪去

既然GPU不能直接访问CPU的内存，那么一段通行的数据流程就是：

1. 在设备端分配显存空间：cudaMalloc。
2. 把CPU端的数据拷贝到设备端：cudaMemcpy。
3. 调用kernel处理数据。
4. 把结果拷回CPU端：cudaMemcpy。
5. 释放显存：cudaFree。

看一个向量加法的完整例子，体会整个过程：

#include <cstdio> #include <cstdlib> __global__ void vectorAdd(const float *a, const float *b, float *c, int n) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < n) { c[idx] = a[idx] + b[idx]; } } int main() { const int n = 1000000; 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] = 1.0f; h_b[i] = 2.0f; } float *d_a, *d_b, *d_c; cudaMalloc(&d_a, bytes); cudaMalloc(&d_b, bytes); cudaMalloc(&d_c, bytes); cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice); int blockSize = 256; int gridSize = (n + blockSize - 1) / blockSize; vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n); cudaDeviceSynchronize(); cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost); cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); free(h_a); free(h_b); free(h_c); return 0; }

这段代码里有几个细节值得展开。

int gridSize = (n + blockSize - 1) / blockSize;是一个取整技巧：数据量n不一定能被blockSize整除，所以算上整数块之外还要多出一个小块。因为kernel里我们写了if (idx < n)做边界判断，所以多出来的线程不会越界访存。这个技巧在几乎所有真实kernel里都会用到。

blockSize = 256是我随机选的。事实上它在很多场景下都是不错的起点，但不是“最优解”。最优值取决于kernel的寄存器用量、共享内存使用和GPU架构，一般用256到1024之间的整数。这块属于性能调优内容，后续会细讲。基础阶段先记住：block的线程数最好是warp大小32的倍数，因为一个warp是32个线程，block不是32的倍数时会产生不满的warp，浪费调度槽位。

cudaMemcpy的最后一个参数指定拷贝方向，常见的有cudaMemcpyHostToDevice和cudaMemcpyDeviceToHost。方向写反了程序会报invalid argument错误，这是新手高频错误之一。还要注意，PCIe总线的拷贝速度通常远低于显存内部带宽，如果频繁地在host和device之前来回搬运小块数据，性能会非常差。这就是“数据搬移比计算贵”的来源。

4.3 内存带宽思维：一次拷贝抵几次计算

我们说GPU快，指的是它计算数据的速度快；但数据进不了GPU，一切都白搭。数据从CPU一路跋涉到GPU寄存器，途经PCIe总线、显存、L2缓存、L1缓存，每一级都是时间成本。Nvidia官方后来提供了统一内存（Unified Memory）机制，目的就是藏起一部分拷贝细节，但它底层仍然在搬数据。

给一个直观数量级：假设你要做100万次浮点加法，这个计算量在GPU上可能只需要几微秒的“纯算力时间”；但如果你用cudaMemcpy把两个数组从CPU拷到GPU，再从GPU拷回去，即使按PCIe 4.0的16GB/s带宽算，也要几毫秒。这中间的差距有上百倍。可以这么说：一个kernel里如果没有足够的计算量去“摊薄”拷贝开销，那么整个程序基本输在了搬数据上。

动手验证这个道理的方法很简单，试着把vectorAdd里的数组长度从100万改成1000万、1个亿，你会发现计算时间增长很小的同时，拷贝时间却按比例涨；而当你优化算法时，最先看到收益的地方可能不是kernel本身，而是减少了不必要的内存拷贝。等讲到进阶调优时你会频繁撞到这句话：先减少数据搬运次数，再谈怎么用共享内存。

5. 错误处理与调试工具链：程序跑飞是常态

5.1 把CUDA错误检查当成习惯

CUDA API调用返回的是一个cudaError_t类型值，表示这次调用是否成功。很多初学者忽略这个返回值，结果程序输出一堆NaN或者直接崩溃，都不知道错在哪。一个标准的做法是把每个返回都包起来检查：

#include <cstdio> #define CUDA_CHECK(call) \ do { \ cudaError_t err = (call); \ if (err != cudaSuccess) { \ fprintf(stderr, "CUDA error %s at %s:%d\n", \ cudaGetErrorString(err), \ __FILE__, __LINE__); \ exit(EXIT_FAILURE); \ } \ } while (0)

之后把cudaMemcpy、cudaMalloc、cudaFree等所有调用全部包上CUDA_CHECK()。唯一不能包的是kernel启动本身，因为kernel是异步的，<<<>>>不会立刻暴露错误。所以在kernel调用完后，再加上一行：

CUDA_CHECK(cudaGetLastError()); CUDA_CHECK(cudaDeviceSynchronize());

cudaGetLastError()会拿到最近一次kernel启动（或异步API调用）丢弃的错误。这样配合起来，你就能快速定位是哪个kernel出错。我在实际项目里几乎每个kernel后面都会写这两行，效率提升明显。

5.2 用nvidia-smi和nsight快速定位问题

nvidia-smi是硬件状态工具，它能显示当前GPU利用率、显存占用、运行中的进程。当你怀疑“GPU到底在忙什么”的时候，先跑它看一眼利用率。如果利用率一直很低，说明程序要么在拷贝数据，要么在同步等待，或者是kernel启动间隔太长，这本身就是一种性能信号。

深入一点的性能分析用NVIDIA Nsight Compute（nsight compute），它是profiling工具，可以告诉你每个kernel的瓶颈是访存还是计算。用法很简单：

ncu --kernel-name regex ./your_executable

它会报告每个kernel的occupancy、memory throughput、compute throughput。这些指标看起来可能懵，但基础阶段只需看一个概念：occupancy（占用率）。占用率是指SM上活跃warp数与理论最大warp数的比值。占用率高不代表性能最好，但占用率过低通常说明资源限制导致调度空间不足，比如共享内存用太多或者寄存器用太多。

nsight输出的详细报告里会有“Warp State”部分，能看到stall原因。比如long scoreboard表示线程在等待全局内存数据，wait表示等待固定延时。这些细分项就是定位瓶颈的钥匙。基础阶段不知道怎么看就截图背下来，后面调优时会自然而然地用到。

5.3 编译错误与运行时错误速查表

整理一份编译和运行过程中高频出现的错误，方便你对着排查：

还有一个容易踩的坑：在kernel里使用C++标准库容器（std::vector等）。它们依赖CPU运行时的内存管理，普通kernel里是不能直接用的。kernel里做数据存储应该用显存指针加手动管理的方式，或者使用CUDA自带的向量类型。

调试cuda-gdb也需要知道一点：cuda-gdb可以设置设备端断点，查看每个线程的值。但它比较依赖经验，新手阶段建议先靠打印和错误检查，把逻辑调对，再考虑用调试器。打印是温和的，cuda-gdb是严厉的。

6. 写在最后的个人体会

我自己最初学CUDA的时候，最深的感悟是：这门技术难的不是API，而是思维方式的转变——从“单核串行”变成“海量线程并行”。前期看不到立刻的收益，因为环境配置、错误处理、数据搬运这些“杂活”看起来都和性能无关。但把这些基本功打牢，后面每走一步都会快很多。

根据我踩过的坑，建议你按这个顺序走：先把今天讲的线程索引公式写熟，能够不看文档就写出正确的全局索引；再亲手跑三个小实验——向量加法、批量求和、矩阵转置；第三个实验特别值得做，因为非连续访问模式和连续访问模式的性能差距会让你对“内存布局”这四个字有深刻印象。

如果你读到这里，已经能独立编译运行一个带cudaMemcpy的完整程序，说明你的CUDA基础框架已经建立起