第5章 共享内存和常量内存

通过安排全局内存访问模式，学会了如何实现良好的性能并且避免浪费了事务内存性能。但在跨全局内存的非合并内存访问，或者未对齐的内存访问，仍然会导致带宽利用率不会达到最佳标准。于是使用共享内存来提高全局内存合并访问是有可能的。

5.1 CUDA共享内存概述

GPU上有两种类型的内存：板载内存和片上内存；

共享内存常见用途;
        1. 块内线程通信的通道；
        2. 用于全局内存数据的可编程管理的缓存，优化全局内存访问模式；

5.1.1 共享内存

共享内存是GPU的一个关键部件。物理上，每个SM都有一个小的低延迟内存池，这个内存池被当前正在该SM上执行的线程块中的所有线程所共享。

共享内存相较于全局内存而言 ，延迟要低20~30倍，而带宽高其约10倍。

共享内存被SM中的所有常驻线程块划分。一个核函数使用的共享内存越多，处于并发活跃状态的线程块就越少。

由于在CUDA中允许手动管理共享内存，所以通过在数据布局上提供更多的细粒度控制和改善片上数据的移动。

5.1.2 共享内存分配

CUDA支持一维、二维和三维共享内存数组的声明。共享内存变量用下列修饰符进行声明：__shared__。如果共享内存的大小在编译时是未知的，那么可以用extern关键字声明一个未知大小的数组。但是只能动态声明一维数组。extern __shared__ int tile[]

5.1.3 共享内存存储体和访问模式

优化内存性能是要度量的两个关键属性：延迟和带宽。

5.1.3.1 内存存储体

为获得高内存带宽，共享内存被分为32个同样大小的内存模型，它们被称为存储体，它们可以被同时访问。有32个存储体是因为在一个线程束中有32个线程。共享内存是一维地址空间。

根据GPU的计算能力，共享内存的地址在不同模式下会映射到不同的存储体中。

5.1.3.2 存储体冲突

在共享内存中当多个地址请求落在相同的存储体中时，就会发生存储体冲突，这会导致冲突访存被重复执行。

当线程发出共享内存请求时，以下3种典型的模式：
        1. 并行访问：多个地址请求访问多个存储体；
        2. 串行访问：多个地址请求访问同一存储体；
        3. 广播访问：单一地址请求读取单一存储体；

并行访问模式意味着，如果不是所有地址，那么至少有一些地址可以在一个内存事务中被服务。最佳情况是，当每个地址都位于一个单独的存储体中，执行无冲突的共享内存访问。

串行访问是最坏的模式，当多个地址属于同一存储体时，必须以串行的方式进行请求。

在广播访问的情况下，线程束所有线程都读取同一存储体中相同的地址。若一个内存事务被执行，那么被访问的字就会被广播到所有请求的线程中。

如果几个线程访问同一存储体时，会产生两种可能的行为：
        1. 如果线程访问同一存储体的相同地址，广播访问无冲突；
        2. 如果线程访问同一存储体的不同地址，会发生存储体冲突。

5.1.3.3 访问模式

共享内存存储体的宽度规定了共享内存地址与共享内存存储体的对应关系。内存存储体的宽度随设备计算能力的不同而变化。有两种不同的存储体宽度：
        1. 计算能力2.x的设备中为4字节；
        2. 计算能力3.x的设备中为8字节；

对于Fermi(2.x)设备而言，从共享内存地址到存储体索引的映射可以按如下公式计算：

        存储体索引 = (字节地址 / 4字节/存储体) % 32存储体

对来自相同线程束中的两个线程访问相同的地址时，不会发生存储体冲突。在这种情况下，对于读访问，这个字被广播到请求的线程中；对于写访问，这个字只能由其中一个线程写入，执行这个写入操作的线程是不确定的。

对于Kepler设备，共享内存有32个存储体，它们有两种地址模式：
        64位和32位模式；从共享内存地址到存储体索引的映射可以按如下公式计算：
        存储体索引 = (字节地址 / 8字节/存储体) % 32存储体

如果两个线程访问64位字中的同一个字节，从线程束发出的共享内存请求就不会产生存储体冲突，因为满足这两个请求只需要一个64位的读操作。当访问64位存储体的不同位置时，发生存储体冲突。线程束内有n线程访问同一存储体不同位置时，称n向存储体冲突。

5.1.3.4 内存填充

内存填充是避免存储体冲突的一种办法。

假设有5个共享内存存储体。如果所有线程访问bank0的不同地址，那么会发生一个5向存储体冲突。解决这种存储体冲突的一个办法是在每N个元素之后添加一个字，这里的N是指存储体数量。

