CONSTANT MEMORY

constant Memory對於device來說只讀但是對於host是可讀可寫。constant Memory和global Memory一樣都位於DRAM，並且有一個獨立的on-chip cache，比直接從constant Memory讀取要快得多。每個SM上constant Memory cache大小限制為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 memory
int idx = threadIdx.x + blockIdx.x * blockDim.x;
// index to shared memory for stencil calculatioin
int sidx = threadIdx.x + RADIUS;
// Read data from global memory into shared memory
smem[sidx] = in[idx];
// read halo part to shared memory
if (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 stencil
float tmp = 0.0f;
#pragma unroll
for (int i = 1; i <= RADIUS; i++) {
tmp += coef[i] * (smem[sidx+i] - smem[sidx-i]);
}
// Store the result
out[idx] = tmp;
}

View Code

Comparing 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 memory
int idx = threadIdx.x + blockIdx.x * blockDim.x;
// index to shared memory for stencil calculatioin
int sidx = threadIdx.x + RADIUS;
// Read data from global memory into shared memory
smem[sidx] = in[idx];
// read halo part to shared memory
if (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 stencil
float tmp = 0.0f;
#pragma unroll
for (int i=1; i<=RADIUS; i++) {
tmp += dcoef[i]*(smem[sidx+i]-smem[sidx-i]);
}
// Store the result
out[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,32
3.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索引相同時，該指令會執行一次廣播操作，如下所示：