Matrix Transpose
338 字
2 分钟
Matrix Transpose
题面
答案
#include <cuda_runtime.h>
__global__ void matrix_transpose_kernel(const float* input, float* output, int rows, int cols) { int row = blockDim.y * blockIdx.y + threadIdx.y; int col = blockDim.x * blockIdx.x + threadIdx.x;
if (row >= rows || col >= cols) return;
int input_index = row * cols + col; int output_index = col * rows + row;
output[output_index] = input[input_index];}
// input, output are device pointers (i.e. pointers to memory on the GPU)extern "C" void solve(const float* input, float* output, int rows, int cols) { dim3 threadsPerBlock(16, 16); dim3 blocksPerGrid((cols + threadsPerBlock.x - 1) / threadsPerBlock.x, (rows + threadsPerBlock.y - 1) / threadsPerBlock.y);
matrix_transpose_kernel<<<blocksPerGrid, threadsPerBlock>>>(input, output, rows, cols); cudaDeviceSynchronize();}两级组织：Grid → Block → Thread
CUDA 线程按两级结构组织，每个线程的全局坐标由 三部分拼成：
row = blockIdx.y × blockDim.y + threadIdx.ycol = blockIdx.x × blockDim.x + threadIdx.x| 部分 | 含义 | 类比 |
|---|---|---|
blockIdx | 你在第几个块 | 第几辆车 |
blockDim | 每块多少线程 | 每辆车多少座位 |
threadIdx | 你在块内排第几 | 车里的第几排第几座 |
示例计算
假设 blockDim = (16, 16)，某线程在 Block(1,0) 中处于 threadIdx = (2, 3)：
col = 1×16 + 2 = 18 ← 前面 1 整块(16列) + 块内第 2 个row = 0×16 + 3 = 3 ← 前面 0 整块 + 块内第 3 个→ 该线程负责计算 output 的 [18][3]（对应 input 的 [3][18]）。
转置中的应用
int input_index = row * cols + col; // input[row][col]int output_index = col * rows + row; // output[col][row] ← 行列互换每个线程读一个元素，写一个元素，无需同步——天然并行。
Matrix Transpose
https://dongyanzhang.com/posts/leetgpu/matrix-transpose/