經歷了很多奇奇怪怪的bug，整理一下。先描述要做的事情以及怎么做：

在項目中，空間中有200w+的點，需要映射到一個grid_map的600*600的網格中，落入到同一個格子的點需要進行一些計算獲得一個值。對于格子與格子之間是并行的，但格子之中的點需要設計為串行。所以在計算某個格子中的點時，需要將格子的值保護起來，只允許一個線程（點）計算并改變。

這里就用到了cuda的通用原子操作。也許有人會問，cuda提供了一些原子操作函數(shù)，能不能直接用呢？cuda提供的原子函數(shù)適用于簡單的單一變量判斷加減，而對于需要復雜的計算操作是力不從心的。但其實，我們要實現(xiàn)的通用原子操作也是基于cuda的原子函數(shù)，我們進行一些設計就可以得到想要的通用原子操作，比如鎖。

方法1.原子鎖

在《GPU高性能編程CUDA實戰(zhàn)》一書中，提到了通用原子操作的鎖的設計，貼上源碼：

struct Lock {int *mutex;Lock(void) {int state = 0;cudaMalloc((void **) &mutex, sizeof(int));cudaMemcpy(mutex, &state, sizeof(int), cudaMemcpyHostToDevice);}~Lock(void) {cudaFree(mutex);}__device__ void lock(void) {while (atomicCAS(mutex, 0, 1) != 0);}__device__ void unlock(void) {atomicExch(mutex, 0);} }; .......__global__ void theKernel(Lock myLock) {myLock.lock();Do_your_job();myLock.unlock(); }

這里通過atomicCAS和atomicExch兩個函數(shù)進行設計，但一個線程lock之后，將mutex置為1，其他線程將在while處循環(huán)等待，直到該線程unlock，將mutex重新置于0，剩下的線程中再次爭奪鎖。

但是這個結構是存在問題的，我在測試時候發(fā)現(xiàn)調用theKernel<<<128, 1>>>(lock)可以正常運行，而theKernel<<<1, 128>>>(lock)出現(xiàn)了死鎖，也就是在block中線程數(shù)大于1情況中，出現(xiàn)死鎖。百思不得其解…后來查到了出現(xiàn)這種情況的原因：

cuda運行是以wrap為單位進行的，也就是說一個wrap中32個線程中的一個獲得了鎖，執(zhí)行完了lock，按理說該線程要繼續(xù)執(zhí)行Do_your_job()再unlock，而現(xiàn)實是線程都卡在了lock處。這就是因為wrap的同步執(zhí)行規(guī)則(locked-step execution)，換句話說，一個wrap的線程是同步執(zhí)行一個函數(shù)，并同步退出一個函數(shù)。獲得鎖的線程在lock函數(shù)結束處苦苦等待其他31個線程兄弟一起進入Do_your_job()，而剩下的31個線程卻等著它unlock釋放鎖，所以出現(xiàn)了死鎖。而每個block中只有一個線程則不會出現(xiàn)死鎖，是因為此時wrap中僅有一個線程。

顯然，這個設計方法并不滿足我的需求。

方法2.通用原子操作

考慮到同一個wrap的線程都是‘同進退共生死’，那么我們只能在那個獲得鎖的線程退出函數(shù)前，就釋放了鎖。看代碼：

__global__ void kernel1(){int index = 0;int mSize = 1;bool blocked = true;while(blocked) {if(0 == atomicCAS(&mLock, 0, 1)) {index = mSize++;doCriticJob();atomicExch(&mLock, 0);blocked = false;}} } int main(){kernel1<<<4,128>>>();cudaDeviceSynchronize(); }

在程序中，獲得鎖的線程進入到if中，并在執(zhí)行完if之前就釋放了鎖，這樣就解決了同一個wrap出現(xiàn)死鎖的情況。當然，這樣的寫法不怎么優(yōu)美且不魯棒…（但是能用）。另外，這個函數(shù)換成這樣寫法就不行了：

__global__ void kernel2(){int index = 0;int mSize = 1;while(true) {if(0 == atomicCAS(&mLock, 0, 1)) {index = mSize++;doCriticJob();atomicExch(&mLock, 0);break}} } int main(){kernel2<<<4,128>>>();cudaDeviceSynchronize(); }

