Accelerated Ray Tracing （八）

https://developer.nvidia.com/blog/using-shared-memory-cuda-cc/

Shared Memory

在上一篇文章中，我讨论了如何将一组线程访问的全局内存合并到一个事务中，以及对齐和跨步如何影响不同代CUDA硬件的合并。对于最新版本的CUDA硬件，未对齐的数据访问不是一个大问题。然而，无论使用哪种CUDA硬件，跨越全局内存都是有问题的，而且在许多情况下似乎是不可避免的，比如在沿着第二个维度和更高维度访问多维数组中的元素时。但是，如果我们使用共享内存，在这种情况下可以合并内存访问。在下一篇文章中，我将向您展示如何避免跨越全局内存，在此之前，我首先需要对共享内存进行一些详细的描述。

Shared Memory

因为共享内存在芯片上，所以它比本地和全局内存快得多。实际上，共享内存延迟大约比未缓存的全局内存延迟低100倍（前提是线程之间不存在存储体冲突，我们将在本文后面进行检查）。共享内存是为每个线程块分配的，因此该块中的所有线程都可以访问同一共享内存。线程可以访问同一线程块中的其他线程从全局内存加载的共享内存中的数据。此功能（与线程同步结合）具有多种用途，例如用户管理的数据缓存，高性能协作并行算法（例如，并行减少），并在其他情况下无法实现全局内存合并。

Thread Synchronization

在线程之间共享数据时，我们需要注意避免出现竞争状况，因为尽管块中的线程在逻辑logically 上并行运行，但并非所有线程都可以在物理physically 上同时执行。假设两个线程A和B各自从全局内存中加载一个数据元素并将其存储到共享内存中。然后，线程A要从共享内存中读取B的元素，反之亦然。假设A和B是两个不同warp中的线程。如果B在A尝试读取其元素之前尚未完成其元素的写入，则我们有一个竞争条件，这可能导致不确定的行为和错误的结果。

为了确保并行线程协作时的正确结果，我们必须同步线程 __syncthreads()。线程的执行只能在其块中的所有线程都已执行__syncthreads（）之后才能通过__syncthreads（）。因此，我们可以通过在存储到共享内存之后以及从共享内存加载任何线程之前调用__syncthreads（）来避免上述竞争情况。请务必注意，在不同代码中调用__syncthreads（）是未定义的，并且可能导致死锁-线程块中的所有线程必须在同一点调用__syncthreads（）。

Shared Memory Example

使用__shared__变量声明说明符在CUDA C / C ++设备代码中声明共享内存。在内核中声明共享内存的方式有多种，具体取决于在编译时还是在运行时知道内存量。

此代码使用共享内存反转64元素数组中的数据。这两个内核非常相似，仅在如何声明共享内存数组以及如何调用内核方面有所不同。

Static Shared Memory

在此示例中使用共享内存的原因是为了便于在较旧的CUDA设备（计算能力1.1或更早版本）上进行全局内存合并。由于始终通过线性、对齐的索引t 访问全局存储器，因此可以实现读写操作的最佳global memory合并。反向索引tr仅用于访问share memory，该共享内存没有全局内存的顺序访问限制以获得最佳性能。共享内存的唯一性能问题是存储区冲突，我们将在后面讨论。（请注意，在Compute Capability 1.2或更高版本的设备上，即使反向索引存储到全局内存，内存系统也可以完全合并。但是，这种技术对其他访问方式仍然有用，我将在下一篇文章中展示。）

#include <stdio.h>

__global__ void staticReverse(int* d, int n)
{
    __shared__ int s[64];//如果像在staticReverse内核中那样在编译时知道共享内存数组的大小，则可以像对array s一样显式声明一个具有该大小的数组。
    int t = threadIdx.x;//在此内核中，t和tr是分别表示原始顺序和反向顺序的两个索引。
    int tr = n - t - 1;
    s[t] = d[t];//线程使用语句s [t] = d [t]将数据从global memory复制到share memory，并在之后两行使用语句d [t] = s [tr]进行反转。
    __syncthreads();//但是在执行最后一行（其中每个线程访问另一个线程写入的共享内存中的数据）之前，请记住我们需要通过调用__syncthreads（）确保所有线程已完成对共享内存的加载。
    d[t] = s[tr];
}

__global__ void dynamicReverse(int* d, int n)
{//动态共享内存内核dynamicReverse（）使用未调整大小的extern数组语法extern shared int s []声明共享内存数组（请注意用方括号和extern说明符的用法）。启动内核时，该大小由第三个执行配置参数隐式确定。内核代码的其余部分与staticReverse（）内核相同。
    extern __shared__ int s[];
    int t = threadIdx.x;
    int tr = n - t - 1;
    s[t] = d[t];
    __syncthreads();
    d[t] = s[tr];
}

