cuda 本地記憶體使用_CUDA學習筆記（二）

一、CUDA代碼的高效政策

1.高效公式

最大化計算強度：Math/Memory ，即:數學計算量/每個線程的記憶體最大化每個線程的計算量；

最小化每個線程的記憶體讀取速度；每個線程讀取的資料量少

每個線程讀取的速度快内地記憶體>共享記憶體>>全局記憶體

合并全局記憶體

2.合并全局記憶體(按順序讀取的方式最好)

3.避免線程發散

線程發散：同一個線程塊中的線程執行不同内容的代碼

(1)kernel中做條件判斷

(2)循環長度不一

二、Kernel加載方式

(1)檢視本機參數

注：Kernel的加載中，自定義的線程數、線程塊的數量等都不要超過系統本身的設定

(2)Kernel加載1D、2D、3D模式

(3)Kernel函數的關鍵字

三、CUDA中的各種記憶體的代碼使用

1.本地記憶體(變量)

2.全局記憶體(變量)

3.共享記憶體(變量)

四、CUDA同步操作

1.原子操作

(1)解決的問題

對于有很多線程需要同時讀取或寫入相同的記憶體時，保證同一時間隻有一個線程能進行操作

(2)特點隻支援某些運算(加、減、最小值、異或運算等，不支援求餘和求幂等)和資料類型(整型)；

運作順序不定；

安排不當，會使速度很慢(因為内部是個串行的運作)。

2.同步函數

(1)_syncthreads()

線程塊内線程同步；保證線程塊内所有線程都執行到統一位置。

(2)_threadfence()

一個線程調用_threadfence後，該線程在該語句前對全局存儲器或共享存儲器的通路已經全部完成，執行結果對grid在的所有線程可見。

(3)_threadfence_bolck()

一個線程調用_threadfence_block後，該線程在該語句前對全局存儲器或共享存儲器的通路已經全部完成，執行結果對block在的所有線程可見。

注：(2)(3)兩個函數的作用：及時通知其他線程，全局記憶體或共享記憶體内的結果已經讀入或寫入完成了。

3.CPU/GPU同步 cudaStreamSynchronize()/cudaEventSynchronize()主機端代碼中使用cudaThreadSynchronize()函數實作CPU與GPU線程同步運作

kernel啟動後控制權将異步傳回，利用該函數可以确定所有裝置端線程均已運作結束

五、并行化高效政策(一)

1.規約Reduce

2.掃描Scan

3.直方圖第一步：并行計算局部直方圖

第二步：把所有局部直方圖每個分組bin使用Reduction(規約)并行累加起來形成一個總的直方圖

4.壓縮與配置設定

(1)壓縮：在輸出中，對真值輸入配置設定1，對假值輸入配置設定0

壓縮的步驟：通過一個并行判斷把需要進行分析的線程先通過一個判斷全部篩選出來

把篩選好的線程做并行分析

(2)配置設定：輸出項數可以動态的從每一個輸入項計算出

5.分段掃描

6.排序

(1)奇偶排序(轉排序、冒泡算法)

通過比較數組中相鄰的(奇-偶)位置數字對，如果該奇偶對是錯誤的順序(第一個大于第二個)，則交換。下一步重複該操作，針對所有的(奇-偶)位置數字對，如此交替進行下去。

(2)歸并排序

将兩個已經排序的序列合并成一個序列的操作

(3)排序網