CUDA程序设计(一)

为什么需要GPU

几年前我启动并主导了一个项目，当时还在谷歌，这个项目叫谷歌大脑。该项目利用谷歌的计算基础设施来构建神经网络。

规模大概比之前的神经网络扩大了一百倍，我们的方法是用约一千台电脑。这确实使深度学习取得了相当大的进展。用到相当多的

计算机。不久之后我发现，之前我并没意识到，用一千台电脑是一项非常昂贵的技术。因此，我和我的朋友，意识到，利用一种

不同的技术，仅用三台电脑，而非一千台，就可以做到这点，而秘诀就是利用GPU技术。

---Andrew Ng [The Big Talk：深度学习与在线教育]

GPGPU是最近几年才流行起来的Project，并且在科学计算上使用频繁。

深度学习在06年立项，当时的数据集小的可怜，只需要CPU仿真一下。

09年，Stanford的CV实验室公布ImageNet后，深度学习开始拥抱大数据。

浅层结构+大数据尚且依赖云计算，深度结构+大数据的需求则是云计算提供不了的。

12年，Hinton组的Alex开源了cuda_convnet，正式大规模使用GPU。衍生出Theano、Caffe等项目。

在今天，如果你要修改这些项目代码，包括自己的创新改良，不会CUDA是不行的。

这也是深度学习带来的一个全新领域，它要求研究者不仅要理论强，建模强，程序设计能力也要过硬，不能纸上谈兵。

单从这点来看，深度学习依然和传统机器学习学派分道而行（多使用CPU多线程、云计算）。

大规模分布式CPU云应该尽量避免，尽管面子大，唬人效果好，但毕竟计算效率低，平台组建麻烦。

高效的程序设计，应当从单机中，榨取更多的计算力，进而考虑多机联合。

CUDA物理计算体系

1.1 CUDA特性

CUDA可以说是介于NVIDIA Driver和DirectX等库的中间产物。

传统游戏开发者，无须考虑太多GPU硬件架构，毕竟A卡、N卡水火不容。

通过设备商提供的Driver，DirectX或者OpenCL可以将不同架构的硬件底层封装，提供给开发者共有的API。

而CUDA，则可以看成是Driver、DX之流的上级。

为了泛型编程(C、C++、Fortran多语言）、以及榨取更多的计算力，NVIDIA对OpenCL进行的改装，

贴合自己的GPU硬件架构，量身定做出CUDA。

较游戏程序员不同，CUDA程序员主要工作，就是把握硬件架构，在算法理论时间复杂度下，

将算法串行执行体系，改组为并行执行体系。违背硬件的并行设计，只会让你的程序跑的百倍之慢。

1.2 NVIDIA GPU三大经典架构（Fermi、Kepler、Maxwell)

与海贼王里CP9道力值 (道力可以反映一个人的战斗力，但不能完全决定一个人的实力)设定相似：

NVIDIA为它在不同成长阶段卖出的产品，规定了计算能力体系：

计算能力1.0是跑CUDA的最低条件，这一时期的代表作是8800GT家族。

Fermi架构的计算能力是2.0, Kepler是3.0，Maxwell是4.0。

1.2.1 年轻有为的Fermi（费米）

Fermi架构，也就是GTX 4XX系卡，这算是CUDA的一个重要转折点——它把流处理器（SP）改名为"CUDA核心"。

NVIDIA的[Fermi架构图]：

#include "cuda_runtime.h"
#include "cstdio"
bool InitCUDA()
{
int cnt, i;
cudaGetDeviceCount(&cnt);
if (!cnt)
{
printf("No GPU Device\n");
return false;
}
for (i = 0; i < cnt; i++)
{
cudaDeviceProp device;
if (cudaGetDeviceProperties(&device, i)==cudaSuccess)
if (device.major >= 1) break;
}
if (i == cnt)
{
printf("GPU's Computing Capability must higher than 1.0\n");
return false;
}
cudaSetDevice(i);
return true;
}

InitCUDA()
只有计算能力>=1.0(8800GT以上)的N卡才能做CUDA计算，默认选择符合条件的第一块卡。

3.2.4 HelloWorld

GPU版HelloWorld，主要目的是演示CUDA基本程序框架：

☻将HelloWorld复制进显存

☻让GPU完成strcpy函数

☻将显存中的HelloWorld转回内存，并且打印

整体代码:

/****kernel.cu****/
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
__global__ void cudaStrcpy(char *des, char *src)  /*内核函数*/
{
while ((*src) != '\0') *des++ = *src++;
*des = '\0';
}

/****gpu_helloworld.cu****/
#include "device.cu"
#include "kernel.cu"
#include "cstring"
void helloworld(char *str1, char *str2)
{
InitCUDA();
char *dev_str1=0, *dev_str2=0;
int size = strlen(str1) + 1;
cudaMalloc((void**)&dev_str1, size);   /*cuda系函数必须放在cu文件里*/
cudaMalloc((void**)&dev_str2, size);
cudaMemcpy(dev_str1, str1, size,cudaMemcpyHostToDevice);
cudaStrcpy<<<1,1>>>(dev_str2, dev_str1); /*单线程块、单线程*/
cudaMemcpy(str2, dev_str1, size, cudaMemcpyDeviceToHost);
}

/****main.cpp****/
#include "cstdio"
#include "cstring"
extern void helloworld(char *str1, char *str2);
int main()
{
char src[] = "HelloWorld with CUDA";
char *des = new char[strlen(src)+1];
helloworld(src, des);
printf("%s\n", des);
}

3.3 GPU版向量加法

向量加法是CUDA 7.0在VS中提供的样例模板，演示了并行算法的经典trick：循环消除。

利用单个线程块中，多个线程并发执行，来消除循环。

时间复杂度估计，不能简单从O(n)迁移到O(1)，因为GPU同时并行量存在限制。

即便是Kepler架构中拥有192SP的SM阵列，理论同时并行量也不过是2048。

/****kernel.cu****/
__global__ void kernel_plus(int *a, int *b, int *c)
{
int x = threadIdx.x;
c[x] = a[x] + b[x];
}

/****gpu_vectoradd.cu****/
void vectorAdd(int *a, int *b, int *c,int size)
{
if (!InitCUDA()) return;
int *dev_a = 0, *dev_b = 0, *dev_c = 0;
cudaMalloc((void**)&dev_a, size*sizeof(int));
cudaMalloc((void**)&dev_b, size*sizeof(int));
cudaMalloc((void**)&dev_c, size*sizeof(int));
cudaMemcpy(dev_a, a, size*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, size*sizeof(int), cudaMemcpyHostToDevice);
kernel_plus << <1, size >> >(dev_a, dev_b, dev_c);
cudaMemcpy(c, dev_c, size*sizeof(int), cudaMemcpyDeviceToHost);
}

/****main.cpp****/
extern void vectorAdd(int *a, int *b, int *c, int size);
int main()
{
int a[5] = { 1, 2, 3, 4, 5 }, b[5] = { 10, 20, 30, 40, 50 }, c[5] = { 0 };
vectorAdd(a, b, c,5);
printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
c[0], c[1], c[2], c[3], c[4]);
}