cuda入门：runtime API创建CUDA程序

CUDA 目前有两种不同的 API：Runtime API 和 Driver API，两种 API 各有其适用的范围。由于 runtime API 较容易使用，一开始我们会以 runetime API 为主。

　　CUDA 的初始化

　　首先，先建立一个档案 first_cuda.cu。要使用 runtime API 的时候，需要 include cuda_runtime.h。所以，在程序的最前面，加上

#include <stdio.h>
#include <cuda_runtime.h>

　　接下来是一个 InitCUDA 函式，会呼叫 runtime API 中，有关初始化 CUDA 的功能：

bool InitCUDA()
{
     int count;

    cudaGetDeviceCount(&count);
     if(count ==  0) {
        fprintf(stderr,  " There is no device.\n ");
         return  false;
    }

     int i;
     for(i =  0; i < count; i++) {
        cudaDeviceProp prop;
         if(cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
             if(prop.major >=  1) {
               break;
            }
        }
    }

     if(i == count) {
        fprintf(stderr,  " There is no device supporting CUDA 1.x.\n ");
         return  false;
    }

    cudaSetDevice(i);

     return  true;
}

　　这个函式会先呼叫 cudaGetDeviceCount 函式，取得支持 CUDA 的装置的数目。如果系统上没有支持 CUDA 的装置，则它会传回 1，而 device 0 会是一个仿真的装置，但不支持 CUDA 1.0 以上的功能。所以，要确定系统上是否有支持 CUDA 的装置，需要对每个 device 呼叫 cudaGetDeviceProperties 函式，取得装置的各项数据，并判断装置支持的 CUDA 版本(prop.major 和 prop.minor 分别代表装置支持的版本号码，例如 1.0 则 prop.major 为 1 而 prop.minor 为 0)。

　　透过 cudaGetDeviceProperties 函式可以取得许多数据，除了装置支持的 CUDA 版本之外，还有装置的名称、内存的大小、最大的 thread 数目、执行单元的频率等等。详情可参考 NVIDIA 的 CUDA Programming Guide。

　　在找到支持 CUDA 1.0 以上的装置之后，就可以呼叫 cudaSetDevice 函式，把它设为目前要使用的装置。

　　最后是 main 函式。在 main 函式中我们直接呼叫刚才的 InitCUDA 函式，并显示适当的讯息：

int main()
{
     if(!InitCUDA()) {
         return  0;
    }

    printf( " CUDA initialized.\n ");

     return  0;
}

　　这样就可以利用 nvcc 来 compile 这个程序了。使用 Visual Studio 的话，若按照先前的设定方式，可以直接 Build Project 并执行。

　　nvcc 是 CUDA 的 compile 工具，它会将 .cu 檔拆解出在 GPU 上执行的部份，及在 host 上执行的部份，并呼叫适当的程序进行 compile 动作。在 GPU 执行的部份会透过 NVIDIA 提供的 compiler 编译成中介码，而 host 执行的部份则会透过系统上的 C++ compiler 编译(在 Windows 上使用 Visual C++ 而在 Linux 上使用 gcc)。

　　编译后的程序，执行时如果系统上有支持 CUDA 的装置，应该会显示 CUDA initialized. 的讯息，否则会显示相关的错误讯息。

利用 CUDA 进行运算

　　到目前为止，我们的程序并没有做什么有用的工作。所以，现在我们加入一个简单的动作，就是把一大堆数字，计算出它的平方和。

　　首先，把程序最前面的 include 部份改成：

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>

#define DATA_SIZE 1048576

int data[DATA_SIZE];

　　并加入一个新函式 GenerateNumbers：

void GenerateNumbers( int *number,  int size)
{
     for( int i =  0; i < size; i++) {
        number[i] = rand() %  10;
    }
}

   　这个函式会产生一大堆 0 ~ 9 之间的随机数。

　　要利用 CUDA 进行计算之前，要先把数据复制到显卡内存中，才能让显示芯片使用。因此，需要取得一块适当大小的显卡内存，再把产生好的数据复制进去。在 main 函式中加入：

    GenerateNumbers(data, DATA_SIZE);

     int* gpudata, *result;
    cudaMalloc(( void**) &gpudata,  sizeof( int) * DATA_SIZE);
    cudaMalloc(( void**) &result,  sizeof( int));
    cudaMemcpy(gpudata, data,  sizeof( int) * DATA_SIZE,
        cudaMemcpyHostToDevice);

