cuda 原子鎖多執行緒操作通用原子操作

在專案中，空間中有200w+的點，需要對映到乙個grid_map的600*600的網格中，落入到同乙個格仔的點需要進行一些計算獲得乙個值。對於格仔與格仔之間是並行的，但格仔之中的點需要設計為序列。所以在計算某個格仔中的點時，需要將格仔的值保護起來，只允許乙個執行緒（點）計算並改變。

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

在《gpu高效能程式設計cuda實戰》一書中，提到了通用原子操作的鎖的設計，貼上原始碼：

struct lock 
~lock
(void
) __device__ void
lock
(void
) __device__ void
unlock
(void)}
;...
....
__global__ void
thekernel
(lock mylock)

這裡通過atomiccas和atomicexch兩個函式進行設計，但乙個執行緒lock之後，將mutex置為1，其他執行緒將在while處迴圈等待，直到該執行緒unlock，將mutex重新置於0，剩下的執行緒中再次爭奪鎖。

但是這個結構是存在問題的，我在測試時候發現呼叫thekernel<<<128, 1>>>(lock)可以正常執行，而thekernel<<<1, 128>>>(lock)出現了死鎖，也就是在block中線程數大於1情況中，出現死鎖。百思不得其解…後來查到了出現這種情況的原因：

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

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

考慮到同乙個wrap的執行緒都是『同進退共生死』，那麼我們只能在那個獲得鎖的執行緒退出函式前，就釋放了鎖。看**：

__global__ void
kernel1()
}}intmain()

在程式中，獲得鎖的執行緒進入到if中，並在執行完if之前就釋放了鎖，這樣就解決了同乙個wrap出現死鎖的情況。當然，這樣的寫法不怎麼優美且不魯棒…（但是能用）。另外，這個函式換成這樣寫法就不行了：

__global__ void
kernel2()
}}intmain()

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

不過我在專案中採取了這種方法，將mlock由int變為int陣列，就可以實現多把鎖並行，提高效率，貼上我執行ok的**：

__device__ void
docriticjob
(int thread_index,
float
* mprocess)
__global__ void
kernel2
(int
* mflag,
float
* mprocess)}}
intmain()
;int h_flag[4]
=;float
*dev_process;
int*dev_flag;
cudastatus =
cudamalloc((
void**
)&dev_process,4*
sizeof
(float))
;if(cudastatus != cudasuccess)
cudastatus =
cudamalloc((
void**
)&dev_flag,4*
sizeof
(int))
;if(cudastatus != cudasuccess)
cudastatus =
cudamemcpy
(dev_process, h_process,4*
sizeof
(float
), cudamemcpyhosttodevice);if
(cudastatus != cudasuccess)
cudastatus =
cudamemcpy
(dev_flag, h_flag,4*
sizeof
(int
), cudamemcpyhosttodevice);if
(cudastatus != cudasuccess)
kernel2<<
<2,
3>>
>
(dev_flag,dev_process)
; cudastatus =
cudadevicesynchronize()
;if(cudastatus != cudasuccess)
float outprocess[4]
; cudastatus =
cudamemcpy
(outprocess, dev_process,4*
sizeof
(float
), cudamemcpydevicetohost);if
(cudastatus != cudasuccess)
for(
float mproces : outprocess)
}

可以實現對多個鎖的控制，鎖之間是並行的。

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

__device__ volatile
int sem =0;
__device__ void
acquire_semaphore
(volatile
int*lock)
__device__ void
release_semaphore
(volatile
int*lock)..
....
....
__global__ void
inkernel()

為什麼叫仲裁中介呢？因為這裡採用了每個block的乙個執行緒作為中介進行仲裁，在acquire_semaphore中爭奪鎖，一旦某個block的第乙個執行緒獲得了鎖，那麼剩下的block第乙個執行緒將陷入while迴圈中，同時因為__syncthreads()，導致整個block停下。這就實現了以block為單位的序列。但是這其實也不完美，若要再進一步在block的執行緒中實現序列，則要繼續加入條件判斷。可以看stackoverflow：鏈結

cuda 中不可避免的遇到需要序列計算的情況，可以每個方案都不是完美的，需要根據情況進行取捨，我也繼續學習，應該是存在更好的方案，日後遇到了再新增進來。

cuda 原子鎖多執行緒操作通用原子操作

多執行緒原子操作

Linux使用者層多執行緒無鎖化原子操作

java 多執行緒原子操作一

cuda 原子鎖 多執行緒操作 通用原子操作

多執行緒 原子操作

Linux使用者層多執行緒無鎖化原子操作

java 多執行緒 原子操作 一

相關推薦

cuda 原子鎖多執行緒操作通用原子操作

多執行緒原子操作

java 多執行緒原子操作一