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.xthreadIdx.x:我在自己 block 内是第几个线程blockIdx.x:我属于第几个 blockblockDim.x:每个 block 有多少线程比喻:3 号车间(blockIdx3)的 5 号工人(threadIdx5),每车间 10 人(blockDim10)→ 全厂工号 3×10535 → 处理第 35 个数据。4. 边界检查 ★★(CUDA 标配,新手必踩的坑)线程总数(block数 × 每block线程数)几乎不可能正好等于数据量,通常向上取整到稍多一点 → 总有多余线程。关键认知:多余线程默认不会自觉空闲,它们照样执行 kernel(SIMT 齐步走),照样算出自己的 idx(如 1000~1023)去访问数组 →越界访问(数组合法下标只到 999)。必须用代码显式挡住:intidxblockIdx.x*blockDim.xthreadIdx.x;if(idxn){// 边界检查:只有合法编号才干活c[idx]a[idx]b[idx];}编号 ≥ n 的线程判断为假 → 跳过 → 不越界。5. 第一个 CUDA 程序:向量加法(c[i]a[i]b[i])CUDA 的 “Hello World”,体现把循环拆成并行。#includestdio.h// __global__ : 在 GPU 执行、由 CPU 调用的 kernel__global__voidvectorAdd(float*a,float*b,float*c,intn){intidxblockIdx.x*blockDim.xthreadIdx.x;if(idxn){c[idx]a[idx]b[idx];// 每个线程只算第 idx 个元素}}intmain(){intn1000;size_t sizen*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(inti0;in;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. 启动 kernelintthreadsPerBlock256;intblocks(nthreadsPerBlock-1)/threadsPerBlock;// 向上取整vectorAddblocks,threadsPerBlock(d_a,d_b,d_c,n);// grid, block CUDA 专有语法// 5. 结果 GPU - CPUcudaMemcpy(h_c,d_c,size,cudaMemcpyDeviceToHost);// 6. 验证for(inti0;i5;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) → kernelgrid,block 并行计算 → cudaMemcpy 拷回 CPU(D2H) → cudaFree 释放拷上去、算、拷回来正是显存层级的体现:数据必须先进全局显存,GPU 才能算。6. 越界实验(亲眼见识越界不一定报错)把边界检查删掉、甚至故意启动 25600 个线程(远超 n1000),结果:“kernel 执行无报错”,结果还全对。为什么 GPU 这么宽容:cudaMalloc 实际可能划一大块对齐内存,越界地址很多落在多给的余量或其他不严格保护的区域。硬件不在每次访问做边界校验(那样太慢,违背 GPU 追求吞吐的设计)。深刻教训:越界访问常常不当场报错,而是埋雷。测试里跑得好好的,换数据规模/换卡后踩到关键内存就崩在最意想不到处,极难复现定位。→边界检查不是可选保险,是必需纪律。不能指望 GPU 兜底,它默认不管。7. 手动错误检查 ★(比工具更基础的必备习惯)CUDA kernel 是异步启动的(CPU 发完命令就往下走,不等 GPU),kernel 内部错误不会自动冒泡,必须主动问:cudaDeviceSynchronize();// 等 GPU 跑完 kernelcudaError_t errcudaGetLastError();// 取最近一次 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的习惯,把防线建在代码里而非工具上。已掌握(模块二上半)单线程心智模型(写一份活,GPU 复制 N 份)Thread/Block/Grid 三层组织全局索引公式 idx blockIdx.x*blockDim.x threadIdx.x边界检查 if(idxn) 及其必要性(亲手验证越界不报错)CUDA 标准骨架:分配/拷贝H2D/计算/拷回D2H/释放手动错误检查 cudaGetLastError cudaGetErrorString知道 WSL 下 sanitizer 不可用