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 = 8

y方向：2 × 4 = 8

刚好覆盖整个矩阵。

三、核心映射公式（最重要）

线程 → 矩阵坐标：

intcol=blockIdx.x*blockDim.x+threadIdx.x;introw=blockIdx.y*blockDim.y+threadIdx.y;

👉 对应关系：

col = 𝐵ₓ ⋅ 𝑏ₓ + 𝑡ₓ

row = 𝐵ᵧ ⋅ 𝑏ᵧ + 𝑡ᵧ

四、直观展开（关键理解）

我们把整个 8×8 矩阵分成 4 个 block 区域

Block 分布

五、一个 block 内的映射

举例：block(1,0)

blockIdx.x = 1

blockIdx.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：

七、矩阵加法的 CUDA 表达

每个线程做一件事：

C[row][col]=A[row][col]+B[row][col];

👉 数学表达：

𝐶ᵢⱼ = 𝐴ᵢⱼ + 𝐵ᵢⱼ

👉 CUDA 展开后本质：

𝐶(𝐵ᵧ𝑏ᵧ+𝑡ᵧ, 𝐵ₓ𝑏ₓ+𝑡ₓ) = 𝐴... + 𝐵...

八、完整 CUDA Kernel 示例

__global__voidmatAdd(float*A,float*B,float*C,intN){intcol=blockIdx.x*blockDim.x+threadIdx.x;introw=blockIdx.y*blockDim.y+threadIdx.y;if(row<N&&col<N){intidx=row*N+col;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 编译运行。

完整代码

#include<iostream>#include<vector>#include<cuda_runtime.h>#defineCHECK_CUDA(call)\do{\cudaError_terr=call;\if(err!=cudaSuccess){\std::cerr<<"CUDA error at "<<__FILE__<<":"<<__LINE__\<<" code="<<err\<<" \""<<cudaGetErrorString(err)<<"\""<<std::endl;\exit(EXIT_FAILURE);\}\}while(0)// 打印矩阵voidprintMatrix(conststd::vector<float>&mat,introws,intcols,constchar*name){std::cout<<name<<" ="<<std::endl;for(intr=0;r<rows;r++){for(intc=0;c<cols;c++){std::cout<<mat[r*cols+c]<<"\t";}std::cout<<std::endl;}std::cout<<std::endl;}// CUDA Kernel: 矩阵加法__global__voidmatAddKernel(constfloat*A,constfloat*B,float*C,introws,intcols){// block 内线程坐标 + block 坐标 -> 全局矩阵坐标intcol=blockIdx.x*blockDim.x+threadIdx.x;// 列introw=blockIdx.y*blockDim.y+threadIdx.y;// 行if(row<rows&&col<cols){intidx=row*cols+col;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(){// 矩阵大小constintrows=8;constintcols=8;constintsize=rows*cols;constsize_tbytes=size*sizeof(float);// Host 端矩阵std::vector<float>h_A(size);std::vector<float>h_B(size);std::vector<float>h_C(size,0.0f);// 初始化矩阵// A: 0,1,2,3...// B: 全部 100for(intr=0;r<rows;r++){for(intc=0;c<cols;c++){intidx=r*cols+c;h_A[idx]=static_cast<float>(idx);h_B[idx]=100.0f;}}printMatrix(h_A,rows,cols,"A");printMatrix(h_B,rows,cols,"B");// Device 指针float*d_A=nullptr,*d_B=nullptr,*d_C=nullptr;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((cols+block.x-1)/block.x,(rows+block.y-1)/block.y);std::cout<<"Launch config:"<<std::endl;std::cout<<"block = ("<<block.x<<", "<<block.y<<", "<<block.z<<")"<<std::endl;std::cout<<"grid = ("<<grid.x<<", "<<grid.y<<", "<<grid.z<<")"<<std::endl;std::cout<<std::endl;// 启动 KernelmatAddKernel<<<grid,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 ok=true;for(inti=0;i<size;i++){floatexpected=h_A[i]+h_B[i];if(h_C[i]!=expected){std::cerr<<"Mismatch at index "<<i<<", got "<<h_C[i]<<", expected "<<expected<<std::endl;ok=false;break;}}if(ok){std::cout<<"Result check passed."<<std::endl;}else{std::cout<<"Result 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), A=30.0, B=100.0, C=130.0

这句话的意思就是：

当前线程属于 blockIdx=(1,0)

它在线程块内的位置是 threadIdx=(2,3)

它最终映射到矩阵元素 (row=3, col=6)

执行了 C[3][6] = A[3][6] + B[3][6]

这里最核心的两行

intcol=blockIdx.x*blockDim.x+threadIdx.x;introw=blockIdx.y*blockDim.y+threadIdx.y;

这就是 grid/block 到矩阵坐标的映射公式。

也就是：

(𝑟𝑜𝑤,𝑐𝑜𝑙)=(blockIdx.y⋅blockDim.y+threadIdx.y, blockIdx.x⋅blockDim.x+threadIdx.x)

这份例子里的映射关系

这里设置的是：

dim3block(4,4);dim3grid(2,2);

所以整个 8×8 矩阵会被分成 4 个 block：

block(0,0) 处理左上角 4×4

block(1,0) 处理右上角 4×4

block(0,1) 处理左下角 4×4

block(1,1) 处理右下角 4×4

也就是：

如果你只是想看最精简版本

把 kernel 里的 printf 去掉就行：

__global__voidmatAddKernel(constfloat*A,constfloat*B,float*C,introws,intcols){intcol=blockIdx.x*blockDim.x+threadIdx.x;introw=blockIdx.y*blockDim.y+threadIdx.y;if(row<rows&&col<cols){intidx=row*cols+col;C[idx]=A[idx]+B[idx];}}

这样更接近实际工程代码。

补一句工程上常见的理解

二维矩阵在内存里其实还是一维存放，所以要写：

intidx=row*cols+col;

这就是把二维坐标 (row, col) 映射成线性地址。

下一步如果你愿意，我可以直接给你一份：

“把每个 block 画成矩阵分块图，并逐个列出 thread 到元素映射表”的版本，那个会更直观。把这个改成一个 markdown 文件下载

（注：文档部分内容可能由 AI 生成）