int main(void)
{
    const int n = 64;
    int a[n], r[n], d[n];

    for (int i = 0; i < n; i++) {
        a[i] = i;
        r[i] = n - i - 1;
        d[i] = 0;
    }

    int* d_d;
    cudaMalloc(&d_d, n * sizeof(int));

    // run version with static shared memory
    cudaMemcpy(d_d, a, n * sizeof(int), cudaMemcpyHostToDevice);
    staticReverse << <1, n >> > (d_d, n);
    cudaMemcpy(d, d_d, n * sizeof(int), cudaMemcpyDeviceToHost);
    for (int i = 0; i < n; i++)
        if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);

    // run dynamic shared memory version
    cudaMemcpy(d_d, a, n * sizeof(int), cudaMemcpyHostToDevice);
    dynamicReverse << <1, n, n * sizeof(int) >> > (d_d, n);//此示例中的其他三个内核使用动态分配的共享内存，当在编译时不知道共享内存的数量时可以使用它们。在这种情况下，必须使用可选的第三个执行配置参数指定每个线程块的共享内存分配大小（以字节为单位），
    cudaMemcpy(d, d_d, n * sizeof(int), cudaMemcpyDeviceToHost);
    for (int i = 0; i < n; i++)
        if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);
}

Dynamic Shared Memory

如果您在单个内核中需要多个动态大小的阵列该怎么办？您必须像以前一样声明一个外部大小未定的数组，并使用其中的指针将其分成多个数组，如以下摘录所示。

extern __shared__ int s[];
int *integerData = s;                        // nI ints
float *floatData = (float*)&integerData[nI]; // nF floats
char *charData = (char*)&floatData[nF];      // nC chars

在内核启动中，指定所需的总共享内存，如下所示。

myKernel<<<gridSize, blockSize, nI*sizeof(int)+nF*sizeof(float)+nC*sizeof(char)>>>(...);

Shared memory bank conflicts

为了实现并发访问的高内存带宽，共享内存被分为大小相等的内存模块（存储体），可以同时访问这些内存模块。因此，可以同时处理跨越b个不同存储体的n个地址的任何存储负载或存储，产生的有效带宽是单个存储体带宽的b倍。

但是，如果多个线程请求的地址映射到同一存储库，则访问将被序列化。硬件根据需要将冲突的内存请求分成尽可能多的单独的无冲突请求，从而将有效带宽减少等于冲突的内存请求数量的因数。异常是warp中的所有线程都使用相同的共享内存地址，从而导致广播。计算能力为2.0和更高版本的设备还具有多播共享内存访问的附加功能，这意味着将同时服务经线中任意数量的线程对同一位置的多次访问。

为了最大程度地减少存储体冲突，了解内存地址如何映射到存储体非常重要。共享存储体的组织方式是，将连续的32位字分配给连续的存储体，带宽为每个时钟周期每个存储体32位。对于计算能力为1.x的设备，warp大小为32个线程，存储区数为16。对warp的共享内存请求被分为一个对warp前半部分的请求和一个对warp后半部分的请求。请注意，如果半个线程束的线程访问每个存储区只有一个内存位置，则不会发生存储区冲突。

对于计算能力为2.0的设备，warp大小为32个线程，并且组数也为32。对于warp的共享内存请求不会像计算能力1.x的设备那样拆分，这意味着在warp的前半部分的线程与同一warp的后半部分的线程之间可能会发生存储区冲突。

计算能力3.x的设备具有可配置的存储区大小，可以使用cudaDeviceSetSharedMemConfig（）将其设置为四个字节（cudaSharedMemBankSizeFourByte，默认值）或八个字节（cudaSharedMemBankSizeEightByte）。将存储体大小设置为八个字节可以帮助避免在访问双精度数据时出现共享存储体冲突。

Configuring the amount of shared memory

在具有计算能力2.x和3.x的设备上，每个多处理器具有64KB的片上内存，可以在L1缓存和共享内存之间进行分区。对于具有2.x计算能力的设备，有两个设置，即48KB共享内存/ 16KB L1缓存和16KB共享内存/ 48KB L1缓存。默认情况下使用48KB共享内存设置。可以在运行时API上从主机为所有内核使用cudaDeviceSetCacheConfig（）或在每个内核的基础上使用cudaFuncSetCacheConfig（）进行配置。它们接受以下三个选项之一：cudaFuncCachePreferNone，cudaFuncCachePreferShared和cudaFuncCachePreferL1。驱动程序将遵循指定的首选项，除非内核要求每个线程块的共享内存多于指定配置中可用的内存。计算能力3.x的设备允许第三种设置32KB共享内存/ 32KB L1缓存，可以使用选项cudaFuncCachePreferEqual获得。

共享内存是用于编写优化优化的CUDA代码的强大功能。对共享内存的访问比全局内存访问快得多，因为它位于芯片上。由于共享内存是由线程块中的线程共享的，因此它提供了线程协作的机制。使用共享内存的一种利用这种线程协作的方法是启用全局内存合并，如本文中的数组反转所示。通过使用共享内存反转阵列，我们能够以跨步执行所有全局内存读取和写入，从而在任何CUDA GPU上实现完全合并。