[CUDA] 多重讀取與寫入的互斥問題

如果 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; }

 

執行結果如下:

這裡可以發現, 雖然每個陣列元速執行的程式碼都一樣, 結果卻不相同. 如果要得到正確答案, 需要用到互斥機制, 在以 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 的警告訊息. 修改後的執行結果如下:

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