CUDA编程入门

CUDA的全称是Compute Unified Device Architecture， 是NVIDIA® 公司的并行计算架构，主要是要来利用GPU的计算能力，来提高计算性能。进一步的说是在GPU上提供标准C编程语言，为在支持CUDA的NVIDIA GPU上进行并行计算而提供了统一的软硬件解决方案。

为什么要用GPU来计算呢，CPU不如GPU吗，这就要从CPU和GPU的架构说起了。

图片中绿色的是ALU（运算器），可以看到，CPU只有4个ALU，而GPU中的ALU多很多，并且GPU中把更多的晶体管用于数据处理方面。之所以这样是因为CPU和GPU的功能不一样，CPU要考虑上下文切换等所以需要复杂的控制单元和缓存来提高执行效率，而GPU不需要那样复杂的控制逻辑，所以GPU可以通过增加并行处理单元和存储控制单元来提供处理能力和存储器带宽。因而同期的GPU的计算能力会比CPU高很多个档次。为了充分利用GPU的运算能力，就有了现在的CUDA编程。

但是利用GPU运算有几个缺点。

• 显示芯片的运算单元数量很多，因此对于不能高度并行化的工作，所能带来的帮助就不大。
• 显示芯片目前通常只支持 32 bits 浮点数，且多半不能完全支持 IEEE 754 规格， 有些运算的精确度可能较低。目前许多显示芯片并没有分开的整数运算单元，因此整数运算的效率较差。
• 显示芯片通常不具有分支预测等复杂的流程控制单元，因此对于具有高度分支的程序，效率会比较差。

所以如果你的计算内容不能高度并行化的话那么就不要去考虑用GPU来加快你的计算了。

————–简介完毕，下面进行配置开发环节—————-

那么我们来搞一下CUDA开发环境的搭建。首先我是在windows下搞的，本来想去linux下搞，但是CUDA对gcc的版本有要求，不能太高，比较囧。

• 下载安装visual studio 2008，网上都有说明告诉你如何把试用版变为正式版，不要用2005，不然可能会出现奇怪的错误。
• 去CUDA的官网去下载安装Toolkit，Drivers ，SDK。
• 使用开勇的CUDA_VS_Wizard配置Visual Studio 2008的CUDA项目
• 安装Visual AssistantX
• 打开VS, 选择 工具->选项->项目与解决方案->VC++项目设置，在“C/C++文件扩展名”后添加*.cu，在“包括的扩展名”后添加.cu
• 打开VS, 选择 工具->选项->项目与解决方案->VC++目录，分别在包含文件(include)，库文件(lib)，源文件(src)里添加toolkit和sdk路径所对应的目录
• 导入注册表，让Visual AssistantX支持CUDA的cu文件和语法高亮

基本配置完成，之后我们在vs中选择新建项目，会在visual c++中看到cuda一项，选CUDAWinApp来新建一个CUDA的程序。建好之后，我们会看到一个sample.cu，代码如下

/********************************************************************
*  sample.cu
*  This is a example of the CUDA program.
*********************************************************************/
#include 
#include 
#include 
//#include 
/************************************************************************/
/* Init CUDA                                                            */
/************************************************************************/
bool InitCUDA(void){
	int count = 0;
	int i = 0;
	cudaGetDeviceCount(&count);
	if(count == 0) {
		fprintf(stderr, "There is no device.\n");
		return false;
	}
	for(i = 0; i < count; i++) {
		cudaDeviceProp prop;
		if(cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
			if(prop.major >= 1) {
				break;
			}
		}
	}
	if(i == count) {
		fprintf(stderr, "There is no device supporting CUDA.\n");
		return false;
	}
	cudaSetDevice(i);
	printf("CUDA initialized.\n");
	return true;
}
/************************************************************************/
/* Example                                                              */
/************************************************************************/
__global__ static void HelloCUDA(char* result, int num){
	int i = threadIdx.x;
	char p_HelloCUDA[] = "Hello CUDA!";
	if(i < num) result[i] = p_HelloCUDA[i];
}
/************************************************************************/
/* HelloCUDA                                                            */
/************************************************************************/
int main(int argc, char* argv[]){
	if(!InitCUDA()) {
		return 0;
	}
	char	*device_result	= 0;
	char	host_result[12]	={0};
	cudaMalloc((void**) &device_result, sizeof(char) * 11);
	HelloCUDA<<<1, 11, 0>>>(device_result, 11);
	cudaMemcpy(host_result, device_result, sizeof(char) * 11, cudaMemcpyDeviceToHost);
	printf("%s\n", host_result);
	cudaFree(device_result);
	return 0;
}

ctrl+F5运行，我们就能看到Hello CUDA!字样了。

让我们简单看一下这个程序，在CUDA编程中用__global__来表示这是个可以有CPU调用的运行在GPU上的一个函数。由于GPU访问数据都是在显存中获取的，所以我们需要调用cudaMalloc在显存中开辟空间，用法类似于在内存中开辟空间的malloc。之后调用函数HelloCUDA<<<1, 11, 0>>>(device_result, 11);尖括号中有三项，分别表示block的数量，thread的数量，和共享内存的大小。用threadIdx.x我们可以获取thread的编号，来达到并行的效果。最后我们需要调用cudaMemcpy把数据从显存中取回到内存中，并释放显存中的数据，这样cpu就可以获得最后的结果了。

前面提到了block还有thread，我们据此继续讲解一下CUDA的编程架构。

用__global__修饰的函数被称作一个kernel函数或者是叫做一个kernel，一个kernel会有一个grid（目前是，据说以后可能会是多个），一个grid包含多个block，一个block包含多个thread，thread就类似于平常cpu运算中的线程thread。这样的排布是因为我们需要不同粒度的并行，处于同一个block的thread可以通过共享内存来交换数据，而不同block的thread不能直接共享数据。为了能更方便的运算一二三维的数组，grid中的block或者是block中的thread都可以按照不同的维度来排布，比如如果block中的thread按照二维排布的话，我们可以通过threadIdx.x和threadIdx.y来获得当前thread的位置，这样的结构使得运算二维或三维数组更加自然，而且可以避免除法和取余操作，加快了速度。

那基本介绍就这样了，如果还想深入了解的话还是需要去看一下相关的书籍和文档才可以。

参考资料：

• 《深入浅出CUDA》
• 《CUDA 编程指南4.0》
• 《GPU高性能运算之CUDA》