填充的的数据不能用于存储数据，其唯一的作用是移动数据元素，以便将原来属于同一存储体中的数据分散到不同的存储体中。虽然导致线程块可用的总的共享内存的数量将减少，以及需要重新计算数组索引以确保能访问到正确的数据元素。

在不同架构设备上填充共享内存时，必须要小心。

5.1.3.5 访问模式配置

Kepler设备支持4字节和8字节的共享内存访问模式。默认是4字节。可采用CUDA运行时API函数查询访问模式：cudaDeviceGetSharedMemConfig()；

在可配置的共享内存存储体的设备上，可以使用一下功能设置一个新的存储体大小：

cudaDeviceSetSharedMemConfig();

在不同的核函数启动之间更改共享内存配置可能需要一个隐式的设备同步点。更改共享内存存储体的大小不会增加共享内存的使用量，也不会影响核函数的占用率，但它对性能可能会有重要影响。一个大的存储体可能为共享内存访问产生更高的带宽，但是可能会导致更多的存储体冲突，这取决于应用程序中共享内存的访问模式。

5.1.4 配置共享内存

若每个SM都有64kb的片上内存。共享内存和一级缓存共享该硬件资源。CUDA为配置一级缓存和共享内存的大小提供了两种方法：按设备进行配置和按核函数进行配置；

使用cudaDeviceSetCacheConfig()可以为在设备上启动的核函数配置一级缓存和共享内存的大小。该函数的参数选择取决核函数使用多少的共享内存：
        1. 当核函数使用较多的共享内存时，倾向于更多的共享内存；
        2. 当核函数使用较多的寄存器时，倾向于更多的一级缓存；(kepler设备一级缓存可用于寄存器溢出)

指定-Xptxas -v选项给nvcc，可以指导核函数使用了多少寄存器。当内核使用的寄存器数量超过了硬件限制所允许的数量时，应该为寄存器溢出配置出一个更大的一级缓存。

每个核函数的配置可以覆盖设备范围的设置，也可使用运行时函数cudaFuncSetCacheConfig()配置。

启动一个不同的优先级的内核比启动有最近优先级设置的内核更可能导致隐式设备同步。对于每个内核，只需调用一次cudaFuncSetCacheConfig函数。每个内核函数启动时，片上内存中的配置不需要重新设定。

共享内存是通过32个存储体进行访问的，而一级缓存则是通过缓存进行访问的。使用共享内存，对存储内容和存放位置有完全的控制权，而使用一级缓存，数据删除工作是由硬件完成的。

GPU使用不同的启发式算法删除数据。数据删除在GPU上可能会发生得频繁而且更不可预知。使用GPU共享内存不仅可以显示管理数据而且还可以利用SM的局部性。

5.1.5 同步

共享内存可以同时被线程块中的多个线程访问。当不同步的多个线程修改同一个共享内存地址时，将导致线程内的冲突。

同步的两个基本方法如下所示：障碍和内存栅栏；

在障碍中，所有调用的线程等待其余调用的线程到达障碍点。

在内存栅栏中，所有调用的线程必须等到全部内存的修改对其余调用线程可见时才能继续执行。

5.1.5.1 弱排序内存模型

现代内存架构有一个宽松的内存模型。内存访问不一定按照它们在程序中出现的顺序执行。

一个线程的写入顺序对其他线程可见时，它可能和写操作被执行的实际顺序不一致。

如果指令之间是相互独立的，线程从不同内存中读取数据的顺序和读指令在程序中出现的顺序不一定相同。

5.1.5.2 显式障碍

在CUDA中，障碍只能在同一个线程块的线程间执行。调用下面的函数来指定一个障碍点：

__syncthreads() 要求块中线程必须等待直到所有线程都到达该点。__syncthreads还确保在障碍点之前，被这些线程访问的所有全局的共享内存对同一个块中所有线程都可见。

__syncthread()用于协调同一块中线程间的通信。当块中的某些线程访问共享内存或全局内存中的同一地址时，会有潜在的写后读、读后写、写后写，这将导致在那些内存位置产生未定义的应用程序行为和未定义的状态。可以通过利用冲突访问间的同步线程来避免这种情况。

如果一个条件能保证对整个线程块进行同等评估，则它是调用__syncthread()的唯一有效条件。

如果不允许跨线程块同步、线程块可能会以任何顺序、并行、串行的顺序在任何SM上执行。线程块执行的独立性质使得CUDA编程在任意数量的核心中都是可扩展的。

如果一个CUDA核函数要求跨线程块全局同步，那么通过同步点分割核函数并执行多个内核启动可能会到达预期的效果。因为每个连续的内核启动必须等待之前的内核启动完成，所以这会产生一个隐式的全局同障碍。

5.1.5.3 内存栅栏

