CUDA线程模型与矩阵运算优化实践

1. CUDA线程模型基础解析

在GPU编程领域，理解线程模型是掌握CUDA编程的核心基础。与CPU编程不同，CUDA的线程组织采用层次化结构，这种设计直接决定了并行计算的效率。当我们启动一个CUDA核函数时，会指定一个由多个线程块(block)组成的网格(grid)，而每个线程块又包含多个线程(thread)。这种三维的线程组织结构（虽然可以只使用一维或二维）为处理不同维度的数据提供了天然的适配性。

以常见的矩阵运算为例，假设我们有一个1024x1024的矩阵，可以将其划分为16x16的线程块网格（共64个block），每个线程块包含16x16=256个线程。这种划分方式使得每个线程都能对应处理矩阵中的一个元素，实现真正的数据并行。在实际编程中，我们通过内置变量threadIdx、blockIdx、blockDim和gridDim来访问线程的索引信息，这些变量在核函数内部自动定义，无需手动声明。

关键提示：blockDim.x * gridDim.x 应该等于或大于问题规模，否则会有部分数据未被处理。同时也要注意blockDim不能超过硬件限制（通常为1024）。

2. 矩阵相加的CUDA实现详解

2.1 内存分配与数据传输

矩阵相加作为最基础的线性代数运算，是理解CUDA编程的绝佳起点。首先需要在主机(host)端分配内存存储输入输出矩阵，然后在设备(device)端分配对应的显存空间：

c复制float *h_A, *h_B, *h_C; // 主机端指针
float *d_A, *d_B, *d_C; // 设备端指针

// 主机内存分配
h_A = (float*)malloc(N * N * sizeof(float));
h_B = (float*)malloc(N * N * sizeof(float)); 
h_C = (float*)malloc(N * N * sizeof(float));

// 设备显存分配
cudaMalloc(&d_A, N * N * sizeof(float));
cudaMalloc(&d_B, N * N * sizeof(float));
cudaMalloc(&d_C, N * N * sizeof(float));

// 数据传输到设备
cudaMemcpy(d_A, h_A, N * N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, N * N * sizeof(float), cudaMemcpyHostToDevice);

2.2 核函数设计与实现

矩阵相加的核函数相对简单，每个线程负责计算一个对应位置的元素和：

c复制__global__ void matrixAdd(float *A, float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (i < N && j < N) {
        int idx = i * N + j;
        C[idx] = A[idx] + B[idx];
    }
}

这里有几个关键点需要注意：

1. 使用二维线程组织更符合矩阵运算的直观理解