GPU开发简史

GPU是现代PC机中的常见设备，采用了最先进的半导体制造工艺，能够实时生成逼真的3D图形。传统上，GPU的强大处理能力只被用于3D图像渲染，应用领域受到了限制，这无疑是对计算资源的极大浪费。随着GPU的可编程性的不断提高，利用GPU完成通用计算的研究渐渐活跃起来。将GPU用于图形渲染以外领域的计算称为GPGPU。GPGPU计算通常采用CPU+GPU异构模式，由CPU负责执行复杂逻辑处理和事务管理等不适合数据并行的计算，由GPU复杂计算密集型的大规模数据并行计算。但是传统的GPGPU受硬件可编程性和开发方式的制约，应用领域受到了限制，开发难度也很大。

2007年6月，NVIDIA推出了CUDA(ComputeUnified Device Architecture,统一计算设备架构)。CUDA不需要借助于图形学API，采用了比较容易掌握的类C语言进行开发。开发人员能够从熟悉的C语言比较平稳地从CPU过渡到GPU，而不必重新学习语法。当然，要开发高性能的GPU通用计算程序，开发人员仍然需要掌握并行算法和GPU架构方面的知识。

GPU硬件架构与CUDA编程模型

GPU硬件架构

GPU硬件结构主要由以下几个关键模块组成：内存（全局的，常量的，共享的）；流处理器簇（SM）；流处理器(SP)。如图所示：

GPU实际上是一个SM的阵列，每个SM包含N个核。一个GPU设备中包含一个或多个SM。SM内部组成结构图如下所示：

CUDA编程模型

主机(host）)设备(device) 内核函数(kernel函数)。CUDA编程模型将CPU称为主机（host）,将GPU称为设备（Device），将运行在GPU上的CUDA并行计算函数称为内核函数（kernel函数）。

在这个模型中，CPU与GPU协同工作，各司其职。CPU负责逻辑性强的事务处理和串行计算，GPU则专注于执行高度线性化的并行处理任务。一旦确定了程序中的并行部分，就可以考虑把这部分计算工作交给GPU（通过内核函数实现），如下图所示：

从上图可以看出，一个kernel函数中存在两个层次的并行，即Grid中的block间并行和block中的thread间并行。Kernel以线程网格的形式组织，每个线程网格由若干线程块组成，而每个线程块又由若干个线程组成。实质上，kernel函数是以block为单位执行的，CUDA引入grid只是用来表示一系列可以被并行执行的block的集合。

CUDA存储器模型

每一个线程拥有自己的私有存储器寄存器（Registers）和局部存储器(Local Memory)；每一个线程块拥有一块共享存储器；最后，grid中所有的线程都可以访问同一块全局存储器(Global Memory)。除此之外，还有两种可以被所有线程访问的只读存储器：常数存储器(Const Memory)和纹理存储器(Texture Memory)，它们分别为不同的应用进行了优化。

不同存储器的位置，缓存情况，访问权限和生存域如下图所示：

硬件映射

流处理器簇(SM)拥有完整前端的处理核心，包含取指，解码，分发逻辑和执行单元。CUDA中的kernel函数实质上是以block为单位执行的，同一block中的线程需要共享数据，因此他们必须在同一个流处理器簇中，而block中的每一个thread则被分配到一个SP上执行。一个block必须被分配到一个SM中，但是一个SM中同一时刻可以有多个活动block在等待执行，如图所示：

CUDA 软件体系

CUDA软件体系由CUDA Library,CUDA runtime API, CUDA driver API组成。CUDA的核心是CUDA C语言，它包含对C语言的最小扩展集和一个运行时库，使用这些扩展和运行时库的源文件必须通过nvcc编译器进行编译。CUDA C语言编译得到的只是GPU端的代码，而要管理GPU资源，在GPU上分配显存并启动内核函数，就必须借助CUDA运行时API（runtime API）或者CUDA驱动API（driver API）来实现。

CUDA编程基础

CUDA C语言

CUDA C语言为程序员提供了一种用C语言编写设备端代码的编程方式，包括对C的一些必要扩展和一个运行时库。CUDA对C的扩展主要包括以下几个方面：

1.      引入函数类型限定符。用来规定函数是在host还是在device上执行，以及这个函数是从host调用还是从device调用。这些限定符有：__device__,__host__和__global__。

2.      引入了变量类型限定符。用来规定变量存储在哪一类型的存储器上。包括：__device__，__shared__和__constant__.

3.      引入了内置矢量类型。如char4,ushort3,double2,dim3等。

4.      引入4个内建变量。blockIdx和threadIdx用于索引线程块和线程，gridDim和blockDim用于描述线程网格和线程块的维度。

5.      引入了<<<>>>运算符。用于指定线程网格和线程块的维度，传递执行参数。

6.      引入了一些函数：memory fence函数，同步函数，数学函数，纹理函数，测时函数，原子函数，warp vote函数。

常用runtime API

cudaGetDeviceCount(int*);             //获取可用设备数量

cudaSetDevice(int)                                                 //设置当前设备

cudaMalloc(void*,int)               //在显存分配存储空间

cudaMemcpy(void*,void*,int,  cudaMemcpyKind kind)        //拷贝数据

示例程序

const int num = 500000;

__global__ voidadd(int* A, int* B, int* C)

{

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

         C[i] = A[i] + B[i];

}

void main()

{

         int *a = new int[num];

         int *b = new int[num];

         int *c = new int[num];

         for (int i = 0; i < num; i++)

         {

                   a[i] = b[i] = i;

         }

         int device_count;

         cudaGetDeviceCount(&device_count);

         if (device_count < 1)

         {

                   cout << "没有支持CUDA的显卡\n";

                   system("pause");

                   return;

         }

         if (cudaSetDevice(0))

         {

                   cout << "设置显卡失败";

                   system("pause");

                   return;

         }

         int* Dev_a, *Dev_b, *Dev_c;

         if (cudaMalloc(&Dev_a,num*sizeof(int)) || cudaMalloc(&Dev_b, num * sizeof(int)) ||cudaMalloc(&Dev_c, num * sizeof(int)))

         {

                   cout << "显卡分配空间失败";

                   return;

         }

         if (cudaMemcpy(Dev_a, a, num *sizeof(int), cudaMemcpyHostToDevice) || cudaMemcpy(Dev_b, b, num * sizeof(int),cudaMemcpyHostToDevice))

         {

                   cout << "向显存拷贝数据失败";

                   return;

         }

         add << <500, 1000 >>>(Dev_a, Dev_b, Dev_c);

         cudaMemcpy(c, Dev_c, num * sizeof(int),cudaMemcpyDeviceToHost);

         for (int i = 0; i < 5; i++)

         {

                   cout << c[i] <<endl;

         }

         cudaFree(Dev_a);

         cudaFree(Dev_b);

         cudaFree(Dev_c);

         system("pause");

}