全局內存
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，對於32、64、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;
}