内存栅栏的功能可确保栅栏前的任何内存写操作对栅栏后的其他线程都是可见的。根据所需的范围，有3种内存栅栏：块，网格，系统

下列内置函数可以在线程块内创建内存栅栏：__threadfence_block(); 它保证了栅栏前被调用的线程产生对共享内存和全局内存的所有写操作对栅栏后同一块中的其他线程都是可见的。回想一下，内存栅栏不执行任何线程同步，所以对于一个块中的所有线程来说，没有必要执行这个指令。

使用__threadfence()挂起调用线程，直到全局内存中所有写操作对相同网格内的所有线程都是可见的。

使用__threadfence_system()可以跨系统(主机和设备)设置内存栅栏。挂起调用的线程，以确保该线程对全局内存、锁页内存和其他设备内存中的所有写操作对全部设备中的线程和主机线程都是可见的。

5.1.5.4 Volatile修饰符

在全局或共享内存使用volatile修饰符声明一个变量，可以防止编译器优化，编译器优化可能会将数据暂时缓存在寄存器或者本地内存中。当使用volatile修饰符，编译器假定任何其他线程在任何时间都可以更改或使用该变量的值。因此，这个变量的任何引用都会直接被编译到全局内存读指令或全局内存写指令中，它们会忽略缓存。

5.2 共享内存的数据布局

其中包含下列主题：

5.2.1 方形共享内存

使用下面语句静态声明一个二维共享内存变量：__shared__ int tile[N][N];

这就需要注意线程与共享内存存储体的映射关系。临近线程在最内层数组维度上访问相邻的阵列单元。

5.2.1.1 行主序

__shared__ int tile[BDIMY][BDIMX]：
        1. 将全局线程索引按行主序写入到一个二维共享内存数组中；
        2. 从共享内存中按行主序读取这些值并将它们存储到全局内存中；

将全局线程索引按行主序顺序写入共享内存块：

        tile[threadIdx.y][threadIdx.x] = idx;

一旦到达同步点(使用syncthreads函数)，所有线程必须将存储的数据送到共享内存存储体。

__global__ void setRowReadRow(int* out)
{
    __shared__ int tile[BDIMY][BDIMX];

    unsigned int idx  = threadIdx.y * blockDim.x + threadIdx.x;

    tile[threadIdx.y][threadIdx.x] = odx;

    __syncthreads();

    out[idx] = tile[threadIdx.y][threadIdx.x];
}

可以使用以下nvprof指标已检查存储体冲突：
        1. shared_load_transactions_per_request;
        2. shared_store_transactions_per_request;      

5.2.1.3 动态共享内存

动态共享内存必须声明为一个未定大小的一维数组。extern __shared__ int tile[];

5.2.1.4 填充静态声明的共享内存

__shared__　int tile[BDIMY][BDIMX + 1];

__global__ void setRowReadColPad(int* out)
{
    __shared__ int tile[BDIMY][BDIMX + IPAD];

    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    tile[threadIdx.y][threadIdx.x] = idx;

    __syncthreads();

    out[idx] = tile[threadIdx.y][threadIdx.x];
}

对于Fermi设备，需要增加一列来解决存储体冲突；对于Kepler设备，并非总是如此。因为在Kepler设备中，每行需要填充的数据元素数量取决于二维共享内存的大小。

5.2.1.5 填充动态声明的共享内存

填充动态声明的共享内存数组更加复杂。当执行从二维线程索引到一维内存索引的转换时，对于每一行必须跳过一个填充的内存空间：

__global__ void setRowReadColDynPad(int* out)
{
    extern __shared__ int tile[];

    unsigned int row_idx = threadIdx.y * (blockDim.x + IPAD) + threadIdx.x;
    unsigned int col_idx = threadIdx.x * (blockDim.x + IPAD) + threadIdx.y;

    unsigned int g_idx = threadIdx.y * blockDim.x + threadIdx.x;
    tile[row_idx] = g_idx;

    __syncthreads();

    out[g_idx] = tile[col_idx];
}

在核函数中用于存储数据的全局内存大于填充的共享内存。所以需要三个索引：一个行主序的写入共享内存，一个列主序读取共享内存，一个用于未填充的全局内存的合并访问。

5.2.1.6 方形共享内存内核的比较

使用共享内存时：
        1. 使用填充的内核可提高性能，因为它减少了存储体冲突；
        2. 带有动态声明的共享内存的内核增加了少量的消耗；

5.2.2 矩形共享内存

__shared__ int tile[Row][Col];

当执行一个转置操作时，不能像在方形共享内存一样，只是通过简单地转换来引用矩形数组的线程坐标。当使用矩形共享内存时，这样会导致内存访问冲突。

5.2.2.2 行主序写操作和列主序读操作

