CUDA grid/block 到矩阵映射示例（矩阵加法）

CUDA grid/block 到矩阵映射示例矩阵加法给一个例子 grid block 到矩阵的映射 以矩阵加法来表示我们用一个非常具体、可以直接运行脑补的例子把 CUDA 里的 grid / block → 矩阵索引彻底讲清楚并用矩阵加法来表达。一、问题设定最标准的场景我们做一个矩阵加法 假设矩阵大小是, , ∈ ⁸ˣ⁸二、CUDA 配置我们这样划分线程dim3block(4,4);// 每个 block 4x4 线程dim3grid(2,2);// 2x2 个 block 总线程数x方向2 × 4 8y方向2 × 4 8刚好覆盖整个矩阵。三、核心映射公式最重要线程 → 矩阵坐标intcolblockIdx.x*blockDim.xthreadIdx.x;introwblockIdx.y*blockDim.ythreadIdx.y; 对应关系col ₓ ⋅ ₓ ₓrow ᵧ ⋅ ᵧ ᵧ四、直观展开关键理解我们把整个 8×8 矩阵分成 4 个 block 区域Block 分布五、一个 block 内的映射举例block(1,0)blockIdx.x 1blockIdx.y 0线程范围threadIdx.x ∈ [0,3]threadIdx.y ∈ [0,3]映射到矩阵col 1*4 tx → [4,7]row 0*4 ty → [0,3] 覆盖矩阵区域行: 0~3列: 4~7六、完整矩阵映射图核心矩阵 C (8x8)矩阵 C (8x8) col → 0 1 2 3 | 4 5 6 7 row --------------------- 0 A A A A | B B B B 1 A A A A | B B B B 2 A A A A | B B B B 3 A A A A | B B B B ----------------- 4 C C C C | D D D D 5 C C C C | D D D D 6 C C C C | D D D D 7 C C C C | D D D D对应 block区域blockIdxA(0,0)B(1,0)C(0,1)D(1,1)七、矩阵加法的 CUDA 表达每个线程做一件事C[row][col]A[row][col]B[row][col]; 数学表达ᵢⱼ ᵢⱼ ᵢⱼ CUDA 展开后本质(ᵧᵧᵧ, ₓₓₓ) ... ...八、完整 CUDA Kernel 示例__global__voidmatAdd(float*A,float*B,float*C,intN){intcolblockIdx.x*blockDim.xthreadIdx.x;introwblockIdx.y*blockDim.ythreadIdx.y;if(rowNcolN){intidxrow*Ncol;C[idx]A[idx]B[idx];}}九、本质总结你要抓住这个你可以把整个 CUDA 映射理解成grid 把大矩阵切块block 每个块内部再细分thread 每个元素一个计算单元数学上就是(,)(blockIdx.y,blockIdx.x)⊗(blockDim.y,blockDim.x)(threadIdx.y,threadIdx.x)十、再往深一层工程本质你其实已经触碰到一个更本质的东西 CUDA 就是在做离散空间 → 并行计算单元 的映射函数这个映射决定memory coalescing是否连续访问warp 效率cache 命中bank conflict如果你愿意我可以帮你再往下一层拆 从这个例子直接推到warp 是怎么切的32线程怎么落在矩阵上为什么 (x方向连续) 才 coalesced为什么很多 kernel 要 transpose这个地方一旦打通你对 GPU 的理解会从“会用”变成“能设计”。给出完整可执行代码已思考若干秒下面给你一份完整可执行的 CUDA 矩阵加法示例重点展示grid 和 block 如何映射到矩阵每个线程如何对应一个矩阵元素最后把结果打印出来验证这份代码可以直接保存为 mat_add.cu 编译运行。完整代码#includeiostream#includevector#includecuda_runtime.h#defineCHECK_CUDA(call)\do{\cudaError_terrcall;\if(err!cudaSuccess){\std::cerrCUDA error at __FILE__:__LINE__\ codeerr\ \cudaGetErrorString(err)\std::endl;\exit(EXIT_FAILURE);\}\}while(0)// 打印矩阵voidprintMatrix(conststd::vectorfloatmat,introws,intcols,constchar*name){std::coutname std::endl;for(intr0;rrows;r){for(intc0;ccols;c){std::coutmat[r*colsc]\t;}std::coutstd::endl;}std::coutstd::endl;}// CUDA Kernel: 矩阵加法__global__voidmatAddKernel(constfloat*A,constfloat*B,float*C,introws,intcols){// block 内线程坐标 block 坐标 - 全局矩阵坐标intcolblockIdx.x*blockDim.xthreadIdx.x;// 列introwblockIdx.y*blockDim.ythreadIdx.y;// 行if(rowrowscolcols){intidxrow*colscol;C[idx]A[idx]B[idx];// 打印映射关系printf(block(%d,%d), thread(%d,%d) - matrix(%d,%d), A%.1f, B%.1f, C%.1f\n,blockIdx.x,blockIdx.y,threadIdx.x,threadIdx.y,row,col,A[idx],B[idx],C[idx]);}}intmain(){// 矩阵大小constintrows8;constintcols8;constintsizerows*cols;constsize_tbytessize*sizeof(float);// Host 端矩阵std::vectorfloath_A(size);std::vectorfloath_B(size);std::vectorfloath_C(size,0.0f);// 初始化矩阵// A: 0,1,2,3...// B: 全部 100for(intr0;rrows;r){for(intc0;ccols;c){intidxr*colsc;h_A[idx]static_castfloat(idx);h_B[idx]100.0f;}}printMatrix(h_A,rows,cols,A);printMatrix(h_B,rows,cols,B);// Device 指针float*d_Anullptr,*d_Bnullptr,*d_Cnullptr;CHECK_CUDA(cudaMalloc(d_A,bytes));CHECK_CUDA(cudaMalloc(d_B,bytes));CHECK_CUDA(cudaMalloc(d_C,bytes));// 拷贝到 GPUCHECK_CUDA(cudaMemcpy(d_A,h_A.data(),bytes,cudaMemcpyHostToDevice));CHECK_CUDA(cudaMemcpy(d_B,h_B.data(),bytes,cudaMemcpyHostToDevice));// 配置 block 和 griddim3block(4,4);// 每个 block 4x4 个线程dim3grid((colsblock.x-1)/block.x,(rowsblock.y-1)/block.y);std::coutLaunch config:std::endl;std::coutblock (block.x, block.y, block.z)std::endl;std::coutgrid (grid.x, grid.y, grid.z)std::endl;std::coutstd::endl;// 启动 KernelmatAddKernelgrid,block(d_A,d_B,d_C,rows,cols);CHECK_CUDA(cudaGetLastError());CHECK_CUDA(cudaDeviceSynchronize());// 拷回 HostCHECK_CUDA(cudaMemcpy(h_C.data(),d_C,bytes,cudaMemcpyDeviceToHost));printMatrix(h_C,rows,cols,C A B);// 验证结果bool oktrue;for(inti0;isize;i){floatexpectedh_A[i]h_B[i];if(h_C[i]!expected){std::cerrMismatch at index i, got h_C[i], expected expectedstd::endl;okfalse;break;}}if(ok){std::coutResult check passed.std::endl;}else{std::coutResult check failed.std::endl;}// 释放显存CHECK_CUDA(cudaFree(d_A));CHECK_CUDA(cudaFree(d_B));CHECK_CUDA(cudaFree(d_C));return0;}编译方式如果你本机有 CUDA 编译器 nvcc直接编译nvcc mat_add.cu-omat_add运行./mat_add你会看到什么这个程序会先打印矩阵 A 和 B然后 GPU 里每个线程都会打印类似block(1,0), thread(2,3) - matrix(3,6), A30.0, B100.0, C130.0这句话的意思就是当前线程属于 blockIdx(1,0)它在线程块内的位置是 threadIdx(2,3)它最终映射到矩阵元素 (row3, col6)执行了 C[3][6] A[3][6] B[3][6]这里最核心的两行intcolblockIdx.x*blockDim.xthreadIdx.x;introwblockIdx.y*blockDim.ythreadIdx.y;这就是 grid/block 到矩阵坐标的映射公式。也就是(,)(blockIdx.y⋅blockDim.ythreadIdx.y, blockIdx.x⋅blockDim.xthreadIdx.x)这份例子里的映射关系这里设置的是dim3block(4,4);dim3grid(2,2);所以整个 8×8 矩阵会被分成 4 个 blockblock(0,0) 处理左上角 4×4block(1,0) 处理右上角 4×4block(0,1) 处理左下角 4×4block(1,1) 处理右下角 4×4也就是如果你只是想看最精简版本把 kernel 里的 printf 去掉就行__global__voidmatAddKernel(constfloat*A,constfloat*B,float*C,introws,intcols){intcolblockIdx.x*blockDim.xthreadIdx.x;introwblockIdx.y*blockDim.ythreadIdx.y;if(rowrowscolcols){intidxrow*colscol;C[idx]A[idx]B[idx];}}这样更接近实际工程代码。补一句工程上常见的理解二维矩阵在内存里其实还是一维存放所以要写intidxrow*colscol;这就是把二维坐标 (row, col) 映射成线性地址。下一步如果你愿意我可以直接给你一份“把每个 block 画成矩阵分块图并逐个列出 thread 到元素映射表”的版本那个会更直观。把这个改成一个 markdown 文件下载注文档部分内容可能由 AI 生成