在高性能计算领域，GPU因为其架构的原因，在并行计算领域正发挥越来越多的用途，比如进行大量计算的游戏、绘图、图像算法等方面，采用GPU进行加速可以得到显著的性能提高。如今，Nvidia显卡在pc上的普及，cuda正是nvidia推出的通用并行计算架构。
下面在学习《深入浅出CUDA》的基础上初次体验下CUDA。
1.工程设置
这个就不多说了，新建一个空的Win32控制台应用程序，设置好工程属性(见前篇博文)。
2.程序初始化
首先加好头文件
1 #include <stdio.h> //C标准输入输出接口 2 #include <stdlib.h> 3 #include <cuda_runtime.h> //使用runtime API
定义CUDA初始化函数InitCUDA(),获得CUDA设备返回true，未获得返回false
1 //CUDA初始化 2 bool InitCUDA() 3 { 4 int count; 5 //传回有计算能力的设备数(≥1)，没有回传回1，device 0是一个仿真装置，不支持CUDA功能 6 cudaGetDeviceCount(&count); 7 8 if(count == 0) //没有cuda计算能力的设备 9 { 10 fprintf(stderr,"There is no device.\n"); 11 return false; 12 } 13 14 int i; 15 for(i=0;i<count;i++) 16 { 17 cudaDeviceProp prop; //设备属性 18 if (cudaGetDeviceProperties(&prop,i)==cudaSuccess) //取得设备数据，brief Returns information about the compute-device 19 { 20 if (prop.major>=1) //cuda计算能力 21 { 22 break; 23 } 24 } 25 } 26 27 if (i==count) 28 { 29 fprintf(stderr,"There is no device supporting CUDA 1.x\n"); 30 return false; 31 } 32 33 cudaSetDevice(i); //brief Set device to be used for GPU executions 34 return true; 35 }
当然，CUDA程序的入口函数也是main()。
1 int main() 2 { 3 if (!InitCUDA()) 4 { 5 return 0; 6 } 7 printf("CUDA initialized.\n"); 8 }
这样，一个简单的可运行初始化程序就完成了，下面用到GPU去计算一个数组的和，为了体现gpu的并行计算能力，将数组长度设大一点
3.产生数组
定义一个数组，并随即产生数组中各个元素的值
1 #define DATA_SIZE 1048576 //定义数据长度 2 3 int data[DATA_SIZE]; 4 5 //产生数组元素值 6 void GenerateNumbers(int *number,int size) 7 { 8 for (int i=0;i<size;i++) 9 { 10 number[i]=rand()%10; //产生0~9的随机数 11 } 12 }
4.显卡上运行的程序
在显卡设备上执行的程序，程序的写法与一般C是一样的，但是gpu上运行的程序目前还是有一些编程规范要求，具体查看相关文档。__global__为函数类型限定符，表示函数为内核函数，在设备(gpu)上执行，只能从主机(cpu)中调用。
要注意的是显卡设备上计算所需的变量都为指针类型，具体的CUDA接口函数可以通过源码查看含义以及各参数意义
需要注意的for循环内没有采用for(i=0;i<DATA_SIZE/(BLOCK_NUM*THREAD_NUM);i++),是因为一个块内的线程是共享内存的，当一个块内线程读取数组数据时，是thread0->thread1->...->thread255顺序读取的，要保持个线程间读取数组数据的连续性，这样能提高性能。数据分配在以后的并行计算中是很重要的。
1 //在显示芯片上执行的函式 2 // 注意：global关键字左右各是两个_ 3 __global__ static void sumOfSquares(int *num,int *result,clock_t* time) 4 { 5 const int tid = threadIdx.x; //取得线程号 6 const int bid = blockIdx.x; //获得块号 7 int sum=0; 8 int i; 9 if(tid==0) 10 time[bid]=clock(); //开始时间 11 for(i=bid*THREAD_NUM+tid;i<DATA_SIZE;i+=BLOCK_NUM*THREAD_NUM) //注意循环内变化 12 { 13 sum += num[i]*num[i]; 14 } 15 result[bid*THREAD_NUM+tid] = sum; //第bid个块内第tid个线程计算的结果 16 if(tid==0) 17 time[bid+BLOCK_NUM] = clock(); //运行时间 18 }
5.分配显存，执行并行
数组中的元素要放到GPU中去运算，因此必须先将内存中的数组拷贝至显卡内存中，这样在显卡中才能读取显存中的数据进行运算。在显卡中，我们利用32个块，每个块开辟256个线程进行并行计算。
分配了显存后就可以调用核函数进行并行计算了，计算完后要将计算的结果从显存中拷贝至内存，然后释放掉分配的显存空间。要注意拷贝时数据长度要保持一致。
每个块将计算出来的结果传到内存上，再在cpu上计算总和。
cpu、gpu、内存间联系如下，gpu中的块和快中的thread只是给出了示例，不是实际数量
cudaMalloc表示在显存上分配内存空间
1 #define BLOCK_NUM 32 //块数量 2 #define THREAD_NUM 256 //每个块中线程数 3 4 int main() 5 { 6 //...CUDA初始化 7 GenerateNumbers(data,DATA_SIZE); //产生随机数 8 9 int *gpudata,*result; //gpu设备上的数 10 clock_t* time; //计算时间(以GPU时钟为基准) 11 12 cudaMalloc((void**)&gpudata,sizeof(int)*DATA_SIZE); //分配显存，Allocate memory on the device 13 cudaMalloc((void**)&result,sizeof(int)*BLOCK_NUM*THREAD_NUM); 14 cudaMalloc((void**)&time,sizeof(clock_t)*BLOCK_NUM*2); 15 // Copies data between host and device 16 cudaMemcpy(gpudata,data,sizeof(int)*DATA_SIZE,cudaMemcpyHostToDevice); //Host->Device
17 // 函式名称<<<block 数目，thread 数目,shared memory 大小>>>(参数...) 18 sumOfSquares<<<BLOCK_NUM,THREAD_NUM,0>>>(gpudata,result,time); 19 20 int sum[THREAD_NUM*BLOCK_NUM]; 21 clock_t time_used[BLOCK_NUM*2]; //运行时间 22 cudaMemcpy(&sum,result,sizeof(int)*THREAD_NUM*BLOCK_NUM,cudaMemcpyDeviceToHost); 23 cudaMemcpy(&time_used,time,sizeof(clock_t)*BLOCK_NUM*2,cudaMemcpyDeviceToHost); 24 cudaFree(gpudata); //释放显存 25 cudaFree(result); 26 cudaFree(time); 27 28 //计算每个线程计算出来和的总和 29 int final_sum=0; 30 for(int i=0;i<THREAD_NUM*BLOCK_NUM;i++) 31 { 32 final_sum += sum[i]; 33 } 34 clock_t min_start,max_end; 35 min_start=time_used[0]; 36 max_end=time_used[BLOCK_NUM]; 37 //计算GPU上总运行时间 38 for (int i=0;i<BLOCK_NUM;i++) 39 { 40 if(min_start>time_used[i]) 41 min_start=time_used[i]; 42 if(max_end<time_used[i+BLOCK_NUM]) 43 max_end=time_used[i+BLOCK_NUM]; 44 } 45 printf("sum:%d time:%d\n",final_sum,max_end-min_start);
6.CPU上验证
下面写个验证程序，看gpu上执行的正确性
1 //在上面的代码后 2 //在cpu上计算验证 3 final_sum=0; 4 for(int i=0;i<DATA_SIZE;i++) 5 { 6 final_sum += data[i]*data[i]; 7 } 8 printf("(CPU) sum:%d\n",final_sum);
7.编译运行
结果如下：
这里的time后的时间是在gpu上运行的时间，单位是gpu周期，即花去了1099014个gpu周期，本机的gpu主频为800MHz，所以花的时间就是10099014/(800*1000) ms
后记：细心的读者会发现，本文并没有体现出使用gpu计算效率有多高，确实，本文只是利用cuda进行并行计算的一个初次体验，意在给出使用gpu计算的大概印象。