使用共享内存执行矩阵转置，通过最大化低延迟的加载和存储来提高性能。

此时内核有3个内存操作：
        1. 写入每个线程束的共享内存行，以避免存储体冲突；
        2. 读取每个线程束中的共享内存列，以完成矩阵转置；
        3. 使用合并访问写入每个线程束的全局内存行；

unsigned int idx = threadIdx.y * blockDim.x + THREADiDX.X 保证全局内存访问是合并的。

__global__ void setRowReadCol(int* out)
{
    __shared__ int tile[BDIMY][BDIMX];

    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    unsigned int irow = idx / blockDim.y;
    unsigned int icol = idx % blockDim.y;

    tile[threadIdx.y][threadIdx.x] = idx;

    __syncthreads();

    out[idx] = tile[icol][irow];
}

5.3  减少了全局内存访问

使用共享内存的主要原因之一是要缓存片上的数据，从而减少核函数中全局内存访问的次数。

第三章使用全局内存的并行归约核函数解释了以下几个问题;
        1. 如何重新安排数据访问模式以避免线程束分化；
        2. 如何展开循环以保证足够的操作使指令和内存带宽饱和；

5.3.1 使用共享内存的并行归约

__global__ void reduceSmem(int* g_data, int* g_odata, unsigned int n)
{
    __shared__ int smem[DIM];

    unsigned int tid = threadIdx.x;

    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (idx >= n) return;

    int* idata = g_idata + blockDim.x * blockIdx.x;

    smem[tid] = idata[tid];
    __syncthreads();

    if (blockDim.x >= 1024 && tid < 512) smem[tid] += smem[tid + 512]; 
    __syncthreads();
    if (blockDim.x >= 512 && tid < 256) smem[tid] += smem[tid + 256];
    __syncthreads();
    if (blockDim.x >= 256 && tid < 128) smem[tid] += smem[tid + 128];
    __syncthreads();
    if (blockDim.x >= 128 && tid < 64) smem[tid] += smem[tid + 64];
    __syncthreads();

    if (tid < 32)
    {
        volatile int* vsmem = smem;
        vsmem[tid] += vsmem[tid + 32];
        vsmem[tid] += vsmem[tid + 16];
        vsmem[tid] += vsmem[tid + 8];
        vsmem[tid] += vsmem[tid + 4];
        vsmem[tid] += vsmem[tid + 2];
        vsmem[tid] += vsmem[tid + 1];
    }

    if (tid == 0) g_odata[blockIdx.x] = smem[0];

}

5.3.2 使用展开的并行归约

当内核展开了4个线程块时，以下的优势是可预期的;
        1. 通过在每个线程中提供更多的并行IO操作，增加全局内存吞吐量；
        2. 全局内存事务减少了1/4；
        3. 整体内核性能提升了；

代码如下：

__global__ void reduceSmemUnroll(int* g_data, int* g_odata, unsigned int n)
{
    __shared__ int smem[DIM];

    unsigned int tid = threadIdx.x;

    unsigned int idx = blockIdx.x * blockDim.x * 4 + threadIdx.x;

    int temSum = 0;

    if (idx + 3 * blockDim.x * blockIdx.x <= n)
    {
        int a1 = g_idata[idx];
        int a2 = g_idata[idx + 1 * blockDim.x * blockIdx.x];
        int a3 = g_idata[idx + 2 * blockDim.x * blockIdx.x];
        int a4 = g_idata[idx + 3 * blockDim.x * blockIdx.x];
        temSum = a1 + a2 + a3 + a4;        
    }

    smem[tid] = temSum;
    __syncthreads();

    if (blockDim.x >= 1024 && tid < 512) smem[tid] += smem[tid + 512]; 
    __syncthreads();
    if (blockDim.x >= 512 && tid < 256) smem[tid] += smem[tid + 256];
    __syncthreads();
    if (blockDim.x >= 256 && tid < 128) smem[tid] += smem[tid + 128];
    __syncthreads();
    if (blockDim.x >= 128 && tid < 64) smem[tid] += smem[tid + 64];
    __syncthreads();

    if (tid < 32)
    {
        volatile int* vsmem = smem;
        vsmem[tid] += vsmem[tid + 32];
        vsmem[tid] += vsmem[tid + 16];
        vsmem[tid] += vsmem[tid + 8];
        vsmem[tid] += vsmem[tid + 4];
        vsmem[tid] += vsmem[tid + 2];
        vsmem[tid] += vsmem[tid + 1];
    }

    if (tid == 0) g_odata[blockIdx.x] = smem[0];

}

要使每个线程处理4个数据元素，第一步是基于每个线程的线程块和线程索引，重新计算全局输入数据的偏移。每个线程块的处理起点现在被偏移为就好像是线程块的四倍。此时核函数的网格维度也必须减少到每个线程执行量的1/4.

