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計算的大概印象。