如果 kernel 執行時, 不同的 thread 對同一個變數進行讀取與寫入, 會發生甚麼事呢? 當不同thread被同時執行時, 同時去讀取同一個變數, 得到相同的值, 如果把此值運算後同時又寫回去, 最後的值就不一定是多少了, 以下程式開啟 BlockSize*ThreadSize 個 thread 去執行把 global memory 陣列變數加1的動作 :

#include "cuda_runtime.h"

#include "device_launch_parameters.h"

#include <stdio.h>

 

#define ThreadSize 1000

#define BlockSize  10000

#define ArraySize 10

 

__global__ void incKernel(int *a)

{

     int i = (blockIdx.x*blockDim.x + threadIdx.x) % ArraySize;

     a[i] = a[i] + 1;

}

 

int main()

{

     int host_a[ArraySize];

     int *dev_a = 0;

     float elapsedTime;

 

     // setup performance meter from CUDA ----------

     cudaEvent_t start, stop;

     cudaEventCreate(&start);

     cudaEventCreate(&stop);

 

     cudaSetDevice(0);

     cudaMalloc((void**)&dev_a, ArraySize * sizeof(int));

 

     for (int run = 0; run < 10; run++) {

 

         cudaMemset(dev_a, 0, ArraySize * sizeof(int));         //clear

 

         cudaEventRecord(start, 0); //keep start time

         incKernel << <BlockSize, ThreadSize >> > (dev_a); //calculate

         cudaEventRecord(stop, 0); //keep stop time

         cudaEventSynchronize(stop); //wait stop event    

         cudaEventElapsedTime(&elapsedTime, start, stop); 

 

         cudaMemcpy(host_a, dev_a, ArraySize * sizeof(int), cudaMemcpyDeviceToHost);

         //Print result

         printf("run {%d}: ",run);

         for (int i = 0; i < ArraySize; i++) {

              printf("%d ", host_a[i]);

         }

         printf(" t=%f\n",elapsedTime);

     }

     //cudaDeviceSynchronize();

     getchar();

 

     cudaFree(dev_a);

     return 0;

 

}

執行結果如下 :

undefined

這裡可以發現, 雖然每個陣列元速執行的程式碼都一樣, 結果卻不相同. 如果要得到正確答案, 需要用到互斥機制, 在以 CPU 計算環境中, 多由作業系統負責管理, 而在 CUDA 中, 用硬體用 atomic 運算來解決這個問題, 程式修改如下:

int atomicAdd(int* address, int val);

 

__global__ void incKernel(int *a)

{

     int i = (blockIdx.x*blockDim.x + threadIdx.x) % ArraySize;

     //a[i] = a[i] + 1;

     atomicAdd(&a[i], 1);

}

 

 因為 atomic 運算屬於 architecture 指令, 如果不自行宣告 header 亦可, compiler 會自動處理, 只是編輯時會看到 Visual Studio 的警告訊息. 修改後的執行結果如下:

undefined

結果就正確了, 速度快慢則不一定, 在這個例子中是變快了

 

 

 

 

 

 

arrow
arrow
    全站熱搜

    ghostyguo 發表在 痞客邦 留言(0) 人氣()

    E-mail轉寄