Accelerated Ray Tracing （七）

https://developer.nvidia.com/blog/how-access-global-memory-efficiently-cuda-c-kernels/

global memory

CUDA 不同的显存，他们各自有不同的范围，生存期和缓存行为。设备DRAM中的global memory，用于主机和设备之间的传输以及用于内核的数据输入和输出。global 表明作用域，可以被host和device一起访问。可以使用__device__声明说明符在global（变量）范围内声明global memory，或使用cudaMalloc（）动态分配并分配给常规C指针变量。Global memory分配可以在应用程序的生存期内持续存在。根据device的计算能力，全局内存可能会或可能不会缓存在芯片上。

__device__ int globalArray[256];
 
void foo()
{
    ...
    int *myDeviceMemory = 0;
    cudaError_t result = cudaMalloc(&myDeviceMemory, 256 * sizeof(int));
    ...
}

我们已经讨论了如何将线程分组为线程块，这些线程块被分配给设备上的多处理器。在执行过程中，有一个更细的线程组成warp。GPU上的多处理器以SIMD(单指令多数据)的方式执行指令。warp大小(有效的SIMD宽度)的所有当前cuda能力gpu是32线程。

Global Memory Coalescing

将线程分组到warp中不仅与计算有关，而且与全局内存访问也有关。该设备将全局内存load和由线程发出的存储warp为尽可能少的事务，以最小化DRAM带宽(在计算能力小于2.0的旧硬件上，事务合并在16个线程的一半范围内，而不是整个范围)。

#include <stdio.h>
#include <assert.h>

// Convenience function for checking CUDA runtime API results
// can be wrapped around any runtime API call. No-op in release builds.
inline
cudaError_t checkCuda(cudaError_t result)
{
#if defined(DEBUG) || defined(_DEBUG)
    if (result != cudaSuccess) {
        fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
        assert(result == cudaSuccess);
    }
#endif
    return result;
}

template <typename T>
__global__ void offset(T* a, int s)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x + s;
    a[i] = a[i] + 1;
}

template <typename T>
__global__ void stride(T* a, int s)
{
    int i = (blockDim.x * blockIdx.x + threadIdx.x) * s;
    a[i] = a[i] + 1;
}

template <typename T>
void runTest(int deviceId, int nMB)
{
    int blockSize = 256;
    float ms;

    T* d_a;
    cudaEvent_t startEvent, stopEvent;

    int n = nMB * 1024 * 1024 / sizeof(T);

    // NB:  d_a(33*nMB) for stride case
    checkCuda(cudaMalloc(&d_a, n * 33 * sizeof(T)));

    checkCuda(cudaEventCreate(&startEvent));
    checkCuda(cudaEventCreate(&stopEvent));

    printf("Offset, Bandwidth (GB/s):\n");

    offset << <n / blockSize, blockSize >> > (d_a, 0); // warm up

    for (int i = 0; i <= 32; i++) {
        checkCuda(cudaMemset(d_a, 0, n * sizeof(T)));

        checkCuda(cudaEventRecord(startEvent, 0));
        offset << <n / blockSize, blockSize >> > (d_a, i);
        checkCuda(cudaEventRecord(stopEvent, 0));
        checkCuda(cudaEventSynchronize(stopEvent));

        checkCuda(cudaEventElapsedTime(&ms, startEvent, stopEvent));
        printf("%d, %f\n", i, 2 * nMB / ms);
    }

    printf("\n");
    printf("Stride, Bandwidth (GB/s):\n");

    stride << <n / blockSize, blockSize >> > (d_a, 1); // warm up
    for (int i = 1; i <= 32; i++) {
        checkCuda(cudaMemset(d_a, 0, n * sizeof(T)));

        checkCuda(cudaEventRecord(startEvent, 0));
        stride << <n / blockSize, blockSize >> > (d_a, i);
        checkCuda(cudaEventRecord(stopEvent, 0));
        checkCuda(cudaEventSynchronize(stopEvent));

        checkCuda(cudaEventElapsedTime(&ms, startEvent, stopEvent));
        printf("%d, %f\n", i, 2 * nMB / ms);
    }

    checkCuda(cudaEventDestroy(startEvent));
    checkCuda(cudaEventDestroy(stopEvent));
    cudaFree(d_a);
}

