CUDA 线程索引计算详解__global__ void naive_gemm(float *A, float *B, float *C, int N){int rowblockIdx.y * blockDim.y threadIdx.y;int colblockIdx.x * blockDim.x threadIdx.x;floatsum0.0f;for(int k0;kN;k){sumA[row * N k]* B[k * N col];// 每次访问全局内存}C[row * N col]sum;}这里的row 和col为什么等于这个是如何计算的让我详细解释这段代码中row和col的计算方式。一、CUDA 线程层次结构可以参考笔者的cutlass学习笔记https://github.com/shizhengLi/cutlass-learning1.1 基本概念Grid网格 ├── Block (0,0) Block (0,1) Block (0,2) ... ├── Block (1,0) Block (1,1) Block (1,2) ... └── Block (2,0) Block (2,1) Block (2,2) ... ↓ 每个 Block 包含多个 Thread Block (0,0) ├── Thread (0,0) Thread (0,1) Thread (0,2) ... ├── Thread (1,0) Thread (1,1) Thread (1,2) ... └── Thread (2,0) Thread (2,1) Thread (2,2) ...1.2 内置变量// Grid 级别Block 的索引blockIdx.x// 当前 Block 在 Grid 中的 X 坐标blockIdx.y// 当前 Block 在 Grid 中的 Y 坐标blockIdx.z// 当前 Block 在 Grid 中的 Z 坐标// Block 级别Thread 的索引threadIdx.x// 当前 Thread 在 Block 中的 X 坐标threadIdx.y// 当前 Thread 在 Block 中的 Y 坐标threadIdx.z// 当前 Thread 在 Block 中的 Z 坐标// Block 维度blockDim.x// Block 在 X 方向的线程数blockDim.y// Block 在 Y 方向的线程数blockDim.z// Block 在 Z 方向的线程数// Grid 维度gridDim.x// Grid 在 X 方向的 Block 数gridDim.y// Grid 在 Y 方向的 Block 数gridDim.z// Grid 在 Z 方向的 Block 数二、row和col计算详解2.1 公式分解introwblockIdx.y*blockDim.ythreadIdx.y;└─────┬─────┘ └───┬───┘ └────┬────┘ Block 偏移 Block 大小 线程偏移intcolblockIdx.x*blockDim.xthreadIdx.x;计算逻辑全局索引 Block索引 × Block大小 线程在Block内的索引2.2 具体例子假设我们要计算一个1024×1024的矩阵乘法// Kernel 启动配置dim3blockDim(16,16);// 每个 Block: 16×16 256 个线程dim3gridDim(64,64);// Grid: 64×64 4096 个 Block// 总线程数: 256 × 4096 1,048,576 个线程正好覆盖 1024×1024naive_gemmgridDim,blockDim(A,B,C,1024);可视化计算Grid 布局64×64 个 Block ┌─────────────────────────────────────┐ │ (0,0) (1,0) (2,0) ... (63,0) │ ← blockIdx.y 0 │ (0,1) (1,1) (2,1) ... (63,1) │ ← blockIdx.y 1 │ (0,2) (1,2) (2,2) ... (63,2) │ ← blockIdx.y 2 │ ... ... ... ... ... │ │ (0,63) (1,63) (2,63) ... (63,63) │ ← blockIdx.y 63 └─────────────────────────────────────┘ ↑ ↑ ↑ ↑ blockIdx.x 0, 1, 2, ..., 63 每个 Block 内部16×16 个 Thread Block (2, 3) → blockIdx.x 2, blockIdx.y 3 ┌────────────────────────────────────┐ │ (0,0) (1,0) (2,0) ... (15,0) │ ← threadIdx.y 0 │ (0,1) (1,1) (2,1) ... (15,1) │ ← threadIdx.y 1 │ (0,2) (1,2) (2,2) ... (15,2) │ ← threadIdx.y 2 │ ... ... ... ... ... │ │ (0,15) (1,15) (2,15) ... (15,15) │ ← threadIdx.y 15 └────────────────────────────────────┘ ↑ ↑ ↑ ↑ threadIdx.x 0, 1, 2, ..., 152.3 具体线程的计算例子 1Block (2, 3) 中的 Thread (5, 7)// 已知blockIdx.x2,blockIdx.y3threadIdx.x5,threadIdx.y7blockDim.x16,blockDim.y16// 计算 row对应矩阵的行索引rowblockIdx.y*blockDim.ythreadIdx.y3*16748755// 计算 col对应矩阵的列索引colblockIdx.x*blockDim.xthreadIdx.x2*16532537// 结论这个线程负责计算矩阵 C[55][37]例子 2Block (0, 0) 中的 Thread (0, 0)// 这是第一个 Block 的第一个线程blockIdx.x0,blockIdx.y0threadIdx.x0,threadIdx.y0row0*1600col0*1600// 负责计算 C[0][0]左上角元素例子 3Block (63, 63) 中的 Thread (15, 15)// 这是最后一个 Block 的最后一个线程blockIdx.x63,blockIdx.y63threadIdx.x15,threadIdx.y15row63*16151023col63*16151023// 负责计算 C[1023][1023]右下角元素三、完整可视化示例3.1 小规模示例更直观假设矩阵64×64使用8×8的 Blockdim3blockDim(8,8);// 每个 Block: 8×8 64 个线程dim3gridDim(8,8);// Grid: 8×8 64 个 Block// 总共 64 × 64 4096 个线程覆盖 64×64 矩阵矩阵 C 与线程的映射矩阵 C (64×64) ┌──────────────────────────────────────────────────┐ │ Block(0,0) │ Block(1,0) │ ...│ │ ┌──────────────────┐ │ ┌──────────────────┐ │ │ │ │ T(0,0) T(1,0)... │ │ │ T(0,0) T(1,0)... │ │ │ │ │ T(0,1) T(1,1)... │ │ │ T(0,1) T(1,1)... │ │ │ │ │ ... ... ... │ │ │ ... ... ... │ │ │ │ │ T(0,7) T(1,7)... │ │ │ T(0,7) T(1,7)... │ │ │ │ └──────────────────┘ │ └──────────────────┘ │ │ │ 0-7 行 │ 0-7 行 │ │ │ 0-7 列 │ 8-15 列 │ │ ├──────────────────────┼──────────────────────┼───┤ │ Block(0,1) │ Block(1,1) │ ...│ │ ┌──────────────────┐ │ ┌──────────────────┐ │ │ │ │ T(0,0) T(1,0)... │ │ │ T(0,0) T(1,0)... │ │ │ │ │ ... ... ... │ │ │ ... ... ... │ │ │ │ └──────────────────┘ │ └──────────────────┘ │ │ │ 8-15 行 │ 8-15 行 │ │ │ 0-7 列 │ 8-15 列 │ │ └──────────────────────┴──────────────────────┴───┘具体计算示例// Block (1, 1), Thread (3, 2)blockIdx.x1,blockIdx.y1threadIdx.x3,threadIdx.y2blockDim.x8,blockDim.y8row1*8210col1*8311// 这个线程负责计算 C[10][11]// 验证// - Block (1,1) 负责矩阵的 [8:15, 8:15] 区域// - Thread (3,2) 在这个区域的相对位置是 (col3, row2)// - 全局位置行 8210, 列 8311 ✓3.2 代码追踪__global__voidnaive_gemm(float*A,float*B,float*C,intN){// 假设Block (2, 3), Thread (5, 7), blockDim (16, 16)introwblockIdx.y*blockDim.ythreadIdx.y;// 3 * 16 7// 55intcolblockIdx.x*blockDim.xthreadIdx.x;// 2 * 16 5// 37// 这个线程计算 C[55][37] A[55][:] · B[:][37]floatsum0.0f;for(intk0;kN;k){// k0: sum A[55*N0] * B[0*N37]// k1: sum A[55*N1] * B[1*N37]// ...// kN-1: sum A[55*N(N-1)] * B[(N-1)*N37]sumA[row*Nk]*B[k*Ncol];}// 将结果写入 C[55][37]C[row*Ncol]sum;}四、为什么这样设计4.1 映射到物理硬件// 二维索引映射row,col → 矩阵元素 C[row][col]// GPU 物理组织SM(Streaming Multiprocessor)└─ 调度多个 Block └─ 每个 Block 包含多个Warp(32threads)└─ 每个 Thread 执行一个元素的计算4.2 优势1. 简单直观// 一维索引需要手动转换inttidblockIdx.x*blockDim.xthreadIdx.x;introwtid/N;intcoltid%N;// 二维索引直接对应introwblockIdx.y*blockDim.ythreadIdx.y;intcolblockIdx.x*blockDim.xthreadIdx.x;2. 内存访问模式// 同一个 Block 内的线程// threadIdx.x: 0, 1, 2, 3, ..., 15// 访问 C[row][col], C[row][col1], C[row][col2], ...// → 连续访问合并内存事务Coalesced Access3. 易于调试if(row10col20){printf(Thread computing C[10][20]\n);printf(blockIdx(%d,%d), threadIdx(%d,%d)\n,blockIdx.x,blockIdx.y,threadIdx.x,threadIdx.y);}五、常见变体5.1 一维线程映射__global__voidkernel_1d(float*data,intN){inttidblockIdx.x*blockDim.xthreadIdx.x;// └───────┬───────┘ └─────┬─────┘// Block 偏移 线程偏移if(tidN){data[tid]tid*2.0f;}}// 调用intN10000;intthreads256;intblocks(Nthreads-1)/threads;// 向上取整kernel_1dblocks,threads(data,N);5.2 三维线程映射__global__voidkernel_3d(float*volume,intX,intY,intZ){intxblockIdx.x*blockDim.xthreadIdx.x;intyblockIdx.y*blockDim.ythreadIdx.y;intzblockIdx.z*blockDim.zthreadIdx.z;if(xXyYzZ){intidxz*(X*Y)y*Xx;// 3D → 1Dvolume[idx]xyz;}}// 调用处理 256×256×256 的体数据dim3blockDim(8,8,8);dim3gridDim(32,32,32);kernel_3dgridDim,blockDim(volume,256,256,256);5.3 Stride 循环处理大数组__global__voidkernel_stride(float*data,intN){// 总线程数inttotal_threadsgridDim.x*blockDim.x;// 起始索引inttidblockIdx.x*blockDim.xthreadIdx.x;// 每个线程处理多个元素间隔 total_threadsfor(intitid;iN;itotal_threads){data[i]i*2.0f;}}// 调用用 1000 个线程处理 1,000,000 个元素kernel_stride10,100(data,1000000);// 每个线程处理约 1000 个元素六、调试技巧6.1 打印线程信息__global__voiddebug_kernel(){introwblockIdx.y*blockDim.ythreadIdx.y;intcolblockIdx.x*blockDim.xthreadIdx.x;// 只打印部分线程的信息避免输出过多if(blockIdx.x0blockIdx.y0threadIdx.x4threadIdx.y4){printf(Block(%d,%d) Thread(%d,%d) → row%d, col%d\n,blockIdx.x,blockIdx.y,threadIdx.x,threadIdx.y,row,col);}}// 输出// Block(0,0) Thread(0,0) → row0, col0// Block(0,0) Thread(1,0) → row0, col1// Block(0,0) Thread(2,0) → row0, col2// Block(0,0) Thread(3,0) → row0, col3// Block(0,0) Thread(0,1) → row1, col0// ...6.2 可视化工具// 使用 CUDA-GDB$ cuda-gdb./program(cuda-gdb)breaknaive_gemm(cuda-gdb)run(cuda-gdb)cudathread(0,0,0),(0,0,0)# 切换到Block(0,0)Thread(0,0)(cuda-gdb)printrow(cuda-gdb)print col// 使用 Nsight Compute$ ncu--set full-o profile./program # 在 GUI 中查看每个线程的执行情况七、总结公式记忆全局索引Block索引 × Block大小Block内线程索引 rowblockIdx.y*blockDim.ythreadIdx.y colblockIdx.x*blockDim.xthreadIdx.x核心思想分层映射Grid → Block → Thread → 矩阵元素并行化每个线程独立计算一个输出元素简单直观二维索引直接对应矩阵坐标记忆技巧blockIdx确定在哪个区域粗粒度threadIdx确定在区域内的具体位置细粒度blockDim是桥梁连接两个层次希望这个详细的解释能帮助你理解 CUDA 线程索引的计算方式补充矩阵乘法核心公式// 数学C[i][j]Σ(k0to K-1)A[i][k]× B[k][j]// 代码sum0;for(k0;kN;k){sumA[row*Nk]*B[k*Ncol];└────┬────┘ └────┬────┘ A[row][k]B[k][col]}C[row*Ncol]sum;关键点row * N k访问 A 矩阵第 row 行第 k 列k * N col访问 B 矩阵第 k 行第 col 列内积计算固定行和列遍历 k 维度求和内存布局行主序二维索引转一维偏移这就是矩阵乘法的本质两个向量的点积 → A 的行 · B 的列后记2026年8月12日于上海在claude opus 4.8辅助下完成。 CUDA 线程索引计算详解
CUDA 线程索引计算详解__global__ void naive_gemm(float *A, float *B, float *C, int N){int rowblockIdx.y * blockDim.y threadIdx.y;int colblockIdx.x * blockDim.x threadIdx.x;floatsum0.0f;for(int k0;kN;k){sumA[row * N k]* B[k * N col];// 每次访问全局内存}C[row * N col]sum;}这里的row 和col为什么等于这个是如何计算的让我详细解释这段代码中row和col的计算方式。一、CUDA 线程层次结构可以参考笔者的cutlass学习笔记https://github.com/shizhengLi/cutlass-learning1.1 基本概念Grid网格 ├── Block (0,0) Block (0,1) Block (0,2) ... ├── Block (1,0) Block (1,1) Block (1,2) ... └── Block (2,0) Block (2,1) Block (2,2) ... ↓ 每个 Block 包含多个 Thread Block (0,0) ├── Thread (0,0) Thread (0,1) Thread (0,2) ... ├── Thread (1,0) Thread (1,1) Thread (1,2) ... └── Thread (2,0) Thread (2,1) Thread (2,2) ...1.2 内置变量// Grid 级别Block 的索引blockIdx.x// 当前 Block 在 Grid 中的 X 坐标blockIdx.y// 当前 Block 在 Grid 中的 Y 坐标blockIdx.z// 当前 Block 在 Grid 中的 Z 坐标// Block 级别Thread 的索引threadIdx.x// 当前 Thread 在 Block 中的 X 坐标threadIdx.y// 当前 Thread 在 Block 中的 Y 坐标threadIdx.z// 当前 Thread 在 Block 中的 Z 坐标// Block 维度blockDim.x// Block 在 X 方向的线程数blockDim.y// Block 在 Y 方向的线程数blockDim.z// Block 在 Z 方向的线程数// Grid 维度gridDim.x// Grid 在 X 方向的 Block 数gridDim.y// Grid 在 Y 方向的 Block 数gridDim.z// Grid 在 Z 方向的 Block 数二、row和col计算详解2.1 公式分解introwblockIdx.y*blockDim.ythreadIdx.y;└─────┬─────┘ └───┬───┘ └────┬────┘ Block 偏移 Block 大小 线程偏移intcolblockIdx.x*blockDim.xthreadIdx.x;计算逻辑全局索引 Block索引 × Block大小 线程在Block内的索引2.2 具体例子假设我们要计算一个1024×1024的矩阵乘法// Kernel 启动配置dim3blockDim(16,16);// 每个 Block: 16×16 256 个线程dim3gridDim(64,64);// Grid: 64×64 4096 个 Block// 总线程数: 256 × 4096 1,048,576 个线程正好覆盖 1024×1024naive_gemmgridDim,blockDim(A,B,C,1024);可视化计算Grid 布局64×64 个 Block ┌─────────────────────────────────────┐ │ (0,0) (1,0) (2,0) ... (63,0) │ ← blockIdx.y 0 │ (0,1) (1,1) (2,1) ... (63,1) │ ← blockIdx.y 1 │ (0,2) (1,2) (2,2) ... (63,2) │ ← blockIdx.y 2 │ ... ... ... ... ... │ │ (0,63) (1,63) (2,63) ... (63,63) │ ← blockIdx.y 63 └─────────────────────────────────────┘ ↑ ↑ ↑ ↑ blockIdx.x 0, 1, 2, ..., 63 每个 Block 内部16×16 个 Thread Block (2, 3) → blockIdx.x 2, blockIdx.y 3 ┌────────────────────────────────────┐ │ (0,0) (1,0) (2,0) ... (15,0) │ ← threadIdx.y 0 │ (0,1) (1,1) (2,1) ... (15,1) │ ← threadIdx.y 1 │ (0,2) (1,2) (2,2) ... (15,2) │ ← threadIdx.y 2 │ ... ... ... ... ... │ │ (0,15) (1,15) (2,15) ... (15,15) │ ← threadIdx.y 15 └────────────────────────────────────┘ ↑ ↑ ↑ ↑ threadIdx.x 0, 1, 2, ..., 152.3 具体线程的计算例子 1Block (2, 3) 中的 Thread (5, 7)// 已知blockIdx.x2,blockIdx.y3threadIdx.x5,threadIdx.y7blockDim.x16,blockDim.y16// 计算 row对应矩阵的行索引rowblockIdx.y*blockDim.ythreadIdx.y3*16748755// 计算 col对应矩阵的列索引colblockIdx.x*blockDim.xthreadIdx.x2*16532537// 结论这个线程负责计算矩阵 C[55][37]例子 2Block (0, 0) 中的 Thread (0, 0)// 这是第一个 Block 的第一个线程blockIdx.x0,blockIdx.y0threadIdx.x0,threadIdx.y0row0*1600col0*1600// 负责计算 C[0][0]左上角元素例子 3Block (63, 63) 中的 Thread (15, 15)// 这是最后一个 Block 的最后一个线程blockIdx.x63,blockIdx.y63threadIdx.x15,threadIdx.y15row63*16151023col63*16151023// 负责计算 C[1023][1023]右下角元素三、完整可视化示例3.1 小规模示例更直观假设矩阵64×64使用8×8的 Blockdim3blockDim(8,8);// 每个 Block: 8×8 64 个线程dim3gridDim(8,8);// Grid: 8×8 64 个 Block// 总共 64 × 64 4096 个线程覆盖 64×64 矩阵矩阵 C 与线程的映射矩阵 C (64×64) ┌──────────────────────────────────────────────────┐ │ Block(0,0) │ Block(1,0) │ ...│ │ ┌──────────────────┐ │ ┌──────────────────┐ │ │ │ │ T(0,0) T(1,0)... │ │ │ T(0,0) T(1,0)... │ │ │ │ │ T(0,1) T(1,1)... │ │ │ T(0,1) T(1,1)... │ │ │ │ │ ... ... ... │ │ │ ... ... ... │ │ │ │ │ T(0,7) T(1,7)... │ │ │ T(0,7) T(1,7)... │ │ │ │ └──────────────────┘ │ └──────────────────┘ │ │ │ 0-7 行 │ 0-7 行 │ │ │ 0-7 列 │ 8-15 列 │ │ ├──────────────────────┼──────────────────────┼───┤ │ Block(0,1) │ Block(1,1) │ ...│ │ ┌──────────────────┐ │ ┌──────────────────┐ │ │ │ │ T(0,0) T(1,0)... │ │ │ T(0,0) T(1,0)... │ │ │ │ │ ... ... ... │ │ │ ... ... ... │ │ │ │ └──────────────────┘ │ └──────────────────┘ │ │ │ 8-15 行 │ 8-15 行 │ │ │ 0-7 列 │ 8-15 列 │ │ └──────────────────────┴──────────────────────┴───┘具体计算示例// Block (1, 1), Thread (3, 2)blockIdx.x1,blockIdx.y1threadIdx.x3,threadIdx.y2blockDim.x8,blockDim.y8row1*8210col1*8311// 这个线程负责计算 C[10][11]// 验证// - Block (1,1) 负责矩阵的 [8:15, 8:15] 区域// - Thread (3,2) 在这个区域的相对位置是 (col3, row2)// - 全局位置行 8210, 列 8311 ✓3.2 代码追踪__global__voidnaive_gemm(float*A,float*B,float*C,intN){// 假设Block (2, 3), Thread (5, 7), blockDim (16, 16)introwblockIdx.y*blockDim.ythreadIdx.y;// 3 * 16 7// 55intcolblockIdx.x*blockDim.xthreadIdx.x;// 2 * 16 5// 37// 这个线程计算 C[55][37] A[55][:] · B[:][37]floatsum0.0f;for(intk0;kN;k){// k0: sum A[55*N0] * B[0*N37]// k1: sum A[55*N1] * B[1*N37]// ...// kN-1: sum A[55*N(N-1)] * B[(N-1)*N37]sumA[row*Nk]*B[k*Ncol];}// 将结果写入 C[55][37]C[row*Ncol]sum;}四、为什么这样设计4.1 映射到物理硬件// 二维索引映射row,col → 矩阵元素 C[row][col]// GPU 物理组织SM(Streaming Multiprocessor)└─ 调度多个 Block └─ 每个 Block 包含多个Warp(32threads)└─ 每个 Thread 执行一个元素的计算4.2 优势1. 简单直观// 一维索引需要手动转换inttidblockIdx.x*blockDim.xthreadIdx.x;introwtid/N;intcoltid%N;// 二维索引直接对应introwblockIdx.y*blockDim.ythreadIdx.y;intcolblockIdx.x*blockDim.xthreadIdx.x;2. 内存访问模式// 同一个 Block 内的线程// threadIdx.x: 0, 1, 2, 3, ..., 15// 访问 C[row][col], C[row][col1], C[row][col2], ...// → 连续访问合并内存事务Coalesced Access3. 易于调试if(row10col20){printf(Thread computing C[10][20]\n);printf(blockIdx(%d,%d), threadIdx(%d,%d)\n,blockIdx.x,blockIdx.y,threadIdx.x,threadIdx.y);}五、常见变体5.1 一维线程映射__global__voidkernel_1d(float*data,intN){inttidblockIdx.x*blockDim.xthreadIdx.x;// └───────┬───────┘ └─────┬─────┘// Block 偏移 线程偏移if(tidN){data[tid]tid*2.0f;}}// 调用intN10000;intthreads256;intblocks(Nthreads-1)/threads;// 向上取整kernel_1dblocks,threads(data,N);5.2 三维线程映射__global__voidkernel_3d(float*volume,intX,intY,intZ){intxblockIdx.x*blockDim.xthreadIdx.x;intyblockIdx.y*blockDim.ythreadIdx.y;intzblockIdx.z*blockDim.zthreadIdx.z;if(xXyYzZ){intidxz*(X*Y)y*Xx;// 3D → 1Dvolume[idx]xyz;}}// 调用处理 256×256×256 的体数据dim3blockDim(8,8,8);dim3gridDim(32,32,32);kernel_3dgridDim,blockDim(volume,256,256,256);5.3 Stride 循环处理大数组__global__voidkernel_stride(float*data,intN){// 总线程数inttotal_threadsgridDim.x*blockDim.x;// 起始索引inttidblockIdx.x*blockDim.xthreadIdx.x;// 每个线程处理多个元素间隔 total_threadsfor(intitid;iN;itotal_threads){data[i]i*2.0f;}}// 调用用 1000 个线程处理 1,000,000 个元素kernel_stride10,100(data,1000000);// 每个线程处理约 1000 个元素六、调试技巧6.1 打印线程信息__global__voiddebug_kernel(){introwblockIdx.y*blockDim.ythreadIdx.y;intcolblockIdx.x*blockDim.xthreadIdx.x;// 只打印部分线程的信息避免输出过多if(blockIdx.x0blockIdx.y0threadIdx.x4threadIdx.y4){printf(Block(%d,%d) Thread(%d,%d) → row%d, col%d\n,blockIdx.x,blockIdx.y,threadIdx.x,threadIdx.y,row,col);}}// 输出// Block(0,0) Thread(0,0) → row0, col0// Block(0,0) Thread(1,0) → row0, col1// Block(0,0) Thread(2,0) → row0, col2// Block(0,0) Thread(3,0) → row0, col3// Block(0,0) Thread(0,1) → row1, col0// ...6.2 可视化工具// 使用 CUDA-GDB$ cuda-gdb./program(cuda-gdb)breaknaive_gemm(cuda-gdb)run(cuda-gdb)cudathread(0,0,0),(0,0,0)# 切换到Block(0,0)Thread(0,0)(cuda-gdb)printrow(cuda-gdb)print col// 使用 Nsight Compute$ ncu--set full-o profile./program # 在 GUI 中查看每个线程的执行情况七、总结公式记忆全局索引Block索引 × Block大小Block内线程索引 rowblockIdx.y*blockDim.ythreadIdx.y colblockIdx.x*blockDim.xthreadIdx.x核心思想分层映射Grid → Block → Thread → 矩阵元素并行化每个线程独立计算一个输出元素简单直观二维索引直接对应矩阵坐标记忆技巧blockIdx确定在哪个区域粗粒度threadIdx确定在区域内的具体位置细粒度blockDim是桥梁连接两个层次希望这个详细的解释能帮助你理解 CUDA 线程索引的计算方式补充矩阵乘法核心公式// 数学C[i][j]Σ(k0to K-1)A[i][k]× B[k][j]// 代码sum0;for(k0;kN;k){sumA[row*Nk]*B[k*Ncol];└────┬────┘ └────┬────┘ A[row][k]B[k][col]}C[row*Ncol]sum;关键点row * N k访问 A 矩阵第 row 行第 k 列k * N col访问 B 矩阵第 k 行第 col 列内积计算固定行和列遍历 k 维度求和内存布局行主序二维索引转一维偏移这就是矩阵乘法的本质两个向量的点积 → A 的行 · B 的列后记2026年8月12日于上海在claude opus 4.8辅助下完成。