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

参与评论 您还未登录,请先 登录 后发表或查看评论

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
©️2022 CSDN 皮肤主题:大白 设计师:CSDN官方博客 返回首页
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、C币套餐、付费专栏及课程。

余额充值