第五篇：CUDA 並行程式中的同步

前言

       在併發，多執行緒環境下，同步是一個很重要的環節。同步即是指程式/執行緒之間的執行順序約定。

       本文將介紹如何通過共享記憶體機制實現塊內多執行緒之間的同步。

       至於塊之間的同步，需要使用到 global memory，代價較為高昂，目前使用的情況也不多，就先不介紹了。

塊內同步函式：__syncthreads ()

       執行緒呼叫此函式後，該執行緒所屬塊中的所有執行緒均執行到這個呼叫點後才會繼續往下執行。

程式碼示例

       使用同步思想優化之前一篇博文中提到的陣列求和程式。在新的程式中，讓每個塊中的第一個執行緒將塊中所有執行緒的運算結果都加起來，然後再存入到結果陣列中。這樣，結果陣列的長度與塊數相等 (原來是和匯流排程數相等)，大大降低了 CPU 端程式求和的工作量以及需要傳遞進/出視訊記憶體的資料 (程式碼下方如果出現紅色波浪線無視之)：

// 相關 CUDA 庫
#include "cuda_runtime.h"
#include "cuda.h"
#include "device_launch_parameters.h"

// 此標頭檔案包含 __syncthreads ()函式
#include "device_functions.h"

#include <iostream>
#include <cstdlib>

using namespace std;

const int N = 100;

// 塊數
const int BLOCK_data = 3; 
// 各塊中的執行緒數
const int THREAD_data = 10; 

// CUDA初始化函式
bool InitCUDA()
{
    int deviceCount; 

    // 獲取顯示裝置數
    cudaGetDeviceCount (&deviceCount);

    if (deviceCount == 0) 
    {
        cout << "找不到裝置" << endl;
        return EXIT_FAILURE;
    }

    int i;
    for (i=0; i<deviceCount; i++)
    {
        cudaDeviceProp prop;
        if (cudaGetDeviceProperties(&prop,i)==cudaSuccess) // 獲取裝置屬性
        {
            if (prop.major>=1) //cuda計算能力
            {
                break;
            }
        }
    }

    if (i==deviceCount)
    {
        cout << "找不到支援 CUDA 計算的裝置" << endl;
        return EXIT_FAILURE;
    }

    cudaSetDevice(i); // 選定使用的顯示裝置

    return EXIT_SUCCESS;
}

// 此函式在主機端呼叫，裝置端執行。
__global__ 
static void Sum (int *data,int *result)
{
    // 宣告共享記憶體 (陣列)
    extern __shared__ int shared[]; 
    // 取得執行緒號
    const int tid = threadIdx.x; 
    // 獲得塊號
    const int bid = blockIdx.x; 

    shared[tid] = 0;
    // 有點像網格計算的思路
    for (int i=bid*THREAD_data+tid; i<N; i+=BLOCK_data*THREAD_data)
    {
        shared[tid] += data[i];
    }
   
    // 塊內執行緒同步函式
    __syncthreads ();

    // 每個塊內索引為 0 的執行緒對其組內所有執行緒的求和結果再次求和
    if (tid == 0) {
        for(int i = 1; i < THREAD_data; i++) { 
            shared[0] += shared[i];
        }
        // result 陣列存放各個塊的計算結果
        result[bid] = shared[0];
    }
}

int main ()
{
    // 初始化 CUDA 編譯環境
    if (InitCUDA()) {
        return EXIT_FAILURE;
    }
    cout << "成功建立 CUDA 計算環境" << endl << endl;

    // 建立，初始化，列印測試陣列
    int *data = new int [N];
    cout << "測試矩陣: " << endl;
    for (int i=0; i<N; i++)
    {
        data[i] = rand()%10;
        cout << data[i] << " ";
        if ((i+1)%10 == 0) cout << endl;
    }
    cout << endl;

    int *gpudata, *result; 
    
    // 在視訊記憶體中為計算物件開闢空間
    cudaMalloc ((void**)&gpudata, sizeof(int)*N); 
    // 在視訊記憶體中為結果物件開闢空間
    cudaMalloc ((void**)&result, sizeof(int)*BLOCK_data);
    
    // 將陣列資料傳輸進視訊記憶體
    cudaMemcpy (gpudata, data, sizeof(int)*N, cudaMemcpyHostToDevice); 
    // 呼叫 kernel 函式 - 此函式可以根據視訊記憶體地址以及自身的塊號，執行緒號處理資料。
    Sum<<<BLOCK_data,THREAD_data,THREAD_data*sizeof (int)>>> (gpudata,result);
    
    // 在記憶體中為計算物件開闢空間
    int *sumArray = new int[BLOCK_data];
    // 從視訊記憶體獲取處理的結果
    cudaMemcpy (sumArray, result, sizeof(int)*BLOCK_data, cudaMemcpyDeviceToHost);
    
    // 釋放視訊記憶體
    cudaFree (gpudata); 
    cudaFree (result);

    // 計算 GPU 每個塊計算出來和的總和
    int final_sum=0;
    for (int i=0; i<BLOCK_data; i++)
    {
        final_sum += sumArray[i];
    }

    cout << "GPU 求和結果為: " << final_sum << endl;

    // 使用 CPU 對矩陣進行求和並將結果對照
    final_sum = 0;
    for (int i=0; i<N; i++)
    {
        final_sum += data[i];
    }
    cout << "CPU 求和結果為: " << final_sum << endl;

    getchar();

    return 0;
}

執行結果

　　

　　PS：矩陣元素是隨機生成的

小結

　　共享記憶體，或者說這個共享陣列是 CUDA 中實現同步最常用的方法。