leetGPU解题笔记（1）

1.题面

题目要求

向量加法
实现一个程序，在GPU上对两个包含32位浮点数的向量执行逐元素加法。该程序应接受两个长度相等的输入向量，并生成一个包含它们和的输出向量。

实现要求

禁止使用外部库
solve函数签名必须保持不变
最终结果必须存储在向量C中
示例1：
输入：A = [1.0, 2.0, 3.0, 4.0]
B = [5.0, 6.0, 7.0, 8.0]
输出：C = [6.0, 8.0, 10.0, 12.0]
示例2：
输入：A = [1.5, 1.5, 1.5]
B = [2.3, 2.3, 2.3]
输出：C = [3.8, 3.8, 3.8]

约束条件

输入向量A和B长度相同
1 ≤ N ≤ 100,000,000

2.已有代码解析

函数参数与目的

void solve(const float* A, const float* B, float* C, int N)

• 该函数接收三个指向GPU内存的指针：A和B是输入向量，C是输出向量。
• N表示向量的长度，即元素个数。

线程与线程块配置

int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

• 线程块大小(threadsPerBlock)：每个线程块包含256个线程。这是CUDA编程中常用的配置，适合大多数GPU架构。
• 网格大小(blocksPerGrid)：根据向量长度N计算所需的线程块数量。使用向上取整公式确保所有元素都被处理。例如：
  • 当N=1000时，blocksPerGrid = (1000 + 256 - 1) / 256 = 4。
  • 每256个线程就打包成一个block，blocksPerGrid就表示打包了多少个

核函数调用

vector_add<<<blocksPerGrid, threadsPerBlock>>>(A, B, C, N);

• 核函数(vector_add)：这是一个在GPU上执行的函数，负责逐元素地将A和B相加，结果存入C。
• 执行配置(<<<...>>>)：指定网格和线程块的维度。这里使用一维网格和一维线程块。
• 核函数会被并行调用，每个线程负责计算一个元素：C[i] = A[i] + B[i]。

同步操作

cudaDeviceSynchronize();

• 确保核函数执行完成后才返回。这对于需要立即使用结果的场景很重要。

3.QA环节

Q1: Grid是什么，为什么bocksPerGrid就是块的数目？

Grid是CUDA编程模型中的一个概念，表示线程块(Block)的集合。在GPU并行计算中，一个Grid是核函数(Kernel)的一次执行实例，它包含多个线程块，这些线程块可以并行或串行执行，具体取决于硬件资源。

blocksPerGrid是根据向量长度N和每个线程块的线程数量threadsPerBlock计算得出的。在CUDA中，网格(Grid)和线程块(Block)都可以是一维、二维或三维的。在这个向量加法的例子中，我们使用了一维网格，因此blocksPerGrid直接表示网格中一维方向上的线程块数量。计算公式为：

blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

这个公式实现了向上取整的效果，确保所有N个元素都有对应的线程来处理。例如：

• 当N=1000，threadsPerBlock=256时，blocksPerGrid = (1000 + 256 - 1) / 256 = 4
• 每个线程块处理256个元素，4个线程块总共可以处理1024个元素，覆盖了N=1000的需求。

Q2: 这里使用“一维网格”和“一维线程块”是什么含义，还有其他选择吗，比如二维的？

一维网格和一维线程块

在CUDA中，线程组织是分级的：

• 线程块(Block)：由多个线程组成，可以是一维、二维或三维
• 网格(Grid)：由多个线程块组成，同样可以是一维、二维或三维

在向量加法的例子中：

• 一维线程块：每个线程块中的线程按一维数组排列，线程索引由threadIdx.x表示
• 一维网格：线程块按一维数组排列，线程块索引由blockIdx.x表示

每个线程的全局索引计算方式为：

int i = blockIdx.x * blockDim.x + threadIdx.x;

其中：

• blockDim.x是每个线程块的线程数量(256)
• blockIdx.x是当前线程块的索引
• threadIdx.x是当前线程在线程块内的索引

二维网格和二维线程块

是的，CUDA支持二维甚至三维的网格和线程块，适用于处理二维或三维数据结构(如图像、矩阵)。例如：

二维线程块

对于图像处理，可以使用二维线程块：

dim3 threadsPerBlock(16, 16);  // 每个线程块有16×16=256个线程

线程索引由threadIdx.x和threadIdx.y表示，线程块内的线程按二维矩阵排列。

二维网格

对应的网格也可以是二维的：

dim3 blocksPerGrid((width + threadsPerBlock.x - 1) / threadsPerBlock.x,(height + threadsPerBlock.y - 1) / threadsPerBlock.y
);

线程块索引由blockIdx.x和blockIdx.y表示，网格中的线程块按二维矩阵排列。

全局索引计算

对于二维网格和二维线程块，全局索引计算方式为：

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int idx = row * width + col;  // 转换为一维索引

为什么向量加法使用一维组织？

向量加法处理的是一维数组，每个元素的计算只依赖于自身索引，使用一维组织更直观且高效：

• 线程索引计算简单
• 内存访问模式连续，提高缓存利用率

对于二维数据(如图像、矩阵)，使用二维组织更自然，能够更好地映射数据结构和计算逻辑。

解题思路

__global__ void vector_add(const float* A, const float* B, float* C, int N) {int i = blockIdx.x * blockDim.x + threadIdx.x;if(i<N){C[i]=A[i]+B[i];}
}

根据第几个block和在block内的局部地址，得到总的id。注意i<N的边界约束，这是因为我们无法保证N是256的倍数，可能会产生数组越界。