前言
在信号处理领域，随着信号处理规模急剧的增大，采用传统的处理机（如FPGA和DSP），其精度虽然可以达到工程要求，但是存在可移植性差、程序开发和调式周期长且硬件成本高等缺点。与传统处理机相比，仅使用中央处理器(CPU)作为计算平台，其底层架构决定其更适合逻辑密集运算，在处理实时数据时会受到很大的限制。而GPU的底层架构决定了其在并行计算方面具有独特的优势，适合大规模数据的并行运算。
采用CPU+GPU异构平台可以获得以下优势：
- 更高的性能：GPU芯片底层设计了大量算术逻辑单元，可以实现大规模数据的并行运算。
- 更强的通用性：CPU和GPU都是通用处理器，支持多种数据类型的运算，而传统的处理方式需要专门的配置及编程才能实现特定的算法功能，同时当数据类型复杂和算法功能多样时，会大大增加开发难度。
- 更高的灵活性：CPU和GPU都是可编程处理器，因此可以通过编程实现参数配置的优化，并且可以通过灵活设置模块参数以适用于不同的工作模式，具有很好的灵活性。
- 更高的计算精度：CPU和GPU都支持单精度和双精度运算，相同的输入信号经过处理后可以获得一致的输出，精度甚至可达10-13。
- 更低的成本：CPU和GPU都是通用处理器，其成本相较于FPGA和DSP低很多，而且由于CPU和GPU使用领域越来越广泛，因此具有更大的市场规模，其价格也相对更低。
1 GPU技术的发展：
1999年，发布标志性产品GeForce256时被首次提出，而这款产品也是全球首款图形处理器(Graphics Processing Unit，GPU)；
2001年，英伟达公司发布了Geforce3系列，不仅提高了运行速度，还首次支持微软公司同年发布DirectX 8.0的标准，同时是首款可编程GPU；
2001年，ATI公司推出了Rafeon 8500，与NVIDIA展开了强有力的竞争，大大推动了GPU技术的发展；
2004年，英伟达不再满足于利用GPU做渲染，而是开始转向通用计算领域，具有代表性的有NVIDIA的GeForce 9600；
2007年，英伟达的“Tesla”（计算能力1.X）是第一代应用于大规模并行计算的架构并实现了统一着色器模型，为后续其他架构的设计奠定了基础；NVIDIA同年正式发布了CUDA（计算统一架构），再一次引发GPU通用革命，它可以利用C/C++完成程序开发，并且采用了统一处理架构及共享存储器，推动了GPU更加广泛的应用
2009年，NVIDIA在首届GPU大会上推出了“Fermi”架构（计算能力2X），引入了单指令多线程(SIMT)执行模型；
2012年秋，再次推出Kepler架构（计算能力3.X），此架构的创新点主要在于每个SM单元的CUDA内核数由32个增加至192个，提高了6倍；
2017年，Volta架构（计算能力7.X）发布，此架构与前面所有系列最大的区别就是将CUDA的内核拆分为了两部分；
……
2 GPU硬件架构
GPU(Graphics Processing Unit)即图形处理单元，作为一种处理器辅助CPU完成大规模密集数据运算，因此GPU的设计目的不再像CPU为了使任务执行单元以最快速度获取数据和程序指令，而是侧重于提高整体的数据吞吐量，由于设计目标的不一致也就决定了两者在底层架构的差异，下图为CPU与GPU的硬件架构示意图的对比。
图1 CPU与GPU的硬件架构示意图的对比
根据上图可以看出，CPU为了实现设计目标设计了相当多的存储单元(Cache)和控制单元(Control)，因此适合执行逻辑比较复杂的串行任务，相比之下负责运算的算术逻辑单元(Arithmetic and Logic Unit, ALU)只占据了一小部分，这也是为何使用CPU进行大规模并行数据运算效率低的原因之一。但是GPU芯片空间有大量的算术逻辑单元以支持并行计算，因此GPU适合大规模密集数据的并行运算。
图2为GPU计算单元简图，可以看出，每个SM包含了8个流处理器(Streaming Processor, SP)及特殊函数单元(Special Function Unit, SFU)、寄存器、共享内存和常量缓存。在执行高密度计算时，每一个SM都可以实现数以千计个线程的并发执行，当核函数被启动时，GPU会根据分配策略将线程块分配到相应的SM执行，但是不同的架构下单个SM的常驻线程数也不尽相同。
图2 GPU计算单元简图
3 CUDA编程模型
CUDA编程模型是CPU+GPU协同的异构模型，提供了所必须的编程接口。在CUDA中，有两个不可或缺的概念：host和device，其中host表示在CPU端运行，称为主机端；device表示在GPU端运行，称为设备端。
在设备端运行的程序称为核函数(kernal)，执行同一个核函数所需的所有线程都分布在同一个网格(grid)内，此网格内的所有线程的全局内存可以实现共享。每一个网格又由许多个线程块(block)构成，每一个线程块又由许多线程构成，线程是程序执行的最小单位。
CUDA的线程层次结构如图3所示，从图中可以看出，CUDA有两层可以实现并行逻辑计算，分别是block层和thread层，其中block和SM构成一一映射，thread和SP(core)构成一一映射。
图3 线程层次结构
4 CUDA编程基本步骤
CUDA编程有一套比较完备的步骤，如图4所示，主要包括五部分：
1. 为设备分配显存空间；
2. 将内存中的数据传输至显存并保存(H2D)；
3. 执行kernal函数；
4. 将生成的结果传输至内存(D2H)；
5. 释放显存空间。
图4 CUDA编程基本步骤
下面以向量（矩阵）的对应元素相加的GPU实现为例对CUDA编程的基本步骤做简要说明：
#include"stdio.h"
#include"cuda.h"
//实现两个矩阵对应元素相加功能
__global__ void addFun(float* a, float* b, float* c, int L) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < L) {
c[idx] = a[idx] + b[idx];
}
}
int main(int argc, char **argv) {
int length = 1024 ;
//主机端(CPU)变量分配内存及初始化
float *a_H = (float*) malloc(sizeof(float) * length);
float *b_H = (float*) malloc(sizeof(float) * length);
float *c_H = (float*) malloc(sizeof(float) * length);
memset(a_H, 0, sizeof(float) * length);
memset(b_H, 0, sizeof(float) * length);
memset(c_H, 0, sizeof(float) * length);
//矩阵赋初值
for (int i = 0; i < length; i++) {
a_H[i] = i+1;
b_H[i] = i+2;
}
for (int i = 0; i < 5; i++) {
printf("输入数据a_H[%d]=%f\n",i, a_H[i]);
printf("输入数据b_H[%d]=%f\n",i, b_H[i]);
}
//设备端(GPU)变量分配内存及初始化
float *a_D;
cudaMalloc((void**) &a_D, length * sizeof(float));
float *b_D;
cudaMalloc((void**) &b_D, length * sizeof(float));
float *c_D;
cudaMalloc((void**) &c_D, length * sizeof(float));
cudaMemset(a_D, 0, length * sizeof(float));
cudaMemset(b_D, 0, length * sizeof(float));
cudaMemset(c_D, 0, length * sizeof(float));
//将数据从主机端复制至设备端
cudaMemcpy(a_D, a_H, sizeof(float) * length, cudaMemcpyHostToDevice);
cudaMemcpy(b_D, b_H, sizeof(float) * length, cudaMemcpyHostToDevice);
//核函数调用
addFun<<<(length + 511) / 512,512>>>(a_D, b_D, c_D, length);
//将数据从设备端复制至主机端
cudaMemcpy(c_H, c_D, sizeof(float) * length, cudaMemcpyDeviceToHost);
for (int i = 0; i < 5; i++) {
printf("经GPU计算后的结果c_H=[%d]%f\n",i, c_H[i]);
}
}如图5为上述代码的执行结果，可以看出，可以通过GPU并行实现矩阵对应元素的相加功能，这也是最简单、最基本的CUDA编程流程。
图5 矩阵元素相加GPU运行结果
扫一扫
2768
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
