如果 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 的警告訊息. 修改後的執行結果如下:
結果就正確了, 速度快慢則不一定, 在這個例子中是變快了
留言列表
雜項記錄 (3)
