在现代计算机编程领域,利用 GPU 的并行计算能力可以显著提升特定任务的处理速度。本文要介绍的这段 C++ 代码,实现了在 GPU 上并行计算两个数组对应元素之和的功能。通过 CUDA 并行计算的方式,充分发挥了 GPU 大量核心同时处理数据的优势,相较于传统的 CPU 串行计算,能够更高效地完成数组元素相加的任务。下面我们将对代码进行详细拆解和分析,以便更好地理解其实现原理和工作流程。
1. 头文件包含和命名空间声明
#include <stdio.h>
#include <stdlib.h>
#include <float.h>
#include <vector>
#include <assert.h>
#include <algorithm>
#include <cublas_v2.h>
#include <cuda_runtime.h>
#include <iostream>
using namespace std;
描述:这段代码主要是包含了程序所需的各种头文件。stdio.h
和 stdlib.h
是标准输入输出和标准库相关的头文件;float.h
提供了浮点数相关的宏定义;vector
用于使用 C++ 的动态数组;assert.h
提供了断言功能;algorithm
包含了一些常用的算法;cublas_v2.h
是 CUDA 的 BLAS 库头文件,不过在这段代码中未实际使用;cuda_runtime.h
是 CUDA 运行时库的头文件,用于使用 CUDA 的各种功能;iostream
用于 C++ 的输入输出流。using namespace std;
则是使用标准命名空间,方便后续使用标准库中的类和函数。
2. 宏定义
#define GPU_BLOCK_THREADS 512
描述:定义了一个宏 GPU_BLOCK_THREADS
,其值为 512。这个宏表示每个 CUDA 线程块中的线程数量,后续代码中会多次使用这个值来设置线程块和网格的维度。
3. 设置 grid 中 block 数的函数
//设置grid中block数
static dim3 grid_dims(int numJobs) {
//GPU_BLOCK_THREADS =512;
//if (需要的总线程数numJobs < 512), 则,grid_dim = 1;
//else 则 grid_dim = numJobs /512 向上取整;
int numBlock = numJobs < GPU_BLOCK_THREADS ? numJobs : GPU_BLOCK_THREADS;
return (numBlock==numJobs) ? 1:std::ceil(numJobs/GPU_BLOCK_THREADS);
}
描述:这个函数 grid_dims
用于根据需要处理的任务数量 numJobs
来计算 CUDA 网格(grid)中需要的线程块(block)数量。首先判断 numJobs
是否小于 GPU_BLOCK_THREADS
,如果是,则 numBlock
等于 numJobs
,否则等于 GPU_BLOCK_THREADS
。然后根据 numBlock
和 numJobs
的关系返回网格维度,如果 numBlock
等于 numJobs
,则网格维度为 1,否则返回 numJobs
除以 GPU_BLOCK_THREADS
向上取整的结果。
4. 设置 block 中线程数的函数
//设置block中线程数
static dim3 block_dims(int numJobs) {
// block的大小:512 ,
//if (需要的总线程numJobs < 512);则,block = numJobs
//else block的大小:512
return numJobs < GPU_BLOCK_THREADS ? numJobs : GPU_BLOCK_THREADS;
}
描述:函数 block_dims
用于根据需要处理的任务数量 numJobs
来确定每个线程块中的线程数量。如果 numJobs
小于 GPU_BLOCK_THREADS
,则每个线程块中的线程数量为 numJobs
,否则为 GPU_BLOCK_THREADS
。
5. CUDA 核函数
// N :数组的长度
__global__ void gpu_add(float* a, float* b, float* sum, int N) {
int index = blockDim.x * blockIdx.x + threadIdx.x;
//超出数组长度的不计算
if (index > N) return;
sum[index] = a[index] + b[index];
}
描述:这是一个 CUDA 核函数 gpu_add