int main(int argc, char** argv)
{
    int nMB = 4;
    int deviceId = 0;
    bool bFp64 = false;

    for (int i = 1; i < argc; i++) {
        if (!strncmp(argv[i], "dev=", 4))
            deviceId = atoi((char*)(&argv[i][4]));
        else if (!strcmp(argv[i], "fp64"))
            bFp64 = true;
    }

    cudaDeviceProp prop;

    checkCuda(cudaSetDevice(deviceId))
        ;
    checkCuda(cudaGetDeviceProperties(&prop, deviceId));
    printf("Device: %s\n", prop.name);
    printf("Transfer size (MB): %d\n", nMB);

    printf("%s Precision\n", bFp64 ? "Double" : "Single");

    if (bFp64) runTest<double>(deviceId, nMB);
    else       runTest<float>(deviceId, nMB);
}

通过传递“fp64”命令行选项，这段代码可以以单精度(默认)或双精度运行offset和stride内核。每个内核接受两个参数，一个输入数array和一个表示用于访问数组元素的offset量或stride的整数。在offset量和stride范围内循环调用内核。

Misaligned Data Accesses

在device memory中分配的array被CUDA驱动对齐到256字节的内存段。device可以通过32字节、64字节或128字节的transactions访问全局内存，这些transactions与它们的大小一致。对于C870或任何其他计算能力为1.0的设备，线程的半warp(或对齐访问，其中半warp的线程不按顺序访问内存)导致16个单独的32字节transactions。因为只有4个字节是要求每32字节的transactions,可能希望有效带宽减少八分之一，这大约是我们在上图中（棕线）看到的，即偏移量不是16个元素的倍数，对应于线程的一半warp。

对于Tesla C1060或其他计算能力为1.2或1.3的设备，未对齐的访问问题较少。基本上，通过半个线程warp对连续数据进行的未对齐访问在“cover”请求的数据的一些transactions 中得到处理。相对于对齐的情况，由于传输了未请求的数据以及不同半warp所请求的数据有一些overlap ，因此仍然存在性能损失，但是损失远远小于C870。

具有计算能力2.0的设备（例如Tesla C2050）在每个多处理器中具有一个128字节的行大小的L1高速缓存。该设备将线程中的线程的访问合并到尽可能少的高速缓存行中，从而导致跨线程的顺序内存访问的对齐方式对吞吐量的影响可忽略不计。

Strided Memory Access

步幅内核的结果如下图所示。

对于跨越式的全局内存访问，我们有不同的看法。对于较大的跨度，无论体系结构版本如何，有效带宽都很差。这并不奇怪：当并发线程同时访问物理内存中相距甚远的内存地址时，硬件就没有机会组合访问。您可以在上图中看到，在Tesla C870上，除1以外的任何步幅都会导致有效带宽大大降低。这是因为计算能力1.0和1.1硬件需要跨线程进行线性，对齐的访问以进行合并，所以我们看到了在偏移内核中也看到的熟悉的1/8带宽。计算能力1.2和更高版本的硬件可以合并访问，这些访问分为对齐的段（在CC 1.2 / 1.3上为32、64或128字节段，在CC 2.0或更高版本上为128字节高速缓存行），因此该硬件可实现平滑的带宽曲线。

当访问多维数组时，线程通常需要索引数组的更高维，因此无法避免跨越式访问。我们可以通过使用一种称为共享内存的CUDA内存来处理这些情况。共享内存是线程块中所有线程共享的片上内存。共享内存的一种用途是以合并的方式从全局内存中提取多维数组的2D切片到共享内存中，然后使相邻线程跨过共享内存块。

设备上的全局内存访问与主机上的数据访问共享性能特征；也就是说，数据局部性非常重要。在早期的CUDA硬件中，内存访问对齐与跨线程的局部性一样重要，但是在最近的硬件对齐方面，并不是一个大问题。另一方面，跨越式的内存访问会损害性能，而使用片上共享内存可以缓解这种情况。