先來點題外話,由於導師的要求使用gpu加速但是又沒有系統的學習就開始了使用pycuda進行gpu的程式設計,變用邊學。
在優化程式的時候shared memory是必用的,所以理解共享記憶體是個什麼東西才可以做到自己對資料是如何跑的有數。
先看看這張gpu的儲存結構圖(偷的圖,哈哈^_^):
在圖中我們關心的重點是,每乙個block都對應的乙個shared memory,執行緒塊之可以相互交流通訊就是通過共享記憶體進行傳遞的。
我們要記住的一點就是每乙個執行緒塊都有自己對用的一塊共享記憶體。這麼說,是感受不到的,下面我就用乙個簡單的程式來形象的看看這個效果。
import numpy as np
import pycuda.autoinit
import pycuda.driver as cuda
from pycuda.compiler import sourcemodule
mod = sourcemodule('''
__global__ void text_gpu(float *a , float *b)
if(tid == 0)
b[bid] = s_data[0];
}''')
text_gpu = mod.get_function("text_gpu")
a = np.ones((32,128),dtype=np.float32)
b = np.ones((32,),dtype=np.float32)
text_gpu(cuda.in(a),cuda.inout(b),grid=(32,1),block=(128,1,1))
print(b)
這個程式非常的簡單,我們可以看到建立了a(32,128)的乙個矩陣再複製給共享記憶體進行乙個規約之後將和存入陣列b。
其中值得體會的是程式中的這句話:
s_data[tid] =a[bid*128 + tid];
s_data只有128 而 a的大小是128*32,要是我們不知道gpu的儲存結構那我們就一頭霧水。
下面講解下這個知識點:
乙個執行緒塊開了128的共享記憶體,而同時呼叫的有32個執行緒塊,每乙個塊都有128的共享記憶體,所以就滿足了和a的大小相等了。
看看輸出的結果:
對a矩陣的每一列進行求和,結果放入b中。結果有32個說明有32共享記憶體一起執行。
好了,這個知識點記錄完了,88。0.0
GPU 共享記憶體位址對映方式
當乙個warp中的不同執行緒訪問乙個bank中的不同的字位址時,就會發生bank衝突。如果沒有bank衝突的話,共享記憶體的訪存速度將會非常的快,大約比全域性記憶體的訪問延遲低100多倍,但是速度沒有暫存器快。然而,如果在使用共享記憶體時發生了bank衝突的話,效能將會降低很多很多。在最壞的情況下,...
GPU記憶體分類
全域性記憶體 通俗意義上的裝置記憶體。共享記憶體 1.位置 裝置記憶體。2.形式 關鍵字 shared 新增到變數宣告中。如 shared float cache 10 3.目的 對於gpu上啟動的每個執行緒塊,cuda c編譯器都將建立該共享變數的乙個副本。執行緒塊中的每個執行緒都共享這塊記憶體,...
CUDA學習 GPU記憶體
通俗意義上的裝置記憶體 位置 裝置記憶體。形式 關鍵字 shared 新增到變數宣告中。如 shared float cache 10 目的 對於gpu上啟動的每個執行緒塊,cuda c編譯器都將建立該共享變數的乙個副本。執行緒塊中的每個執行緒都共享這塊記憶體,但執行緒卻無法看到也不能修改其他執行緒...