CUDA 內(nèi)存概述

GPU的內(nèi)存包括：

• 全局內(nèi)存（global memory）
• 常量?jī)?nèi)存（constant memory）
• 紋理內(nèi)存核表面內(nèi)存（texture memory）
• 寄存器（register）
• 局部?jī)?nèi)存（local memory）
• 共享內(nèi)存（shared memory）
• L1、L2緩存（從費(fèi)米架構(gòu)開(kāi)始有了SM層次的 L1 cache 和設(shè)備層次的 L2 cache）

速度快慢如下圖所示：

CUDA 內(nèi)存詳解

Global Memory

Global Memory在某種意義上等同于GPU顯存，kernel函數(shù)通過(guò)Global Memory來(lái)讀寫(xiě)顯存。Global Memory是kernel函數(shù)輸入數(shù)據(jù)和寫(xiě)入結(jié)果的唯一來(lái)源。

Rigisters 寄存器

寄存器是GPU最快的memory，kernel中沒(méi)有什么特殊聲明的自動(dòng)變量都是放在寄存器中的。當(dāng)數(shù)組的索引是constant類型且在編譯期能被確定的話，就是內(nèi)置類型，數(shù)組也是放在寄存器中。

• 寄存器變量是每個(gè)線程私有的，一旦thread執(zhí)行結(jié)束，寄存器變量就會(huì)失效
• 寄存器是稀有資源。(省著點(diǎn)用，能讓更多的block駐留在SM中，增加Occupancy)
• --maxrregcount 可以設(shè)置大小
• 不同設(shè)備架構(gòu)，數(shù)量不同

Shared Memory

Shared Memory位于GPU芯片上，訪問(wèn)延遲僅次于寄存器。Shared Memory是可以被一個(gè)Block中的所有Thread來(lái)進(jìn)行訪問(wèn)的，可以實(shí)現(xiàn)Block內(nèi)的線程間的低開(kāi)銷通信。在SMX中，L1 Cache跟Shared Memory是共享一個(gè)64KB的告訴存儲(chǔ)單元的，他們之間的大小劃分不同的GPU結(jié)構(gòu)不太一樣；
用 __shared__ 修飾符修飾的變量存放在shared memory:

• On-chip
• 擁有高的多bandwidth和低很多的latencyo
• 同一個(gè)Block中的線程共享一塊Shared Memoryo
• 需要使用 __syncthreads() 同步。
• 比較小，要節(jié)省著使用，不然會(huì)限制活動(dòng)warp的數(shù)量

Local Memory

Local Memory本身在硬件中沒(méi)有特定的存儲(chǔ)單元，而是從Global Memory 虛擬出來(lái)的地址空間。** Local Memory 是為寄存器無(wú)法滿足存儲(chǔ)需求的情況而設(shè)計(jì)的，主要是用于存放單線程的大型數(shù)組和變量。** Local Memory是線程私有的，線程之間是不可見(jiàn)的。由于GPU硬件單位沒(méi)有Local Memory的存儲(chǔ)單元，所以，針對(duì)它的訪問(wèn)是比較慢的。

但是更多在以下情況，會(huì)使用 Local Memory：

• 無(wú)法確定其索引是否為常量的數(shù)組。
• 會(huì)消耗太多寄存器空間的大型結(jié)構(gòu)或數(shù)組。
• 如果內(nèi)核使用了多于可用寄存器的任何變量(這也稱為寄存器溢出)
• --ptxas-options=-v

Constant Memory

Constant Memory （常量?jī)?nèi)存） 類似于 Local Memory，也是沒(méi)有特定的存儲(chǔ)單元的，只是Global Memory 的虛擬地址。因?yàn)樗侵蛔x的，所以簡(jiǎn)化了緩存管理，硬件無(wú)需管理復(fù)雜的回寫(xiě)策略。Constant Memory 啟動(dòng)的條件是同一個(gè)warp所有的線程同時(shí)訪問(wèn)同樣的常量數(shù)據(jù)。

其具有以下幾個(gè)特點(diǎn)：

• constant的范圍是全局的，針對(duì)所有kernel。
• 在同一個(gè)編譯單元，constant對(duì)所有kernel可見(jiàn)。
• kernel只能從constant Memory讀取數(shù)據(jù)，因此其初始化必須在host端使用下面的函數(shù)調(diào)用: cudaError_t cudaMemcpyToSymbollconst void* symbol, const void* src,size t count);
• 當(dāng)一個(gè)warp中所有thread都從同一個(gè)Memory地址讀取數(shù)據(jù)時(shí)，constant Memory表現(xiàn)會(huì)非常好會(huì)觸發(fā)廣播機(jī)制。

Texture Memory

Texture Memory是GPU的重要特性之一，也是GPU編程優(yōu)化的關(guān)鍵。Texture Memory實(shí)際上也是Global Memory的一部分，但是它有自己專用的只讀cache。這個(gè)cache在浮點(diǎn)運(yùn)算很有用，Texture Memory是針對(duì)2D空間局部性的優(yōu)化策略，所以thread要獲取2D數(shù)據(jù)就可以使用texture Memory來(lái)達(dá)到很高的性能。從讀取性能的角度跟Constant Memory類似。

Host Memory

主機(jī)端存儲(chǔ)器主要是內(nèi)存可以分為兩類：可分頁(yè)內(nèi)存（Pageable）和頁(yè)面 （Page-Locked 或 Pinned）內(nèi)存。

可分頁(yè)內(nèi)存通過(guò)操作系統(tǒng) API(malloc/free) 分配存儲(chǔ)器空間，該內(nèi)存是可以換頁(yè)的，即內(nèi)存頁(yè)可以被置換到磁盤(pán)中?？煞猪?yè)內(nèi)存是不可用使用DMA（Direct Memory Acess)來(lái)進(jìn)行訪問(wèn)的，普通的C程序使用的內(nèi)存就是這個(gè)內(nèi)存

例子

下面例子講解如何使用統(tǒng)一內(nèi)存：

__device__ __managed__ int x[2];
__device__ __managed__ int y;
__global__ void kernel(){x[1] = x[0] + y;
}int main(){x[0] = 3;y = 5;kernel<<< 1, 1 >>>();cudaDeviceSynchronize();printf("result=%d\n", x[1]);return 0;
}