這是因為break在不同的機器和編譯器中，不能都保證是先釋放了鎖再break出來,可能被編譯器優(yōu)化成其他形式。可以看出這種cuda通用原子操作確實比較蛋疼。

不過我在項目中采取了這種方法，將mLock由int變?yōu)閕nt數(shù)組，就可以實現(xiàn)多把鎖并行，提高效率，貼上我運行ok的代碼：

__device__ void doCriticJob(int thread_index, float* mProcess) {mProcess[thread_index] += 0.222;printf("thread is: %d \n", threadIdx.x); } __global__ void kernel2(int* mFlag, float* mProcess) {bool blocked = true;int thread_index = (threadIdx.x + blockDim.x*blockIdx.x) % 4;while (blocked) {if (0 == atomicCAS(&mFlag[thread_index], 0, 1)) {doCriticJob(thread_index, mProcess);atomicExch(&mFlag[thread_index], 0);blocked = false;}} }int main() {cudaError_t cudaStatus;float h_Process[4] = {0};int h_Flag[4] = {0};float *dev_Process;int *dev_Flag;cudaStatus = cudaMalloc((void **)&dev_Process, 4*sizeof(float));if(cudaStatus != cudaSuccess){ fprintf(stderr,"malloc 1 failed\n");}cudaStatus = cudaMalloc((void **)&dev_Flag, 4*sizeof(int));if(cudaStatus != cudaSuccess){ fprintf(stderr,"malloc 2 failed\n");}cudaStatus = cudaMemcpy(dev_Process, h_Process, 4*sizeof(float), cudaMemcpyHostToDevice);if(cudaStatus != cudaSuccess){ fprintf(stderr,"malloc 3 failed\n");}cudaStatus = cudaMemcpy(dev_Flag, h_Flag, 4*sizeof(int), cudaMemcpyHostToDevice);if(cudaStatus != cudaSuccess){ fprintf(stderr,"malloc 4 failed\n");}kernel2<<<2, 3>>>(dev_Flag,dev_Process);cudaStatus = cudaDeviceSynchronize();if (cudaStatus != cudaSuccess) {fprintf(stderr, "ffffff is %d\n",cudaStatus);}float outProcess[4];cudaStatus = cudaMemcpy(outProcess, dev_Process, 4*sizeof(float), cudaMemcpyDeviceToHost);if (cudaStatus != cudaSuccess) {fprintf(stderr, "kkkkkk is %d\n",cudaStatus);}for (float mProces : outProcess) {std::cout << mProces << std::endl;} }

可以實現(xiàn)對多個鎖的控制，鎖之間是并行的。

方法3.仲裁中介

這個方法的名字是我瞎取的。直接上代碼：

__device__ volatile int sem = 0;__device__ void acquire_semaphore(volatile int *lock){while (atomicCAS((int *)lock, 0, 1) != 0);}__device__ void release_semaphore(volatile int *lock){*lock = 0;__threadfence();}..........__global__ void inKernel(){...__syncthreads();if (threadIdx.x == 0)acquire_semaphore(&sem);__syncthreads();//begin critical section// ... your critical section code goes here//end critical section__threadfence(); // not strictly necessary for the lock, //but to make any global updates in the critical //section visible to other threads in the grid__syncthreads();if (threadIdx.x == 0)release_semaphore(&sem);__syncthreads();... }

為什么叫仲裁中介呢？因為這里采用了每個block的一個線程作為中介進行仲裁，在acquire_semaphore中爭奪鎖，一旦某個block的第一個線程獲得了鎖，那么剩下的block第一個線程將陷入while循環(huán)中，同時因為__syncthreads()，導致整個block停下。這就實現(xiàn)了以block為單位的串行。但是這其實也不完美，若要再進一步在block的線程中實現(xiàn)串行，則要繼續(xù)加入條件判斷。可以看StackOverflow：鏈接

總結：

cuda 中不可避免的遇到需要串行計算的情況，可以每個方案都不是完美的，需要根據(jù)情況進行取舍，我也繼續(xù)學習，應該是存在更好的方案，日后遇到了再添加進來。

總結

以上是生活随笔為你收集整理的cuda 原子锁多线程操作通用原子操作的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。