CUDA ---- Constant Memory

標籤：

CONSTANT  MEMORY

constant Memory對於device來說唯讀但是對於host是可讀可寫。constant Memory和global Memory一樣都位於DRAM，並且有一個獨立的on-chip cache，比直接從constant Memory讀取要快得多。每個SM上constant Memory大小限制為64KB。

constant Memory的擷取方式不同於其它的GPU記憶體，對於constant Memory來說，最佳擷取方式是warp中的32個thread擷取constant Memory中的同一個地址。如果擷取的地址不同的話，只能串列的服務這些擷取請求了。

constant Memory使用__constant__限定符修飾變數。

constantMemory的生命週期伴隨整個應用程式，並且可以被同一個grid中的thread和host中調用的API擷取。因為constant Memory對device來說是可讀的，所以只能在host初始化，使用下面的API：

cudaError_t cudaMemcpyToSymbol(const void *symbol, const void * src, size_t count, size_t offset, cudaMemcpyKind kind)

Implementing a 1D Stencil with Constant Memory

實現一個1維Stencil（數值分析領域的東，卷積神經網路處理映像的時候那個stencil），簡單說就是計算一個多項式，係數放到constant Memory中，即y=f(x)這種東西，輸入是九個點，如下：

{x − 4h, x − 3h, x − 2h, x − h, x, x + h, x + 2h, x + 3h, x + 4h}

在記憶體中的過程如下：

 

公式如下：

 

那麼要放到constant Memory中的便是其中的c0、c1、c2 ……

因為每個thread使用九個點來計算一個點，所以可以使用shared memory來降低延遲。

__shared__ float smem[BDIM + 2 * RADIUS];

RADIUS定義了x兩邊點的個數，對於本例，RADIUS就是4。如所示，每個block需要RADIUS=4個halo（暈）左右邊界：

 

#pragma unroll用來告訴編譯器，自動延伸迴圈。

