CUDA_全局內存及訪問優化

全局內存

GPU全局內存，CPU和GPU均可以進行讀寫操做。任何設備均可以經過PCI-E總線對其進行訪問，GPU之間不經過CPU，直接將數據從一塊GPU卡上的數據傳輸到另外一塊GPU上。

點對點的特性實在DUDA4.x SDK中引入。只對特定平臺進行支持（特斯拉硬件經過TCC驅動模型可以支持windows7和windows Vista平臺，對於linux或windowsXP平臺，消費機GPU卡和特斯拉卡都支持）。

CPU主機端處理器能夠經過如下三種方式對GPU上的內存進行訪問：

• 顯式地阻塞傳輸；
• 顯式地非阻塞傳輸；
• 隱式地使用零拷貝內存複製。

一旦數據進入到GPU，主要問題就成了如何在GPU中進行高效訪問。經過建立一個每十次計算只需一次訪存的模式，內存延遲能明顯的被隱藏，但前提是對全局內存的訪問必須是以合併的方式進行訪問。

對全局內存的訪問是否知足合併訪問條件是對CUDA程序性能影響最明顯的因素之一。

合併訪問--全局存儲器訪問優化

全部線程訪問連續的對齊的內存塊。

若是咱們對內存進行一對一連續對齊訪問，則每一個線程的訪問地址能夠合併起來，只需一次存儲食物便可解決問題。假設咱們訪問一個單精度或者整型值，每一個線程將訪問一個4字節的內存塊。內存會基於線程束的方式進行合併（老式的G80硬件上使用半個線程束），也就是說訪問一次內存將獲得32*4=128個字節的數據。

合併大小支持32字節、64字節、128字節，分貝標識線程束中每一個線程一個字節、16位以及32位爲單位讀取數據，但前提是訪問必須連續，而且以32字節位基準對其。

將標準的cudaMalloc替換爲cudaMallocPitch，能夠分配到對齊的內存塊。

extern __host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height);

該方法的第一個參數表示指向設備內存指針的指針，第二個參數表示指向對齊以後每行真實字節數的指針，第三個參數爲須要開闢的數據的寬度，單位爲字節，最後一個參數爲數組的高度。

合併訪問條件要求同一warp或者同一half-warp中的線程要按照必定字長訪問通過對齊的段。

不一樣設備中合併訪問的具體要求：

• 計算能力1.0、1.1設備上，一個half-warp中的第k個線程必須訪問段裏面的第k個字，而且half-warp訪問的段的地址必須對齊到每一個線程訪問的字長的16倍。只支持對字長32bit、64bit、128bit的數據的合併訪問。
• 在1.2及更高能力的設備上，合併訪問要求大大放寬，支持字長爲8bit（對應段長32Byte）、16bit（對應段長64Byte）、32bit/64bit/128bit（對應段長128Byte）的數據進行合併訪問。

下面描述1.2/1.3能力硬件的一個half-warp是如何完成一次合併訪問的。

• 首先，找到有最低線程號活動線程（前half-warp中的線程0，或者後half-warp中的線程16）請求訪問的地址所在段。對於8bit數據來講，段長爲32Byte，對於16bit數據來講段長爲64Byte，對於3二、6四、128bit數據來講段長爲128Byte。
• 而後，找到所請求訪問的地址也在這個段內的活動線程。若是全部線程訪問的數據都處於段的前半部分或者後半部分，那麼還能夠減小一次傳輸的數據大小。例如，若是一個段的大小爲128Byte，但只有上半部分或下半部分被使用了，那麼實際傳輸的數據大小就能夠進一步減少到64Byte，同理，對於64Byte的段的合併傳輸，在只有前半或者後半被使用的狀況下也能夠繼續減少到32Byte。
• 進行傳輸，此時，執行訪存指令的線程將處於不活動狀態，執行資源被釋放供SM中處於就緒態的其餘warp使用。
• 重複上述過程，知道half-warp全部線程均訪問結束。

須要注意的是，經過運行時API(如cudaMalloc())分配的存儲器，已經能保證其首地址至少會按256Byte進行對齊。所以，選擇合適的線程塊大小（例如16的整數倍），能使half-warp的訪問請求按段長對齊。使用__align__(8)和__align__(16)限定符來定義結構體，可使對結構體構成的數組進行訪問時可以對齊到段。

訪問時段不對齊或者間隔訪問都會要成有效帶寬的大幅度下降。對於間隔訪問顯存的狀況，能夠藉助shared memory來實現。

全局內存分配

當使用CUDA運行時時，設備指針與主機指針類型均爲void*。

動態內存分配

大多數CUDA中的全局內存經過動態分配獲得，使用cuda運行時，經過如下函數分別進行全局內存的分配和釋放。

cudaError_t cudaMalloc(void **, size_t);
cudaError_t cudaFree(void);

對應的驅動程序API函數爲：

CUresult CUDAAPI cuMemAlloc(CUdeviceptr *dptr, size_t bytesize);
CUresult CUDAAPI cuMemFree(CUdeviceptr dptr);

分配全局內存成本較大，CUDA驅動程序實現了一個CUDA小型內存請求的子分配器（suballocator），可是若是這個suballocator必須建立一個新的內存塊，這須要調用操做系統的一個成本很高的內核模式驅動程序。若是這種狀況發生，CUDA驅動程序必須與GPU同步，這可能會中斷CPU、GPU的併發，所以，在性能要求很高的代碼中避免分配或釋放全局內存時一個較好的作法。

靜態內存分配

經過使用__device__關鍵字標記在內存聲明中進行標記便可。這一內存是由cuda驅動程序在模塊加載時分配的。

運行時API：

cudaError_t cudaMemcpyToSymbol(
	char *symbol,
	const void *src,
	size_t count,
	size_t offset=0,
	enum cudaMemcpyKind kind=cudaMemcpyHostToDevice
);

cudaError_t cudaMemcpyFromSymbol(
	void *dst,
	char *symbol,
	size_t count,
	size_t offset,
	enum cudaMemcpyKind kind=cudaMemcpyDeviceToHost
);

cuda運行時應用程序能夠經過調用函數cudaGetSymbolAddress()查詢關聯到靜態分配的內存上的指針。

cudaError_t cudaGetSymbolAddress(void **devPtr, char *symbol);

驅動程序API：

CUresult CUDAAPI cuModuleGetGlobal(CUdeviceptr *dptr, size_t *bytes, CUmodule hmod, const char *name);

該函數返回基指針和對象大小。若是咱們不須要大小，能夠在bytes參數傳入NULL。

指針查詢

cuda跟蹤全部內存分配，並提供API使應用程序能夠查詢CUDA中的全部指針。函數庫和插件能夠在基礎之上使用不一樣的處理策略。

struct cudaPointerAttributes{
  	enum cudaMemoryType memoryType;
  	int device;
  	void *devicePointer;
  	void *hostPointer;
}