CUDA编程值Kernal

一、什么是GPU KernelKernel内核是在GPU上并行执行的函数。与CPU上只执行一次的传统C/C函数不同GPU kernel会被成千上万个线程同时执行每个线程执行相同的代码但处理不同的数据。在CUDA中kernel通过__global__关键字标识调用时使用特殊的...语法指定线程数量。二、Kernel的线程层次结构2.1 三层架构Thread → Block → GridGPU将线程组织成三级层次结构Thread线程最小的执行单元每个线程执行kernel代码的一个副本拥有独立的寄存器和程序计数器。Block线程块一组线程的集合同一block内的线程可以通过共享内存高效交换数据使用__syncthreads()进行同步必须在同一个流式多处理器SM上执行每个block最多包含1024个线程。Grid网格所有block的集合一个kernel启动的所有线程构成一个grid。Grid的维度可以是一维、二维或三维。2.2 内置变量与索引计算每个线程通过内置变量获取自己的身份信息变量含义threadIdx线程在block内的索引blockIdxblock在grid内的索引blockDimblock的维度线程数gridDimgrid的维度block数全局线程ID计算公式一维情况global_thread_id blockIdx.x * blockDim.x threadIdx.x三、Kernel的启动与执行流程3.1 总体流程一个完整的CUDA程序执行流程如下1. 分配主机内存和GPU显存2. 初始化主机数据3. 将数据从主机内存复制到GPU显存4. 启动kernel执行计算5. 将结果从GPU显存复制回主机内存6. 释放内存3.2 Kernel启动的底层机制“邮局模型”Kernel的启动可以类比为CPU向GPU“投递邮包”第一步准备命令包CPU将kernel的函数指针、网格/线程块维度、参数等打包成特定格式的命令包。第二步投递到命令队列命令包被放入驻留在内存中的命令队列环形缓冲区。第三步敲响门铃DoorbellCPU通过写特定的内存映射I/O寄存器通知GPU“有新任务”。第四步硬件调度器接管GPU前端的硬件调度器NVIDIA的GigaThread Engine通过DMA从队列中拉取命令包kernel正式进入设备端执行。四、硬件执行模型SM与Warp4.1 流式多处理器SMGPU由多个流式多处理器SM组成每个SM包含CUDA核心阵列执行实际算术运算寄存器文件每个线程的私有存储共享内存block内线程共享调度单元负责分配任务当kernel启动时block被分发到不同的SM上执行。一个block不能被拆分到多个SM但多个block可以分配到同一个SM。4.2 Warp真正的执行单元Warp是GPU调度的基本单位。每个warp包含32个连续线程这些线程同步执行锁步同一时刻执行相同指令共享程序计数器所有线程从同一地址取指令拥有独立寄存器每个线程有自己的数据当block被分配给SM后SM将其中的线程划分为warp。如果block大小不是32的倍数最后一个warp会有空闲线程。Warp调度SM中的warp调度器每个时钟周期选择一个就绪的warp执行。warp有三种状态Selected正在执行Eligible准备就绪等待执行Stalled因等待数据等原因停滞4.3 SIMT架构GPU采用SIMT单指令多线程架构一条指令控制多个线程执行。Warp分歧Divergence当warp内线程遇到条件分支走向不同路径时GPU会串行执行各分支路径禁用不在当前路径上的线程。这会导致性能下降因此应尽量避免warp内的分支分歧。4.4 延迟隐藏Latency HidingGPU通过大量active warp实现延迟隐藏当某些warp因内存访问停滞时调度器立即切换到其他就绪warp执行。这就像流水线——某个工位等待材料时其他工位继续工作整体效率不受影响。每个SM可驻留的warp数量有限通常64个更多的active warp意味着更好的延迟隐藏能力。五、GPU内存层次结构GPU的内存层次结构遵循“离ALU越近越快、越小、越贵”的原则5.1 各级内存概览内存类型作用域访问速度典型大小寄存器每线程私有最快1周期每SM数百KB共享内存每block共享极快几周期每SM 64-164KBL1缓存每SM快几十KBL2缓存全局中等几MB到几十MB全局内存所有线程可访问慢数百周期几GB到几十GB常量内存所有线程只读缓存命中时快64KB5.2 各级内存详解寄存器每个线程私有的最快存储用于存放频繁使用的局部变量。共享内存Shared Memory同一block内所有线程共享是线程间通信延迟最小的方法。由程序员显式管理。全局内存Global Memory所有线程均可读写是所有数据的“仓库”容量最大但延迟最高。常量内存Constant Memory只读适合所有线程访问相同数据的场景。六、完整示例向量加法6.1 Kernel代码__global__ void vectorAdd(const float* A, const float* B, float* C, int N) { // 计算全局线程ID int idx blockIdx.x * blockDim.x threadIdx.x; // 边界检查防止越界 if (idx N) { C[idx] A[idx] B[idx]; } }6.2 主机端调用代码int main() { int N 1000000; size_t bytes N * sizeof(float); // 1. 分配主机内存 float *h_A (float*)malloc(bytes); float *h_B (float*)malloc(bytes); float *h_C (float*)malloc(bytes); // 2. 初始化数据 for (int i 0; i N; i) { h_A[i] rand() / (float)RAND_MAX; h_B[i] rand() / (float)RAND_MAX; } // 3. 分配GPU显存 float *d_A, *d_B, *d_C; cudaMalloc(d_A, bytes); cudaMalloc(d_B, bytes); cudaMalloc(d_C, bytes); // 4. 将数据从主机复制到GPU cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice); // 5. 配置并启动kernel int threadsPerBlock 256; int blocksPerGrid (N threadsPerBlock - 1) / threadsPerBlock; vectorAddblocksPerGrid, threadsPerBlock(d_A, d_B, d_C, N); // 6. 等待GPU完成 cudaDeviceSynchronize(); // 7. 将结果从GPU复制回主机 cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost); // 8. 释放内存 cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); free(h_A); free(h_B); free(h_C); return 0; }6.3 执行过程分解Step 1 - 数据准备主机分配内存并初始化A、B向量。Step 2 - 数据传输通过PCIe总线将数据从主机内存复制到GPU全局内存。Step 3 - Kernel启动blocksPerGrid, threadsPerBlock配置启动网格。Step 4 - Block分发GigaThread Engine将block分发到各SM。假设有100万个元素、每block256线程约3907个block被分发。Step 5 - Warp划分每个SM收到block后将256线程划分为8个warp256/328。Step 6 - 指令执行warp调度器每个周期选择一个warp执行。32个线程同时执行C[idx] A[idx] B[idx]。Step 7 - 结果返回GPU计算完成数据从全局内存复制回主机。七、优化关键点7.1 内存合并访问Coalesced Access相邻线程应访问连续的内存地址使硬件能将多个访问合并为一次事务。7.2 共享内存使用将频繁访问的全局内存数据加载到共享内存减少对慢速全局内存的访问。7.3 占用率Occupancy每个SM上活跃warp与最大可能warp数的比率。更高的占用率有助于更好地隐藏延迟。7.4 避免Warp分歧尽量让warp内32个线程走相同的执行路径。八、总结GPU kernel的执行是一个从软件抽象到硬件实现的完整流程软件层面程序员编写__global__函数通过指定线程组织结构运行时层面CPU将kernel打包为命令包通过队列投递给GPU硬件层面GPU的SM将block划分为warpwarp调度器以SIMT方式执行内存层面数据在寄存器、共享内存、全局内存间流动理解这一完整链条是写出高性能GPU程序的关键。从线程网格的精心设计到warp分歧的谨慎规避再到内存访问模式的深度优化每一步都直接影响着最终的计算效率。