CUDA

http://blog.sina.com.cn/s/blog_c1fccc770101do3q.html

CUDA是NVIDA公司提出的一种并行计算的结构,能够应用到科学计算的各个领域.笔者的电脑正好有一块支持CUDA的显卡,于是正好试用了一下.

 

1.CUDA的下载和配置:

关于支持CUDA的GPU型号，最好在NVIDA英文官网上确认，像笔者的Geforce530GT（OEM）就没有出现在中文官网支持CUDA型号的列表里。

 

笔者的环境是64bitWin7系统和64bit VS2008。下载可以在这个界面完成https://developer.nvidia.com/cuda-downloads. 下载的5.0 production release比较全，包括了原来的SDK等。正常安装即可，最好保证电脑安装了最新的clean的NVIDA显卡驱动。然后就可以在VS的新建project中找到CUDA项目了。CUDA安装包中包括了很多example，可以在C:\ProgramData\NVIDIA Corporation\CUDA Samples 中找到它们。

 

关于高亮显示可以把C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.0\doc\syntax_highlighting\visual_studio_7的usertype.dat文件拷贝到\Microsoft Visual Studio 9.0\Common7\IDE目录下。

也别忘了在Tools->Options->Text Editor->File Extension。在“Extension：”处添加(ADD).cu

 

2.CUDA的基本概念

可以在http://cudazone.nvidia.cn/cuda-education-training/ 找到很多资料。作者推荐Stanford大学的CS193G课程。相关ppt可以在上个网址的链接中找到。

笔者在这个只介绍一点最基本的概念.GPU运算的最基本单位是thread，thread可以构成一个运算的阵列block，block可以进一步构成grid。

 

函数定义:
__global__void my_kernel() { } 这种类型的函数是被CPU(Host)调用的
__device__float my_device_func() { } 这种类型的函数是被GPU(Device)调用的
变量定义:

__shared__float my_shared_array[32]; 这种数据类型可以被一个block共享，数据调用速度很快，相当于CPU中的L1 cache，GPU计算的一个问题就在于要把需要处理的线程分成合适的block，并且指定每一块block对应的显存，相当于人工分配缓存。
定义Block大小和程序运行:

dim3 grid_dim;

grid_dim.x = 64;

grid_dim.y = 64;//每个grid由64*64个block组成

dim3 block_dim; //类似的我们可以定义block的大小
my_kernel <<<grid_dim, block_dim >>>(...input...)//函数调用

 

进行计算时，我们首先要在显存上分配空间，来保存我们用以计算的数据，语句如下

cudaMalloc( (void**) &d_A, N * sizeof(float));

以下代码可以把内存中的数据拷入显存

cudaMemcpy( d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice) );

在计算结束后，我们可以类似地把数据拷回电脑内存。

 

我们可以发现函数调用只是由一句命令完成的，那么我们怎样使不同的thread完成不同的功能呢，以下的函数便是一个例子。

__global__void vector_add(float* A, float* B, float* C)
{
int i = threadIdx.x+ blockDim.x* blockIdx.x;
C[i] = A[i] + B[i];
}

其中threadIdx.x(y), blockIdx.x(y,z)唯一地标志了每个线程。

 

基于这些简单的概念，笔者写了一个简单的利用CUDA进行PDE求解的小程序。但事实是这个程序的运行效率并不是很高，笔者也懒得去优化block size和诸如此类的一些东西了，Geforce 530GT的性能比较捉急，大家可能最好采购价格高于400的独立显卡来玩计算。