CUDA
示例 > CUDA Driver API

下面將示範以最底層的CUDA Driver API （页面存档备份，存于）調用GPU做列向量的加法，以下為 CPU 端的程式碼
// 本範例修改自 Andrei de A. Formiga (2012-06-04) 寫的範例: https://gist.github.com/tautologico/2879581
// 編譯指令 nvcc -O3 -lcuda add.c -o add.exe

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

#include <cuda.h>
#include <builtin_types.h> // Driver api 的型態定義

#define N 1024 //列向量長度

// 利用CUDA函數的錯誤傳回做例外處理
inline void checkCudaErrors( CUresult err)
{
    if( CUDA_SUCCESS != err) {
        printf("CUDA Driver API error = %04d from file <%s>, line %i.\n",
                err, __FILE__, __LINE__ );
        exit(-1); // 直接終止程式
    }
}

CUdevice   device; // CUDA 裝置(也就是GPU)物件
CUcontext  context; // CUDA 內容物件
CUmodule   module; // 代表GPU程式碼的物件
CUfunction function; // CUDA GPU 函數
size_t     totalGlobalMem; // CUDA 裝置記憶體總量

// Driver API 只能自外部檔案讀取 GPU 程式，可以為 PTX 中間碼也可以是 cubin 機器碼(或是混合各種架構機器碼的fatbin)
char       *module_file = (char*) "matSumKernel.cubin";

// GPU 函數名稱
char       *kernel_name = (char*) "matSum";

// 初始化 CUDA 的手續
void initCUDA()
{
    int deviceCount = 0; // 當前可使用的 CUDA 裝置(GPU)數
    CUresult err = cuInit(0); // 初始化 CUDA API

    if (err == CUDA_SUCCESS) // 取得可用裝置數
        checkCudaErrors(cuDeviceGetCount(&deviceCount));

    if (deviceCount == 0) { // 確定有可用的裝置
        fprintf(stderr, "Error: no devices supporting CUDA\n");
        exit(-1);
    }

    // get first CUDA device
    checkCudaErrors(cuDeviceGet(&device, 0)); // 取編號為 0 的裝置
    char name[100];
    cuDeviceGetName(name, 100, device); // 印出裝置名稱
    printf("> Using device 0: %s\n", name);

    checkCudaErrors( cuDeviceTotalMem(&totalGlobalMem, device) ); 
    // 印出裝置可用記憶體
    printf("  Total amount of global memory:   %llu bytes\n",
           (unsigned long long)totalGlobalMem);
    // GPU 記憶體是否是為64bits定址
    printf("  64-bit Memory Address:           %s\n",
           (totalGlobalMem > (unsigned long long)4*1024*1024*1024L)?
           "YES" : "NO");

    // 創建 CUDA 內容
    err = cuCtxCreate(&context, 0, device);
    if (err != CUDA_SUCCESS) {
        fprintf(stderr, "* Error initializing the CUDA context.\n");
        cuCtxDetach(context);
        exit(-1);
    }

    // 讀取編譯好的cubin GPU程式碼
    err = cuModuleLoad(&module, module_file);
    if (err != CUDA_SUCCESS) {
        fprintf(stderr, "* Error loading the module %s\n", module_file);
        cuCtxDetach(context); // 釋放 CUDA 內容物件
        exit(-1);
    }

    // 獲取GPU程式裡函數"matSum"的指標
    err = cuModuleGetFunction(&function, module, kernel_name);

    if (err != CUDA_SUCCESS) {
        fprintf(stderr, "* Error getting kernel function %s\n", kernel_name);
        cuCtxDetach(context);
        exit(-1);
    }
}

int main(int argc, char **argv)
{
    int a[N], b[N], c[N];
    CUdeviceptr d_a, d_b, d_c;

    // 注意 GPU 變數指標的型態是 CUdeviceptr
    // typedef unsigned int CUdeviceptr_v2
    // typedef CUdeviceptr_v2 CUdeviceptr

    // 初始化主記憶體變數
    for (int i = 0; i < N; ++i) {
        a[i] = i;
        b[i] = N - i;
    }

    initCUDA();

    // 動態分配 GPU 記憶體
    // CUresult cuMemAlloc ( CUdeviceptr* dptr, size_t bytesize )
    checkCudaErrors( cuMemAlloc(&d_a, sizeof(int) * N) ); // 
    checkCudaErrors( cuMemAlloc(&d_b, sizeof(int) * N) );
    checkCudaErrors( cuMemAlloc(&d_c, sizeof(int) * N) );

    // 將列向量傳入裝置
    // CUresult cuMemcpyHtoD ( CUdeviceptr dstDevice, const void* srcHost, size_t ByteCount ) 
    checkCudaErrors( cuMemcpyHtoD(d_a, a, sizeof(int) * N) );
    checkCudaErrors( cuMemcpyHtoD(d_b, b, sizeof(int) * N) );

    void *args[3] = { &d_a, &d_b, &d_c }; // 包裝放入GPU 函數的引數
    
    // 運行 GPU 函數
    // CUresult cuLaunchKernel ( CUfunction f, unsigned int  gridDimX, unsigned int  gridDimY, unsigned int  gridDimZ,
    //                           unsigned int  blockDimX, unsigned int  blockDimY, unsigned int  blockDimZ,
    //                           unsigned int  sharedMemBytes, CUstream hStream, void** kernelParams, void** extra )
    
    checkCudaErrors( cuLaunchKernel(function, N, 1, 1,  // Nx1x1 blocks
                                    1, 1, 1,            // 1x1x1 threads
                                    0, 0, args, 0) );

    // 將運算結果送回主記憶體
    // CUresult cuMemcpyDtoH ( void* dstHost, CUdeviceptr srcDevice, size_t ByteCount ) 
    checkCudaErrors( cuMemcpyDtoH(c, d_c, sizeof(int) * N) );
    
    // 將 CPU 和 GPU 運算結果做對照 
    for (int i = 0; i < N; ++i) {
        if (c[i] != a[i] + b[i])
            printf("* Error at array position %d: Expected %d, Got %d\n",
                   i, a[i]+b[i], c[i]);
    }

    // 釋放 GPU 記憶體
    // CUresult cuMemFree ( CUdeviceptr dptr ) 
    checkCudaErrors( cuMemFree(d_a) );
    checkCudaErrors( cuMemFree(d_b) );
    checkCudaErrors( cuMemFree(d_c) );

    cuCtxDetach(context); // 釋放 CUDA 內容物件
    return 0;
}
而以下是GPU端的程式碼
// 本範例修改自 Andrei de A. Formiga (2012-06-04) 寫的範例: https://gist.github.com/tautologico/2879581
// 此部分要先編譯成 cubin 後才可以被 CPU 端程式使用
// 編譯指令 nvcc -O3 -cubin -arch=native matSumKernel.cu -o matSumKernel.cubin

#define N 1024 //列向量長度

extern "C" __global__ void matSum(int *a, int *b, int *c)
{
    int tid = blockIdx.x; // thread 的 x 座標
    if (tid < N)
        c[tid] = a[tid] + b[tid]; //每個 thread 做一次加法
}