当前位置：首页 > news >正文

【infra之路】阶段二 · 模块二:CUDA 编程入门(上)— 基本功与向量加法

news 2026/6/1 9:19:27

AI Infra 学习路线 · 阶段二 · 模块二(上半部分)
目标:把模块一的理论(SIMT/线程/显存)落到代码 —— 亲手写、编译、运行第一个 CUDA 程序
环境:WSL2 + CUDA Toolkit 12.8 (nvcc) + RTX 5060 Ti,PyTorch cu130(写独立 .cu 不受版本差异影响)

1. 核心心智模型:写"一个线程的活",GPU 复制几千份同时跑

普通编程:一个线程从头到尾循环处理 1000 个数据。
CUDA:你写的 kernel 描述单个线程要干的那一小份活,然后告诉 GPU"启动 N 个线程",GPU 把这份活复制 N 份,几千核心同时执行(SIMT,一声令下齐步走)。

关键问题:几千线程跑同一段代码,每个怎么知道自己处理哪块数据?→ 靠唯一编号。

2. 线程的组织:Thread → Block → Grid

层级	说明	对应硬件
Thread线程	最小单位,执行 kernel 里那份活	工人
Block线程块	一组线程打包;整个分配到一个 SM上;块内线程可共享该 SM 的共享显存	派给一个车间的工作包
Grid网格	所有 block 合起来 = 本次 kernel 启动的全部线程	整个工厂这次的订单

启动 kernel 要指定"多少 block、每 block 多少线程",合起来 = 总线程数。

3. 全局索引公式 ★(把"齐步走的线程"和"各处理不同数据"连起来的桥梁)

全局索引 idx = blockIdx.x * blockDim.x + threadIdx.x

threadIdx.x:我在自己 block 内是第几个线程
blockIdx.x:我属于第几个 block
blockDim.x:每个 block 有多少线程

比喻:3 号车间(blockIdx=3)的 5 号工人(threadIdx=5),每车间 10 人(blockDim=10)→ 全厂工号 3×10+5=35 → 处理第 35 个数据。

4. 边界检查 ★★(CUDA 标配,新手必踩的坑)

线程总数(block数 × 每block线程数)几乎不可能正好等于数据量,通常向上取整到稍多一点 → 总有多余线程。

关键认知:多余线程默认不会自觉空闲,它们照样执行 kernel(SIMT 齐步走),照样算出自己的 idx(如 1000~1023)去访问数组 →越界访问(数组合法下标只到 999)。

必须用代码显式挡住:

intidx=blockIdx.x*blockDim.x+threadIdx.x;if(idx<n){// 边界检查:只有合法编号才干活c[idx]=a[idx]+b[idx];}

编号 ≥ n 的线程判断为假 → 跳过 → 不越界。

5. 第一个 CUDA 程序:向量加法(c[i]=a[i]+b[i])

CUDA 的 “Hello World”,体现"把循环拆成并行"。

#include<stdio.h>// __global__ : 在 GPU 执行、由 CPU 调用的 kernel__global__voidvectorAdd(float*a,float*b,float*c,intn){intidx=blockIdx.x*blockDim.x+threadIdx.x;if(idx<n){c[idx]=a[idx]+b[idx];// 每个线程只算第 idx 个元素}}intmain(){intn=1000;size_t size=n*sizeof(float);// 1. CPU(host)准备数据,h_ 前缀 = hostfloat*h_a=(float*)malloc(size);float*h_b=(float*)malloc(size);float*h_c=(float*)malloc(size);for(inti=0;i<n;i++){h_a[i]=i;h_b[i]=i*2;}// 2. GPU(device)分配显存(全局显存),d_ 前缀 = devicefloat*d_a,*d_b,*d_c;cudaMalloc(&d_a,size);cudaMalloc(&d_b,size);cudaMalloc(&d_c,size);// 3. 数据 CPU -> GPUcudaMemcpy(d_a,h_a,size,cudaMemcpyHostToDevice);cudaMemcpy(d_b,h_b,size,cudaMemcpyHostToDevice);// 4. 启动 kernelintthreadsPerBlock=256;intblocks=(n+threadsPerBlock-1)/threadsPerBlock;// 向上取整vectorAdd<<<blocks,threadsPerBlock>>>(d_a,d_b,d_c,n);// <<<grid, block>>> CUDA 专有语法// 5. 结果 GPU -> CPUcudaMemcpy(h_c,d_c,size,cudaMemcpyDeviceToHost);// 6. 验证for(inti=0;i<5;i++)printf("c[%d]=%.1f\n",i,h_c[i]);// 7. 释放cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);free(h_a);free(h_b);free(h_c);return0;}

编译运行:

nvcc vector_add.cu-ovector_add ./vector_add

结果:c[i]=3i 全部正确(c[999]=2997)。1000 个加法被一批线程同时算完。

CUDA 程序标准骨架(记牢这条主线)

CPU 准备数据 → cudaMalloc 分配显存 → cudaMemcpy 拷上 GPU(H2D) → kernel<<<grid,block>>> 并行计算 → cudaMemcpy 拷回 CPU(D2H) → cudaFree 释放

"拷上去、算、拷回来"正是显存层级的体现:数据必须先进全局显存,GPU 才能算。

6. 越界实验(亲眼见识"越界不一定报错")

把边界检查删掉、甚至故意启动 25600 个线程(远超 n=1000),结果:
“kernel 执行无报错”,结果还全对。

为什么 GPU 这么宽容:

cudaMalloc 实际可能划一大块对齐内存,越界地址很多落在"多给的余量"或其他不严格保护的区域。
硬件不在每次访问做边界校验(那样太慢,违背 GPU 追求吞吐的设计)。

深刻教训:越界访问常常不当场报错,而是埋雷。测试里跑得好好的,换数据规模/换卡后踩到关键内存就崩在最意想不到处,极难复现定位。
→边界检查不是可选保险,是必需纪律。不能指望 GPU 兜底,它默认不管。

7. 手动错误检查 ★(比工具更基础的必备习惯)

CUDA kernel 是异步启动的(CPU 发完命令就往下走,不等 GPU),kernel 内部错误不会自动冒泡,必须主动问:

cudaDeviceSynchronize();// 等 GPU 跑完 kernelcudaError_t err=cudaGetLastError();// 取最近一次 CUDA 错误if(err!=cudaSuccess)printf("CUDA 错误: %s\n",cudaGetErrorString(err));

正经 CUDA 代码里到处是这种检查。WSL 用不了 sanitizer(见下),这套手动检查永远可用,更该掌握。

8. WSL 环境限制(记一笔)

compute-sanitizer(CUDA 自带的内存错误检测利器,老版叫 cuda-memcheck)在 WSL 下用不了:
报错Failed to initialize WDDM debugger interface/Device not supported。
原因:它需要直接访问 GPU 调试接口,WSL 隔着 Windows 驱动(WDDM)这层拿不到。

结论:WSL 能跑 CUDA,但部分底层调试/分析工具受限。深度使用这类工具需原生 Linux。
→ 反而更要养成"手写边界检查 + cudaGetLastError"的习惯,把防线建在代码里而非工具上。