__global__ void stencil_1d(float *in, float *out) {// shared memory__shared__ float smem[BDIM + 2*RADIUS];// index to global memoryint idx = threadIdx.x + blockIdx.x * blockDim.x;// index to shared memory for stencil calculatioinint sidx = threadIdx.x + RADIUS;// Read data from global memory into shared memorysmem[sidx] = in[idx];// read halo part to shared memoryif (threadIdx.x < RADIUS) {smem[sidx - RADIUS] = in[idx - RADIUS];smem[sidx + BDIM] = in[idx + BDIM];}// Synchronize (ensure all the data is available)__syncthreads();// Apply the stencilfloat tmp = 0.0f;#pragma unrollfor (int i = 1; i <= RADIUS; i++) {tmp += coef[i] * (smem[sidx+i] - smem[sidx-i]);}// Store the resultout[idx] = tmp;}

View CodeComparing with the Read-only Cache

Kepler系列的GPU允許使用texture pipeline作為一個global Memory唯讀緩衝。因為這是一個獨立的使用單獨的頻寬的唯讀緩衝，所以對頻寬節流設定的kernel效能有很大的提升。

Kepler的每個SM有48KB大小的唯讀緩衝，一般來說，在讀地址比較分散的情況下，這個唯讀緩衝比L1表現要好，但是在讀同一個地址的時候，一般不適用這個唯讀緩衝，唯讀緩衝的讀取粒度為32位元。

有兩種方式來使用唯讀緩衝：

• 使用__ldg限定
• 指定特定global Memory稱為唯讀緩衝

下面程式碼片段對於第一種情況：

__global__ void kernel(float* output, float* input) {    ...    output[idx] += __ldg(&input[idx]);    ...}

下面代碼對應第二種情況，使用__restrict__來指定該資料的要從唯讀緩衝中擷取：

void kernel(float* output, const float* __restrict__ input) {    ...    output[idx] += input[idx];}

一般使用__ldg是更好的選擇。通過constant緩衝儲存的資料必須相對較小而且必須擷取同一個地址以便擷取最佳效能，相反，唯讀緩衝則可以存放較大的資料，且不必地址一致。

下面的代碼是之前stencil的翻版，使用過了唯讀緩衝來儲存係數，二者唯一的不同就是函數的聲明：

__global__ void stencil_1d_read_only (float* in, float* out, const float *__restrict__ dcoef) {// shared memory__shared__ float smem[BDIM + 2*RADIUS];// index to global memoryint idx = threadIdx.x + blockIdx.x * blockDim.x;// index to shared memory for stencil calculatioinint sidx = threadIdx.x + RADIUS;// Read data from global memory into shared memorysmem[sidx] = in[idx];// read halo part to shared memoryif (threadIdx.x < RADIUS) {smem[sidx - RADIUS] = in[idx - RADIUS];smem[sidx + BDIM] = in[idx + BDIM];}// Synchronize (ensure all the data is available)__syncthreads();// Apply the stencilfloat tmp = 0.0f;#pragma unrollfor (int i=1; i<=RADIUS; i++) {tmp += dcoef[i]*(smem[sidx+i]-smem[sidx-i]);}// Store the resultout[idx] = tmp;}

View Code

由於係數原本是存放在global Memory中的，然後讀進緩衝，所以在調用kernel之前，我們必須分配和初始化global Memory來儲存係數，代碼如下：

const float h_coef[] = {a0, a1, a2, a3, a4};cudaMalloc((float**)&d_coef, (RADIUS + 1) * sizeof(float));cudaMemcpy(d_coef, h_coef, (RADIUS + 1) * sizeof(float), cudaMemcpyHostToDevice);

下面是運行在TeslaK40上的結果，從中可知，使用唯讀緩衝效能較差。

Tesla K40c array size: 16777216 (grid, block) 524288,323.4517ms stencil_1d(float*, float*)3.6816ms stencil_1d_read_only(float*, float*, float const *)

總的來說，constant緩衝和唯讀緩衝對於device來說，都是唯讀。二者都有大小限制，前者每個SM只能有64KB，後者則是48KB。對於讀同一個地址，constant緩衝表現好，唯讀緩衝則對地址較分散的情況表現好。

The Warp Shuffle Instruction

之前我們有介紹shared Memory對於提高效能的好處，在CC3.0以上，支援了shuffle指令，允許thread直接讀其他thread的寄存器值，只要兩個thread在 同一個warp中，這種比通過shared Memory進行thread間的通訊效果更好，latency更低，同時也不消耗額外的記憶體資源來執行資料交換。

這裡介紹warp中的一個概念lane，一個lane就是一個warp中的一個thread，每個lane在同一個warp中由lane索引唯一確定，因此其範圍為[0,31]。在一個一維的block中，可以通過下面兩個公式計算索引：

laneID = threadIdx.x % 32

warpID = threadIdx.x / 32

例如，在同一個block中的thread1和33擁有相同的lane索引1。

Variants of the Warp Shuffle Instruction

有兩種設定shuffle的指令：一種針對整型變數，另一種針對浮點型變數。每種設定都包含四種shuffle指令變數。為了交換整型變數，使用過如下函數：

int __shfl(int var, int srcLane, int width=warpSize);

該函數的作用是將var的值返回給同一個warp中lane索引為srcLane的thread。選擇性參數width可以設定為2的n次冪，n屬於[1,5]。

eg：如果shuffle指令如下：

int y = shfl(x, 3, 16);

則，thread0到thread15會擷取thread3的資料x，thread16到thread31會從thread19擷取資料x。

當傳送到shfl的lane索引相同時，該指令會執行一次廣播操作，如下所示：

 

另一種使用shuffle的形式如下：

int __shfl_up(int var, unsigned int delta, int width=warpSize)

該函數通過使用調用方的thread的lane索引減去delta來計算源thread的lane索引。這樣源thread的相應資料就會返回給調用方，這樣，warp中最開始delta個的thread不會改變，如下所示：

 

第三種shuffle指令形式如下：

int __shfl_down(int var, unsigned int delta, int width=warpSize)

該格式是相對__shfl_down來說的，具體形式如所示：

 

最後一種shuffle指令格式如下：

int __shfl_xor(int var, int laneMask, int width=warpSize)

這次不是加減操作，而是同laneMask做抑或操作，具體形式如所示：

 

所有這些提及的shuffle函數也都支援單精確度浮點值，只需要將int換成float就行，除此外，和整型的使用方法完全一樣。

 

 參考書：《professional cuda c programming》

NVIDIA CUDA板塊：https://developer.nvidia.com/cuda-zone

CUDA線上文檔：http://docs.nvidia.com/cuda/index.html#

CUDA ---- Constant Memory