最后检查全局内存吞吐量。加载吞吐量增加了2.57倍，而存储吞吐量下降了1.56倍。加载吞吐量的增加归因于大量的同时加载请求。吞吐量的下降是因为较少的存储请求使总线达到了饱和。

5.3.3 使用动态共享内存的并行归约

extern __shared__ int smem[];

发现用动态共享内存实现的核函数和利用静态分配的共享内存实现的核函数之间没有显著差异。

5.3.4 有效带宽

由于归约核函数是受内存带宽约束的。

有效带宽是在核函数的完整执行时间内I/O的数量(以字节为单位)。可以表示为：

        有效带宽 = （读字节 + 写字节）/ （运行时间 * 10^9）GB/s

5.4 合并的全局内存访问

使用共享内存也可以帮助避免对未合并的全局内存的访问。

第4章已经表明，交叉访问是全局内存中最糟糕的访问模式，因为它浪费了总线带宽。

5.4.1 基准转置内核

__global__ void naiveGmem(float* out, float* in, const int nx, const int ny)
{
    unsigned int ix = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int iy = blockDim.y * blockDim.y + threadIdx.y;

    if (ix < nx && iy < ny)
        out[ix * ny + iy] = in[iy * nx + ix];
}

5.4.2 使用共享内存的矩阵转置

为了避免交叉全局内存访问，可以使用二维共享内存来缓存原始矩阵的数据。

虽然朴素实现会导致共享内存存储体冲突，但这个结果将比非合并的全局内存访问好得多。

__global__ void transposeSmem(float* out, float* in, const int nx, const int ny)
{
    __shared__ float tile[BDIMY][BDIMX];

    unsigned int ix, iy, ti, to;

    unsigned int ix = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int iy = blockDim.y * blockDim.y + threadIdx.y;

    ti = iy * nx + ix;

    unsigned int bidx, irow, icol;
    bidx = threadIdx.y * blockDim.x + threadIdx.x;
    irow = bidx / blockDim.y;
    icol = bidx % blockDim.y;

    ix = blockIdx.y * blockDim.y + icol;
    iy = blockIdx.x * blockDim.x + irow;

    to = iy * ny + ix;

    if (ix < nx && iy < ny)
    {
        tile[threadIdx.y][threadIdx.x] = in[ti];

        __syncthreads();

        out[to] = tile[icol][irow];
    }
        
}

以上kernel可被分解为以下几个步骤：
        1. 线程束执行合并读取一行，该行存储在全局内存中的原始矩阵块中；
        2. 该线程束按行读取将数据写入共享内存；
        3. 因为线程块的读和写是同步的，所以会有一个填满全局内存数据的二维共享内存；
        4. 该线程束从二维共享内存数组中读取一列。再加上共享内存没有被填充，所以会发生存储体冲突；
        5. 该线程束执行数据的合并写入操作，将其写入到全局内存的转置矩阵中的某行。

对于每一个线程，若想要从全局内存和共享内存中取得正确的数据，都必须计算多个索引。原始矩阵全局内存读，共享内存写。转置后的共享内存读，转置矩阵的写。

5.4.3 使用填充共享内存的矩阵转置

通过给二维共享内存数组tile中的每一行添加列填充，可以将原矩阵相同列中的数据元素均匀地划分到共享内存存储体中。需要填充的列数取决于设备的计算能力和线程块的大小。

__shared__ float tile[BDIMY][BDIMX + IPAD];

5.4.4 使用展开的矩阵转置

展开的目标是通过创造更多的同时加载和存储以提高设备内存带宽利用率。

__global__ void transposeSmemUnrollPad(float* out, float* in, const int nx, const int iy)
{
    __shared__ float tile[BDIMY * (BDIMX * 2 + IPAD)];

    unsigned int ix = blockIdx.x * blockDim.x * 2 + threadIdx.x;
    unsigned int iy = blockIdx.y * blockDim.y + threadIdx.y;

    unsigned int ti = ix + iy * nx;

    unsigned int bidx = threadIdx.y * blockDim.x + threadIdx.x;
    unsigned int icol = bidx % blockDim.y;
    unsigned int irow = bidx / blockDim.y;

    unsigned int ix2 = icol + blockDim.y * blockIdx.y;
    unsigned int iy2 = irow + blockDim.x * blockDim.x * 2;

    unsigned int to = ix + iy * ny;

    if (ix + blockDim.x < nx && iy < ny)
    {
        unsigned int row_idx = threadIdx.y * (blockDim.x * 2 + IPAD) + threadIdx.x;

        tile[row_idx] = in[ti];
        tile[row_idx + BDIMX] = in[ti + BDIMX];

        __syncthreads();

        unsigned int col_idx = icol * (blockDim.x * 2 + IPAD) + irow;
        out[to] = tile[col_idx];
        out[to + ny * BDIMX] = tile[col_idx + BDIMX];
        
    }

}

