??首先來回顧一下GPU中的內(nèi)存：

• 每個(gè)線程都有自己的私有本地內(nèi)存（Local Memory）和Resigter

• 每個(gè)線程都包含共享內(nèi)存（Shared Memory）,可以被線程中所有的線程共享，其生命周期與線程快一致

• 所有的線程都可以訪問全局內(nèi)存（Global Memory）

• 只讀內(nèi)存塊：常量內(nèi)存（Constant Memory）和紋理內(nèi)存（Texture Memory）

• 每個(gè)SM有自己的L1 cache，SM通過L2 cache連接到Global Memory

內(nèi)存訪問速度

??這一點(diǎn)跟CPU比較像，就是存儲(chǔ)空間越大，訪問速度越慢，GPU內(nèi)存的訪問速度從快到慢依次為：Registers->Caches->Shared Memory->Gloabl Memory。

    

Registers

??寄存器是訪問速度最快的空間，寄存器變量在程序中宏如何使用？

??當(dāng)我們在核函數(shù)中不加修飾的聲明一個(gè)變量，那該變量就是寄存器變量，如果在核函數(shù)中定義了常數(shù)長度的數(shù)組，那也會(huì)被分配到Registers中；寄存器變量是每個(gè)線程私有的，當(dāng)這個(gè)線程的核函數(shù)執(zhí)行完成后，寄存器變量也就不能訪問了。

??寄存器是比較稀缺的資源，空間很小，F(xiàn)ermi架構(gòu)中每個(gè)線程最多63個(gè)寄存器，Kepler架構(gòu)每個(gè)線程最多255個(gè)寄存器；一個(gè)線程中如果使用了比較少的寄存器，那么SM中就會(huì)有更多的線程塊，GPU并行計(jì)算速度也就越快。

??如果一個(gè)線程中變量太多，超出了Registers的空間，這時(shí)寄存器就會(huì)發(fā)生溢出，就需要其他內(nèi)存（Local Memory）來存儲(chǔ)，當(dāng)然程序的運(yùn)行速度也會(huì)降低。

??因此，在程序中，對(duì)于那種循環(huán)操作的變量，我們可以放到寄存器中；同時(shí)要盡量減少寄存器的使用數(shù)量，這樣線程塊的數(shù)量才能增多，整個(gè)程序的運(yùn)行速度才能更快。 

Local Memory

??Local Memory也是每個(gè)線程私有的，在核函數(shù)中符合存儲(chǔ)在寄存器中但不能進(jìn)入核函數(shù)分配的寄存器空間中的變量將被存儲(chǔ)在Local Memory中，Local Memory中可能存放的變量有以下幾種：

• 使用未知索引的本地?cái)?shù)組

• 較大的本地?cái)?shù)組或結(jié)構(gòu)體

• 任何不滿足核函數(shù)寄存器限定條件的變量 

Shared Memory

??每個(gè)SM中都有共享內(nèi)存，使用__shared__關(guān)鍵字（CUDA關(guān)鍵字的下劃線一般都是兩個(gè)）定義，共享內(nèi)存在核函數(shù)中聲明，生命周期和線程塊一致。

??同樣需要注意的是，SM中共享內(nèi)存使用太多，會(huì)導(dǎo)致SM上活躍的線程數(shù)量減少，也會(huì)影響程序的運(yùn)行效率。

??數(shù)據(jù)的共享肯定會(huì)導(dǎo)致線程間的競爭，可以通過同步語句來避免內(nèi)存競爭，同步語句為：

void __syncthreads();

當(dāng)所有線程都執(zhí)行到這一步時(shí)，才能繼續(xù)向下執(zhí)行；頻繁調(diào)用__syncthreads()也會(huì)影響核函數(shù)的執(zhí)行效率。

??共享內(nèi)存被分成了不同個(gè)Bank，我們知道一個(gè)Warp中有32個(gè)SM，在比較老的GPU中，16個(gè)Bank可以同時(shí)訪問，即一條指令就可以讓半個(gè)Warp同時(shí)訪問16個(gè)Bank，這種并行訪問的效率可以極大的提高GPU的效率。比較新的GPU中，一個(gè)Warp即32個(gè)SM可以同時(shí)訪問32個(gè)Bank，效率又提升了一倍。

??下面這個(gè)圖中，左邊的圖每個(gè)線程訪問一個(gè)Bank，不存在內(nèi)存沖突，通過一個(gè)指令即可完成訪問所有的訪問操作；右邊的圖雖然看起來有些亂，但還是一個(gè)線程對(duì)應(yīng)一個(gè)Bank，也不存在沖突，一個(gè)指令即可完成。

    

??下面這個(gè)圖中，存在多個(gè)Thread訪問一個(gè)Bank的情況，如果是讀操作，那么GPU底層可以通過廣播的方式將數(shù)據(jù)傳給各個(gè)Thread，延遲不會(huì)很大，但如果是寫操作，就必須要等上一個(gè)線程寫完成后才能進(jìn)行下一個(gè)線程的寫操作，延時(shí)會(huì)比較大。

     

Constant Memroy

??常量內(nèi)存駐留在設(shè)備內(nèi)存中，每個(gè)SM都有專用的常量內(nèi)存空間，使用__constant__關(guān)鍵字來聲明。

??常量內(nèi)存存在于核函數(shù)之外，在全局范圍內(nèi)聲明，常量內(nèi)容的訪問速度也是很快的，對(duì)所有的核函數(shù)都可見，在Host端進(jìn)行初始化后，核函數(shù)不能再修改。 

Texture Memory

??紋理內(nèi)存的使用并不多，它是為了GPU的顯示而設(shè)計(jì)的，這里不多講了。 

Global Memory

??全局內(nèi)存，就是我們常說的顯存，就是GDDR的空間，全局內(nèi)存中的變量，只要不銷毀，生命周期和應(yīng)用程序是一樣的。

??在訪問全局內(nèi)存時(shí)，要求是對(duì)齊的，也就是一次要讀取指定大?。?2、64、128）整數(shù)倍字節(jié)的內(nèi)存，數(shù)據(jù)對(duì)齊就意味著傳輸效率降低，比如我們想讀33個(gè)字節(jié)，但實(shí)際操作中，需要讀取64字節(jié)的空間。 

GPU緩存

??每個(gè)SM都有一個(gè)一級(jí)緩存，所有SM公用一個(gè)二級(jí)緩存，GPU讀操作是可以使用緩存的，但寫操作不能被緩存。

??每個(gè)SM有一個(gè)只讀常量緩存，只讀紋理緩存，它們用于設(shè)備內(nèi)存中提高來自于各個(gè)內(nèi)存空間內(nèi)的讀取性能。