第2节线程组织与索引计算【CUDA高性能编程实战‑板块一】 专栏《CUDA 高性能编程实战从 Kernel 到 FlashAttention》✨ 本篇为板块一第2节详解Grid/Block层级线程架构、三维索引、全局唯一线程ID计算配套完整可运行代码示例 阅读目标掌握任意维度下线程索引推导能独立实现一维、二维数据并行读写学习目标学完本节你将掌握理解 Grid、Block、Thread 三层层级架构逻辑分清gridIdx/blockIdx/threadIdx三组内置变量含义熟练计算一维、二维场景下全局唯一线程ID写出通用索引计算模板避免数组越界访问理解dim3三维启动参数的使用规则1. CUDA三层线程架构GPU并行执行单元分为三层嵌套结构从上至下Grid网格一次Kernel启动的全部线程集合由多个Block组成Block线程块一组共享片上共享内存的线程单个Block最大1024线程Thread线程最小执行单元对应SIMT流水线单条执行通道层级关系示意图Grid ├─ Block(0,0) │ ├─ Thread(0,0,0) │ ├─ Thread(1,0,0) │ └─ ... ├─ Block(1,0) │ ├─ Thread(0,0,0) │ └─ ... └─ ...核心特性同一个Block内线程可同步、访问共享内存不同Block线程无通信机制Grid、Block均支持三维定义x/y/z日常一维计算只用x维度2. 内置索引变量只读CUDA自动为每个线程分配三组内置dim3类型索引变量名作用域含义gridDim全局Grid整体尺寸gridDim.xBlock总数blockDim单个Block每个Block内线程总数blockDim.x单块线程数blockIdxBlock内所有线程当前线程所在Block编号threadIdx单个线程当前线程在Block内部的局部编号一维场景索引逻辑最常用一维下仅使用x分量全局唯一线程ID计算公式int global_tid blockIdx.x * blockDim.x threadIdx.x;公式拆解blockIdx.x * blockDim.x当前Block前面所有Block的线程总量 threadIdx.x叠加当前Block内部局部偏移3. dim3 三维启动参数dim3是CUDA内置三维结构体默认未赋值维度自动填充为1。两种启动写法等价// 写法1纯int一维启动 kernel32, 128(); // 写法2dim3标准三维启动 dim3 grid_size(32, 1, 1); dim3 block_size(128, 1, 1); kernelgrid_size, block_size();二维矩阵计算标准dim3示例// 二维场景每行一个Block每个Block 256线程 dim3 grid(rows, 1); dim3 block(256, 1);4. 完整代码示例1一维数组索引遍历#include cstdio __global__ void indexTestKernel(int data_len) { // 计算全局唯一线程ID int tid blockIdx.x * blockDim.x threadIdx.x; // 边界判断防止线程超出数组长度避免越界 if (tid data_len) { printf(Block:%d Thread:%d GlobalTid:%d\n, blockIdx.x, threadIdx.x, tid); } } int main() { int data_size 20; // 启动配置4个Block每个Block 8线程总32线程 indexTestKernel4, 8(data_size); cudaDeviceSynchronize(); cudaError_t err cudaGetLastError(); if (err ! cudaSuccess) { printf(Kernel启动异常%s\n, cudaGetErrorString(err)); return -1; } return 0; }编译运行指令nvcc thread_index.cu -o thread_index ./thread_index5. 二维矩阵索引计算图像处理、矩阵运算常用二维Grid/Block全局坐标推导// 当前像素在矩阵中的x、y坐标 int x blockIdx.x * blockDim.x threadIdx.x; int y blockIdx.y * blockDim.y threadIdx.y; // 一维展平数组下标行优先存储 int global_idx y * matrix_width x;二维Kernel完整示例片段__global__ void matrixKernel(float* mat, int width, int height) { int x blockIdx.x * blockDim.x threadIdx.x; int y blockIdx.y * blockDim.y threadIdx.y; if (x width y height) { int idx y * width x; mat[idx] x y * 0.5f; } }6. 索引避坑核心要点必须加边界判断if(tid length)总线程数 gridDim.x * blockDim.x通常会大于数据长度多余线程不处理数据不加判断会数组越界、内存报错。Block单块线程上限1024blockDim.x * blockDim.y * blockDim.z ≤ 1024超过直接启动失败。索引变量只读不可修改blockIdx.x 5;这类代码编译报错。全局ID公式固定不变一维计算统一使用blockIdx.x * blockDim.x threadIdx.x不要颠倒乘数顺序。7. 课后练习修改一维示例启动参数为5, 6手动计算总线程数观察输出的最大global_tid。改写代码实现二维512×512矩阵遍历dim3设置为dim3 block(16,16)推导grid尺寸。移除边界判断语句传入较小data_len观察程序是否出现内存错误。尝试将blockDim设置为1025查看Kernel启动报错信息巩固Block线程限制知识点。专栏上下篇上一篇第1节CUDA 编程模型与核函数【CUDA高性能编程实战‑板块一】下一篇第3节SIMT与Warp基础【CUDA高性能编程实战‑板块一】提示所有代码复制后可直接nvcc编译运行建议手动修改启动参数直观感受线程索引变化。