CUDA: 数组求和

在高性能计算领域，GPU因为其架构的原因，在并行计算领域正发挥越来越多的用途，比如进行大量计算的游戏、绘图、图像算法等方面，采用GPU进行加速可以得到显著的性能提高。如今，Nvidia显卡在pc上的普及，cuda正是nvidia推出的通用并行计算架构。

下面在学习《深入浅出CUDA》的基础上初次体验下CUDA。

1.工程设置

这个就不多说了，新建一个空的Win32控制台应用程序，设置好工程属性(见前篇博文)。

2.程序初始化

首先加好头文件

#include <stdio.h>  //C标准输入输出接口#include <stdlib.h>#include <cuda_runtime.h>  //使用runtime API

定义CUDA初始化函数InitCUDA(),获得CUDA设备返回true，未获得返回false

//CUDA初始化bool InitCUDA(){    int count;     //传回有计算能力的设备数(≥1)，没有回传回1，device 0是一个仿真装置，不支持CUDA功能    cudaGetDeviceCount(&count);    if(count == 0) //没有cuda计算能力的设备    {        fprintf(stderr,"There is no device.\n");        return false;    }    int i;    for(i=0;i<count;i++)    {        cudaDeviceProp prop; //设备属性        if (cudaGetDeviceProperties(&prop,i)==cudaSuccess) //取得设备数据，brief Returns information about the compute-device        {            if (prop.major>=1) //cuda计算能力            {                break;            }        }    }    if (i==count)    {        fprintf(stderr,"There is no device supporting CUDA 1.x\n");        return false;    }    cudaSetDevice(i); //brief Set device to be used for GPU executions    return true;}

当然，CUDA程序的入口函数也是main()。

int main(){    if (!InitCUDA())    {        return 0;    }    printf("CUDA initialized.\n");}

这样，一个简单的可运行初始化程序就完成了，下面用到GPU去计算一个数组的和，为了体现gpu的并行计算能力，将数组长度设大一点

3.产生数组

定义一个数组，并随即产生数组中各个元素的值

#define DATA_SIZE 1048576 //定义数据长度int data[DATA_SIZE];//产生数组元素值void GenerateNumbers(int *number,int size){    for (int i=0;i<size;i++)    {        number[i]=rand()%10; //产生0~9的随机数    }}

4.显卡上运行的程序

在显卡设备上执行的程序，程序的写法与一般C是一样的，但是gpu上运行的程序目前还是有一些编程规范要求，具体查看相关文档。__global__为函数类型限定符，表示函数为内核函数，在设备(gpu)上执行，只能从主机(cpu)中调用。

要注意的是显卡设备上计算所需的变量都为指针类型，具体的CUDA接口函数可以通过源码查看含义以及各参数意义

需要注意的for循环内没有采用for(i=0;i<DATA_SIZE/(BLOCK_NUM*THREAD_NUM);i++),是因为一个块内的线程是共享内存的，当一个块内线程读取数组数据时，是thread0->thread1->...->thread255顺序读取的，要保持个线程间读取数组数据的连续性，这样能提高性能。数据分配在以后的并行计算中是很重要的。

//在显示芯片上执行的函式// 注意：global关键字左右各是两个___global__ static void sumOfSquares(int *num,int *result,clock_t* time){    const int tid = threadIdx.x; //取得线程号    const int bid = blockIdx.x; //获得块号    int sum=0;    int i;    if(tid==0)        time[bid]=clock(); //开始时间    for(i=bid*THREAD_NUM+tid;i<DATA_SIZE;i+=BLOCK_NUM*THREAD_NUM) //注意循环内变化    {        sum += num[i]*num[i];    }    result[bid*THREAD_NUM+tid] = sum; //第bid个块内第tid个线程计算的结果    if(tid==0)        time[bid+BLOCK_NUM] = clock(); //运行时间}

5.分配显存，执行并行

数组中的元素要放到GPU中去运算，因此必须先将内存中的数组拷贝至显卡内存中，这样在显卡中才能读取显存中的数据进行运算。在显卡中，我们利用32个块，每个块开辟256个线程进行并行计算。

分配了显存后就可以调用核函数进行并行计算了，计算完后要将计算的结果从显存中拷贝至内存，然后释放掉分配的显存空间。要注意拷贝时数据长度要保持一致。

每个块将计算出来的结果传到内存上，再在cpu上计算总和。

cpu、gpu、内存间联系如下，gpu中的块和快中的thread只是给出了示例，不是实际数量

cudaMalloc表示在显存上分配内存空间

#define BLOCK_NUM 32 //块数量#define THREAD_NUM 256 //每个块中线程数int main(){    //...CUDA初始化    GenerateNumbers(data,DATA_SIZE); //产生随机数    int *gpudata,*result; //gpu设备上的数    clock_t* time; //计算时间(以GPU时钟为基准)    cudaMalloc((void**)&gpudata,sizeof(int)*DATA_SIZE); //分配显存，Allocate memory on the device    cudaMalloc((void**)&result,sizeof(int)*BLOCK_NUM*THREAD_NUM);    cudaMalloc((void**)&time,sizeof(clock_t)*BLOCK_NUM*2);    // Copies data between host and device    cudaMemcpy(gpudata,data,sizeof(int)*DATA_SIZE,cudaMemcpyHostToDevice); //Host->Device    // 函式名称<<<block 数目，thread 数目,shared memory 大小>>>(参数...)    sumOfSquares<<<BLOCK_NUM,THREAD_NUM,0>>>(gpudata,result,time);    int sum[THREAD_NUM*BLOCK_NUM];    clock_t time_used[BLOCK_NUM*2]; //运行时间    cudaMemcpy(&sum,result,sizeof(int)*THREAD_NUM*BLOCK_NUM,cudaMemcpyDeviceToHost);    cudaMemcpy(&time_used,time,sizeof(clock_t)*BLOCK_NUM*2,cudaMemcpyDeviceToHost);    cudaFree(gpudata); //释放显存    cudaFree(result);    cudaFree(time);    //计算每个线程计算出来和的总和    int final_sum=0;    for(int i=0;i<THREAD_NUM*BLOCK_NUM;i++)    {        final_sum += sum[i];    }    clock_t min_start,max_end;    min_start=time_used[0];    max_end=time_used[BLOCK_NUM];    //计算GPU上总运行时间    for (int i=0;i<BLOCK_NUM;i++)    {        if(min_start>time_used[i])            min_start=time_used[i];        if(max_end<time_used[i+BLOCK_NUM])            max_end=time_used[i+BLOCK_NUM];    }    printf("sum:%d  time:%d\n",final_sum,max_end-min_start); }

6.CPU上验证

下面写个验证程序，看gpu上执行的正确性

//在上面的代码后//在cpu上计算验证    final_sum=0;    for(int i=0;i<DATA_SIZE;i++)    {        final_sum += data[i]*data[i];    }    printf("(CPU) sum:%d\n",final_sum);

7.编译运行

结果如下：

这里的time后的时间是在gpu上运行的时间，单位是gpu周期，即花去了1099014个gpu周期，本机的gpu主频为800MHz，所以花的时间就是10099014/(800*1000) ms

后记：细心的读者会发现，本文并没有体现出使用gpu计算效率有多高，确实，本文只是利用cuda进行并行计算的一个初次体验，意在给出使用gpu计算的大概印象。

http://www.cnblogs.com/Romi/archive/2012/04/20/2459817.html