因为填充的内存不是用来存储数据的，所以计算索引时必须跳过填充列。

5.4.5 增大并行性

一个简单而有效的优化技术是调整线程块的维度，找出最佳的执行配置。

5.5 常量内存

常量内存是一种专用的内存，它用于只读数据和统一访问线程束中的线程的数据。

常量内存读内核代码而言是只读的，但它对主机而言是可读又是可写的。

每个SM常量内存缓存大小为64KB。

相比于其他内存类型，常量内存有一个不同的最优访问模式。在常量内存中，如果线程束中的所有线程都访问相同的位置，那么这个访问模式就是最优的。如果线程束中的线程访问不同的地址，则访问需要串行。

在全局作用域中必须用一下修饰符声明常量变量：__constant__；

常量内存变量的生存期与应用程序的生存期相同，对网格内的所有线程都是可访问的，并且通过运行时函数对主机可访问。

当使用CUDA独立编译能力时，常量内存变量跨多个源文件是可见的。因为设备只能读取常量内存，所以常量内存中的值必须使用以下运行时函数进行初始化。cudaMemcpyToSymbol()将src指向的数据复制到设备上由symbol指定的常量内存中。

5.5.1 使用常量内存实现一维模板

为实现一维模板计算，由于每个线程需要9个点来计算一个点，所以要使用共享内存来缓存数据，从而减少对全局内存的冗余访问。

__shared__ float smem[BDIM + 2 * RADIUS]; RADIUS = 4;

__constant__ float coef[RADIUS + 1];
void setup_coef_constant(void)
{
    const float h_coef[] = {a1, a1, a2, a3, a4};
    cudaMemcpyToSymbol(coef, h_coef, (RADIUS + 1) * sizeof(float));
}

__global__ void stencil_1d(float* in, float* out)
{
    __shared__ float smem[BDIM + 2 * RADIUS];   

    int idx = threadIdx.x + blockDim.x * blockIdx.x;

    int side = threadIdx.x + RADIUS;

    smem[side] = in[idx];

    if (threadIdx.x < RADIUS)
    {
        smem[side - RADIUS] = in[idx - RADIUS];
        smem[side + BDIM] = in[idx + BDIM];
    }

    __syncthreads();

    float temp = 0.0f;
    #pragma unroll
    for (int i = 1; i <= RADIUS; i++)
        temp += coef[i] * (smem[side + i] - side[side - x]);
    
    out[idx] = temp;
}

5.5.2 与只读缓存相比

每个Kepler SM都有48KB的只读缓存。一般来说，只读缓存分散读取方面比一级缓存更好，当线程束中的线程都读取相同地址时，不应该使用只读缓存。只读缓存的粒度为32字节。

当通过只读缓存访问全局内存时，需要向编译器指出在内核的持续时间里数据是只读的。有两种方法可以实现这一点:
        1. 使用内部函数__ldg：内部函数__ldg用于代替标准指针解引用__ldg(&input[idx])。
        2. 全局内存的限定指针；const __restrict__ 表明它们应该通过只读缓存被访问。

在只读缓存机制需要更多显式控制的情况下，或者在代码非常复杂以至于编译器无法检测到只读缓存的使用是否安全的情况下，内部函数__ldg()是一个更好的选择。

只读缓存是独立的，而且区别于常量缓存。通过常量缓存加载的数据必须相对较小，而且访问必须一致以获得良好的性能，而通过只读缓存加载的数据可以比较大，而且能够在一个非统一的模式下进行访问。

常量缓存和只读缓存：
        1. 在设备上常量缓存和只读缓存都是只读的。
        2. 每个SM资源都有限：常量缓存是64KB，而只读缓存是48KB。
        3. 常量缓存在统一读取中可更好地执行(统一读取是线程束中的每一线程都访问相同的地址)。
        4. 只读缓存更适合分散读取。

5.6 线程束洗牌指令

之前介绍了如何使用共享内存执行线程块中线程间低延迟数据的交换。从kepler架构开始，洗牌指令(shuffle instruction)作为一种机制被加入其中，只要两个线程在相同的线程束中，那么就允许这两个线程直接读取另一个线程的寄存器。

洗牌指令使得线程束中的线程彼此之间交换数据，而不是通过共享内存或全局内存来进行。洗牌指令比访问共享内存指令有更低的延迟，并且该指令在执行数据交换时不消耗额外的内存。因此，洗牌指令为应用程序快速交换线程束内的数据提供了一个有吸引力的方法。