　　上面这段程序会先呼叫 GenerateNumbers 产生随机数，并呼叫 cudaMalloc 取得一块显卡内存(result 则是用来存取计算结果，在稍后会用到)，并透过 cudaMemcpy 将产生的随机数复制到显卡内存中。cudaMalloc 和 cudaMemcpy 的用法和一般的 malloc 及 memcpy 类似，不过 cudaMemcpy 则多出一个参数，指示复制内存的方向。在这里因为是从主内存复制到显卡内存，所以使用 cudaMemcpyHostToDevice。如果是从显卡内存到主内存，则使用 cudaMemcpyDeviceToHost。这在之后会用到。

　　接下来是要写在显示芯片上执行的程序。在 CUDA 中，在函式前面加上 __global__ 表示这个函式是要在显示芯片上执行的。因此，加入以下的函式：

__global__  static  void sumOfSquares( int *num,  int* result)
{
     int sum =  0;
     int i;
     for(i =  0; i < DATA_SIZE; i++) {
        sum += num[i] * num[i];
    }

    *result = sum;
}

　　在显示芯片上执行的程序有一些限制，例如它不能有传回值。其它的限制会在之后提到。

　　接下来是要让 CUDA 执行这个函式。在 CUDA 中，要执行一个函式，使用以下的语法：

　　函式名称<<>>(参数...);

　　呼叫完后，还要把结果从显示芯片复制回主内存上。在 main 函式中加入以下的程序：

    sumOfSquares<<< 1,  1,  0>>>(gpudata, result);

     int sum;
    cudaMemcpy(&sum, result,  sizeof( int), cudaMemcpyDeviceToHost);
    cudaFree(gpudata);
    cudaFree(result);

    printf( " sum: %d\n ", sum);

　　因为这个程序只使用一个 thread，所以 block 数目、thread 数目都是 1。我们也没有使用到任何 shared memory，所以设为 0。编译后执行，应该可以看到执行的结果。

　　为了确定执行的结果正确，我们可以加上一段以 CPU 执行的程序代码，来验证结果：

    sum =  0;
     for( int i =  0; i < DATA_SIZE; i++) {
        sum += data[i] * data[i];
    }
    printf( " sum (CPU): %d\n ", sum);

　　编译后执行，确认两个结果相同。

计算运行时间

　　CUDA 提供了一个 clock 函式，可以取得目前的 timestamp，很适合用来判断一段程序执行所花费的时间(单位为 GPU 执行单元的频率)。这对程序的优化也相当有用。要在我们的程序中记录时间，把 sumOfSquares 函式改成：

__global__  static  void sumOfSquares( int *num,  int* result,
    clock_t* time)
{
     int sum =  0;
     int i;
    clock_t start = clock();
     for(i =  0; i < DATA_SIZE; i++) {
        sum += num[i] * num[i];
    }

    *result = sum;
    *time = clock() - start;
}

　　把 main 函式中间部份改成：

     int* gpudata, *result;
    clock_t* time;
    cudaMalloc(( void**) &gpudata,  sizeof( int) * DATA_SIZE);
    cudaMalloc(( void**) &result,  sizeof( int));
    cudaMalloc(( void**) &time,  sizeof(clock_t));
    cudaMemcpy(gpudata, data,  sizeof( int) * DATA_SIZE,
        cudaMemcpyHostToDevice);

    sumOfSquares<<< 1,  1,  0>>>(gpudata, result, time);

     int sum;
    clock_t time_used;
    cudaMemcpy(&sum, result,  sizeof( int), cudaMemcpyDeviceToHost);
    cudaMemcpy(&time_used, time,  sizeof(clock_t),
        cudaMemcpyDeviceToHost);
    cudaFree(gpudata);
    cudaFree(result);

    printf( " sum: %d time: %d\n ", sum, time_used);

　　编译后执行，就可以看到执行所花费的时间了。

　　如果计算实际运行时间的话，可能会注意到它的执行效率并不好。这是因为我们的程序并没有利用到 CUDA 的主要的优势，即并行化执行。