CUDA：在GPU上实现向量加法的两种方法

pqin

已于 2023-12-25 12:06:55 修改

阅读量642

点赞数 9

文章标签：算法 c++ c语言开发语言

于 2023-12-25 12:05:17 首次发布

本文链接：https://blog.csdn.net/weixin_43308899/article/details/135195521

版权

写在前面：本文主要向量加法的多种实现，包括在CPU上的一种实现和在GPU上的两种实现。

一. 向量加法在CPU上的实现

void vec_add_cpu(FLOAT *x, FLOAT *y, FLOAT *z, int N)
{
    for (int i = 0; i < N; i++) z[i] = y[i] + x[i];
}

在CPU上的代码其实是比较简单的，输入是x和y两个向量，输出是向量z，每个向量包含的元素个数是N。在函数内部，通过for循环依次对向量的每个元素进行加法操作。

二. 向量加法在GPU上的实现（实现一）

__global__ void vec_add_one(float *x, float *y, float *z, int N)
{
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    if (idx < N) z[idx] = y[idx] + x[idx];
}

优点:
- 简单直接: 这个函数更简洁，易于理解和维护。
- 灵活性: 对数据大小和对齐的要求不像使用 float4 那么严格。
- 更广泛的适用性: 因为它简单，所以更容易适应不同的需求和环境。
缺点:
- 单一元素处理: 每个线程只处理一个元素，可能在大数据集上效率不如循环处理。
- 可能的内存带宽限制: 如果内存访问是瓶颈，使用单个 float 而非 float4 可能导致性能受限。

三. 向量加法在GPU上的实现（实现二）

__global__ void vec_add_two (float* A,  float* B, float* C){
	// block and thread index
	int idx = blockIdx.x * blockDim.x + threadIdx.x;

	for(int i = idx; i < MEMORY_OFFSET / 4; i += blockDim.x * gridDim.x) {
		float4 a1 = reinterpret_cast<float4*>(A)[i];
		float4 b1 = reinterpret_cast<float4*>(B)[i];
		float4 c1;

		c1.x = a1.x + b1.x;
		c1.y = a1.y + b1.y;
		c1.z = a1.z + b1.z;
		c1.w = a1.w + b1.w;

		reinterpret_cast<float4*>(C)[i] = c1;
	}
}

优点:
- 使用 float4 类型: 这种方式可能提高内存吞吐量，因为每次访问可以加载和存储更多数据。float4 是一个包含四个浮点数的向量类型，可能允许更有效的内存访问和计算。
- 循环并行处理: 函数内的循环使得每个线程可以处理多个元素，这在大数据集上可能更有效率。
缺点:
- 复杂性较高: 代码相对复杂，难以理解和维护。
- 对齐和大小要求: 使用 float4 需要确保数据适当对齐，且数据大小必须是 4 的倍数。
- 可能的性能瓶颈: 如果 MEMORY_OFFSET / 4 不是 blockDim.x * gridDim.x 的整数倍，可能会导致线程执行不均衡，降低效率。

小知识

float4数据类型

float4 是一种在许多编程环境中使用的数据类型，特别是在进行 GPU 编程（如 CUDA 或 OpenCL）时。它是一个简单的数据结构，包含四个浮点数（通常是 32 位的 float 类型）。这种类型在处理图形、科学计算和其他需要大量数值运算的应用中非常有用。以下是一些 float4 类型的关键特点和用途：

组成：float4 由四个浮点数组成，通常表示为 x, y, z, w。这可以方便地用于表示如3D空间中的点或向量（使用 x, y, z），以及其他需要四个数值的数据（例如颜色的红、绿、蓝和 alpha 值）。
内存对齐：float4 类型自然地对齐为 128 位，这对于 GPU 和 SIMD（单指令多数据）编程非常重要，因为它允许更高效的内存访问和数据吞吐量。
在 GPU 编程中的应用：在 CUDA 和 OpenCL 等 GPU 编程框架中，float4 类型常被用于加速计算。例如，在处理图像或进行物理模拟时，可以利用 float4 同时处理多个数据点，从而提高性能。
操作：大多数情况下，你可以对 float4 类型进行标准的算术运算，就像处理单个 float 一样。许多编程环境提供了对这种类型的内置支持，包括加法、减法、乘法和除法等操作。
效率提升：在适当的场合使用 float4 可以显著提升程序的效率，特别是在需要大量并行运算的应用中。由于可以一次性处理四个浮点数，因此在处理向量化数据时可以减少必要的运算周期和内存访问次数。