一个束内线程指的是线程束内的单一线程。[lane index]---[0, 31]

束内线程索引没有内置变量，因为线程索引有内置变量。在一维线程块中，对于一个给定线程的束内线程索引和线程束索引可按照公式计算; laneID = threadIdx.x % 32和warpID = threadIdx.x / 32;

5.6.1 线程束洗牌指令的不同形式

有两组洗牌指令，一组用于整型变量，一组用于浮点型变量。每组有4种形式的洗牌指令。在线程束内交换整型变量，其基本函数标记如下;

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

        __shfl(val, 2):一条从束内线程2到线程束中所有线程的广播；

这个函数能使得线程束中的每个线程都可以直接从一个特定的线程中获取某个值。线程束内所有活跃的线程都同时产生此操作，这将导致每个线程中有4个字节数据的移动。

变量width可以被设置为2~32之间任何2的指数。当width设置为32时，洗牌指令跨整个线程束来执行，并且srcLane指定源线程的束内线程索引。然而，设置width允许将线程束细分为段，使每段包含有width个线程，并且在每个段上执行独立的洗牌操作。对于不是32的其他width值，线程的束内线程ID和其在洗牌操作中的ID不一定相同时。在这种情况下，一维线程块中的线程洗牌ID可以按下了公式进行计算：suffleID = threadIdx.x % width；

当传递给shuf的束内线程索引与线程束中所有线程的值相同时，指令从特定的束内线程到线程束中所有线程都执行线程束广播操作。

洗牌操作的另一种形式是从与调用线程相关的线程中复制数据：

int __shfl_up(int var, unsigned int delta, int width = warpsize)；

__shfl_up(val, 2)将值转移到右边两个通道内。__shfl_up周围若没有线程束，所以线程束中最低的delta个线程delta将保持不变。

__shfl_down(bal, 2)将值转移到左边两个通道内。__shfl_down周围没有线程束，所以线程束中最高的delta线程将保持不变。

洗牌指令的最后一种形式是根据调用束内线程索引自身的按位异或来传输束内线程中的数据；

int __shfl_xor(int var, int laneMask, int width=warpSize)使用laneMask执行调用束内线程索引的按位异或，内部指令可计算源束内线程索引。该指令适合蝴蝶寻址模式(a butterfly addressing pattern)。

5.6.2 线程束内的共享数据

洗牌指令将被应用到以下3中整型变量类型中：
        1. 标量变量；
        2. 数组；
        3. 向量型变量

5.6.2.1 线程束内值的广播

__global__ void test_shfl_broadcast(int* d_out, int* d_in, int const srcLane)
{
    int val = d_in(threadIdx.X);
    value = __shfl(val, srcLane, BDIMX);
    d_out[threadIdx.x] = value;
}

该kernel使用16个线程的一维线程块。

5.6.2.2 线程束内上移

__global__ void test_shfl_up(int* d_out, int* d_in, int const delta)
{
    int val = d_in(threadIdx.X);
    value = __shfl_up(val, delta, BDIMX);
    d_out[threadIdx.x] = value;
}

5.6.2.3 线程束内下移

__global__ void test_shfl_down(int* d_out, int* d_in, int const delta)
{
    int val = d_in(threadIdx.X);
    value = __shfl_down(val, delta, BDIMX);
    d_out[threadIdx.x] = value;
}

5.6.2.4 线程束内环绕移动

每个线程的源束内线程是不同的，并且由它自身的束内线程索引加上偏移量来确定的。

// offset 左移为正，右移为负；
__global__ void test_shfl_warp(int* d_out, int* d_in, int const offset)
{
    int val = d_in(threadIdx.X);
    value = __shfl(val, threadIdx.x + offset, BDIMX);
    d_out[threadIdx.x] = value;
}

5.6.2.5 线程束的蝴蝶交换

__global__ void test_shfl_xor(int* d_out, int* d_in, int const mask)
{
    int val = d_in(threadIdx.X);
    value = __shfl_xor(val, mask, BDIMX);
    d_out[threadIdx.x] = value;
}

调用掩码为1的内核将导致相邻线程交换它们的值。

5.6.2.6 线程束交换数组值

考虑内核中使用寄存器数组的情况，若想要在线程束的线程间交换数据的某些部分，则可以使用洗牌指令交换线程束中线程间的数组元素。

在下面的内核中，每个线程都有一个寄存器数组val，大小为SEGM。每个线程从全局内存d_in中读取数据块到val中，使用由掩码确定的相邻线程交换该块，然后将接收数据写回全局内存数组中：

