CUDA代碼筆記（二） cudaOpenMP

cudaOpenMP項目展示了如何在cuda項目中運用openmp技術。

該項目位于cuda samples檔案夾下的0_Simple/cudaOpenMP檔案夾下。

在正式開始剖析代碼之前，讓我們先來了解一下openmp的背景知識。

OpenMP

根據百度百科，OpenMP是用于共享記憶體并行系統的多線程程式設計的一套指導性的編譯處理方案。 簡單來講，我們可以把OpenMP了解為一個并行線程庫，它跟linux下的pthread庫很類似，但是相比後者而言，OpenMP具有更加簡單的程式設計邏輯。 OpenMP目前支援兩類語言：C/C++和Fortran，對于C/C++程式員來說，使用OpenMP非常容易，目前市面上流行的各種編譯器：gcc、visual c++ cl以及clang均不同程度地實作了對OpenMP的支援，更多的資訊可以從下面幾個連結中獲得： OpenMP官網： http://openmp.org GCC OpenMP： http://gcc.gnu.org/wiki/openmp Visual C++ OpenMP：http://msdn.microsoft.com/zh-cn/library/vstudio/tt15eb9t.aspx Clang OpenMP： http://clang-omp.github.io/

代碼筆記

注：代碼筆記中的代碼并非完整的項目代碼，有一些部落客認為不太重要的部分被省略了，想要檢視完整代碼的讀者請直接檢視CUDA Samples

首先我們來關注在代碼中GPU做了哪些事情，為此我們檢視相應的__global__函數，其代碼如下：

__global__ void kernelAddConstant(int *g_a, const int b) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    g_a[idx] += b;
}
           

可以看到，每一個CUDA線程僅僅從全局記憶體空間g_a中讀取其線程ID對應位置的數值，并将其與一個常量b相加。 這個函數非常簡單，這樣我們更容易将注意力集中在如何結合CUDA和OpenMP上。

接下來，我們将目光轉向main函數，main函數主要包括3個部分： 1、初始化 2、執行計算 3、正确性檢查

先來看"初始化"部分：

int num_gpus = 0;

printf("%s Starting...\n\n", argv[0]);

cudaGetDeviceCount(&num_gpus);

printf("number of host CPUs:\t%d\n", omp_get_num_procs());
printf("number of CUDA devices:\t%d\n", num_gpus);

unsigned int n = num_gpus * 8192;
unsigned int nbytes = n * sizeof(int);
int *a = 0;
int b = 3;
a = (int *)malloc(nbytes);

for (unsigned int i = 0; i < n; ++i)
    a[i] = i;
           

上面的代碼中，函數cudaGetDeviceCount()用來獲得系統中可用的GPU數目，它的值儲存在變量num_gpus中。omp_get_num_procs()是OpenMP提供的庫函數，它的作用是擷取函數調用時系統中可用的CPU數目。變量n儲存待處理的整數的數目，nbytes儲存這些整數所需要的記憶體空間。指針a訓示這部分空間的起始位址，代碼最後對這部分記憶體的每一個元素賦初值。至此，初始化過程結束。

下面進入到"執行計算"過程：

omp_set_num_threads(2 * num_gpus);
#pragma omp parallel
{
    unsigned int cpu_thread_id = omp_get_thread_num();
    unsigned int num_cpu_threads = omp_get_num_threads();
    
    int gpu_id = -1;
    checkCudaErrors(cudaSetDevice(cpu_thread_id % num_gpus));
    checkCudaErrors(cudaGetDevice(&gpu_id));
    printf("CPU thread %d (of %d) uses CUDA device %d\n", cpu_thread_id, num_cpu_threads, gpu_id);

    int *d_a = 0;
    int *sub_a = a + cpu_thread_id * n / num_cpu_threads;
    unsigned int nbytes_per_kernel = nbytes / num_cpu_threads;
    dim3 gpu_threads(128);
    dim3 gpu_blocks(n / (gpu_threads.x * num_cpu_threads));

    checkCudaErrors(cudaMalloc((void **)&d_a, nbytes_per_kernel));
    checkCudaErrors(cudaMemset(d_a, 0, nbytes_per_kernel));
    checkCudaErrors(cudaMemcpy(d_a, sub_a, nbytes_per_kernel, cudaMemcpyHostToDevice));
    kernelAddConstant<<<gpu_blocks, gpu_threads>>>(d_a, b);

    checkCudaErrors(cudaMemcpy(sub_a, d_a, nbytes_per_kernel, cudaMemcpyDeviceToHost));
    checkCudaErrors(cudaFree(d_a));
}
           

計算過程一開始首先調用openmp函數omp_set_num_threads()設定執行并行代碼片段的CPU線程數目。 然後，使用#pragma omp parallel訓示下面的代碼塊将會使用openmp多線程執行。

omp_get_thread_num()函數傳回執行目前代碼的線程編号。 omp_get_num_threads()傳回執行該代碼塊的線程總數。 也就是說，對于該代碼片段，總共有2 * num_gpus個CPU線程将會執行它。 之後，代碼對總的工作量進行了劃分，将所有工作均勻配置設定給每一個GPU核心，并啟動核心執行，然後将執行結果拷貝回來。 指針d_a指向GPU全局記憶體空間，sub_a指向執行該代碼塊的CPU線程所對應的主機記憶體空間的開始位置。代碼中為每個CUDA block配置設定了128個CUDA線程，是以每個Grid需要的block數目為n / (num_cpu_threads * 128)，至此完成劃分開始拷貝資料和計算。 最後，使用cudaFree()函數釋放GPU記憶體。

最後是"正确性檢查"部分：

if (cudaSuccess != cudaGetLastError())
    printf("%s\n", cudaGetErrorString(cudaGetLastError()));

bool bResult = correctResult(a, n, b);

if (a) free(a);

cudaDeviceReset();

exit(bResult ? EXIT_SUCCESS : EXIT_FAILURE);
           

這部分代碼首先檢查GPU執行是否出錯，然後調用函數correctResult()檢查執行結果是否正确，最後釋放記憶體并退出。 correctResult()函數功能很簡單，就是對比GPU執行結果和CPU執行結果：

int correctResult(int *data, const int n, const int b) {
    for (int i = 0; i < n; ++i)
        if (data[i] !== i + b) return 0;

    return 1;
}