__global__ void test_shfl_xor_array(int* d_out, int* d_in, int const mask)
{
    int idx = threadIdx.x * SEGM;

    int value[SEGM];

    for (int i = 0; i < SEGM; i++) value[i] = d_in[idx + i];

    value[0] = __shfl_xor(value[0], mask, BDIMX);
    value[1] = __shfl_xor(value[1], mask, BDIMX);
    value[2] = __shfl_xor(value[2], mask, BDIMX);
    value[3] = __shfl_xor(value[3], mask, BDIMX);

    for (int i = 0; i < SEGM; i++)  d_out[idx + i] = value[i];
}

因为每个线程有4个元素，所以线程块被缩小到原来的1/4。

5.6.2.7 线程束使用数组索引交换数值

如果想在两个线程各自的数组中以不同的偏移量交换元素，需要有基于洗牌指令的交换函数。

__inline__ __device__ void swap(int* value, int laneIdx, int mask, int firstIdx, int secondIdx)
{
    bool pred = ((laneIdx / mask + 1) == 1)
    if (pred)
    {
        int temp = value[secondIdx];
        value[firstIdx] = value[secondIdx];
        value[secondIdx] = temp;
    }

    value[secondIdx] = __shfl_xor(value[secondIdx], mask, BDIMX);

    if (pred)
    {
        int tem = value[firstIdx];
        value[firstIdx] = value[secondIdx];
        value[secondIdx] = temp;
    }

}

__global__ void test_shfl_swap(int* d_out, int* d_in, int const mask, int firstIdx, int secondIdx)
{
    int idx = threadIdx * SEGM;
    int value[SEGM];

    for (int i = 0; i < SEGM; i++) value[i] = d_in[idx + i];

    swap(value, threadIdx.x, mask, firstIdx, secondIdx);

    for (int i = 0; i < SEGM; i++) d_out[idx + i] = value[i];
}

此swap仅能交换本地数组。上述内核基于上述swap函数，交换两个线程间不同偏移的两个元素。

5.6.3 使用线程束洗牌指令的并行归约

基本思路非常简单，它包括三个层面的归约；
        1. 线程束级归约；
        2. 线程块级归约；
        3. 网格级归约；

对于线程束级归约来说，各个线程束执行自己的归约。每个线程不使用共享内存，而是使用寄存器存储一个全局内存中读取的数据元素。

__inline__ __device__ int warpReduce(int mySum)
{
    mySum += __shfl_xor(mySum, 16);
    mySum += __shfl_xor(mySum, 8);
    mySum += __shfl_xor(mySum, 4);
    mySum += __shfl_xor(mySum, 2);
    mySum += __shfl_xor(mySum, 1);
    return mySum;
}

__global__ void reduceShfl(int* g_idata, int* g_odata, unsigned int int n)
{
    __shared__ int smem[SMEMDIM];

    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= n) return;

    int mySum = g_idata[idx];

    int laneIdx = threadIdx.x % warpSize;
    int warpIdx = threadIdx.x / warpSize;

    mySum = warpReduce(mySum);

    if (laneIdx == 0) smem[warpIdx] = mySum;

    __syncthreads();

    // last warp reduce
    mySum = (threadIdx.x < SMEMDIM) ? smem[laneIdx] : 0;
    if (warpIdx==0) mySum = warpReduce(mySum);

    if (threadIdx.x == 0) g_odata[blockIdx.x] = mySum;
}

5.7 总结

为了获得最大的应用性能，需要有一个显式管理的内存层次结构。

共享内存可以被声明为一维或二维数组。避免存储体冲突是在共享内存应用优化过程中一个重要的因素。共享内存被分配在所有常驻线程块中，因此，它是一个关键资源，可能会限制内核占用率。

在内核中使用共享内存有两个主要原因：
        1. 用于缓存片上内存并且减少全局内存访问次数；
        2. 传输共享内存中数据的安排方式，避免非合并的全局内存访问；

常量内存对只读数据进行了优化，这些数据每次都将数据广播到许多线程中。常量内存也使用自己的SM缓存，防止常量内存的读操作通过一级缓存干扰全局内存访问。因此，对合适的数据使用常量内存，不仅可优化特定项目的访问，还可能提高整体全局内存吞吐量。

只读纹理缓存提供了常量内存的替换方案，该方案优化了数据的分散读取。只读缓存访问全局内存中的数据，但它使用一个独立的内存访问流水线和独立的缓存，以使SM可以访问数据。

洗牌指令是线程束级的内部功能，能使线程束中内线程快速直接地共享数据。洗牌指令比共享内存有更低的延迟，并且不需要分配额外的资源。

使用洗牌指令可以减少内核中线程束同步优化的数目。

适当的使用各种内存类型可以提高带宽利用率，降低整体的内存延迟。如果你正在研究优化的因素，那么牢记共享内存，常量内存，只读缓存和洗牌指令是非常重要的。