库达与全球内存共享，可能加速

Question

我相信我的CUDA应用程序可能会从共享内存中受益，以便将数据保存在GPU核心附近。现在，我有一个内核，向它传递一个指向先前分配的设备内存块的指针，以及一些常量。内核完成后，设备内存包括结果，结果被复制到主机内存中。该方案工作良好，并与CPU上运行的相同算法进行了交叉验证.

这些文档非常清楚地表明，与共享内存相比，全局内存要慢得多，并且具有更高的访问延迟，但要获得最佳性能，您应该使线程合并并对齐任何访问。我的GPU具有计算能力6.1 "Pascal"，每个线程块有48 kiB共享内存和2 GiB DRAM。如果我重构代码以使用共享内存，如何确保避免银行冲突？

共享内存被组织在32个银行中，因此来自同一个块的32个线程可以同时访问不同的银行而不必等待。假设我从上面获取内核，启动一个内核配置，其中包含一个块和32个线程，并静态地在内核外部分配48 kiB共享内存。另外，每个线程只会从(共享)内存中的同一个单一内存位置读取和写入，这是我正在研究的算法所特有的。在此情况下，我将访问具有48 kiB / 32 banks / sizeof(double)偏移量(等于192个)的32个共享内存位置：

__shared__ double cache[6144];

__global__ void kernel(double *buf_out, double a, double b, double c)
{
    for(...)
    {
       // Perform calculation on shared memory
       cache[threadIdx.x * 192] = ...
    }

    // Write result to global memory
    buf_out[threadIdx.x] = cache[threadIdx.x * 192];
}

我的推理是:当threadIdx.x运行在0到31之间时，偏移量加上cache是一个double，确保每个线程同时访问不同银行的第一个元素。我还没有修改和测试代码，但是这是对SM访问的正确方式吗？

MWE补充道：这是一个天真的CPU到CUDA端口的算法，只使用全局内存.Visual报告内核执行时间为10.3秒。环境: Win10，MSVC 2019，x64发行版构建，CUDA v11.2。

#include "cuda_runtime.h"

#include <iostream>
#include <stdio.h>

#define _USE_MATH_DEFINES
#include <math.h>
    

__global__ void kernel(double *buf, double SCREEN_STEP_SIZE, double APERTURE_RADIUS,
    double APERTURE_STEP_SIZE, double SCREEN_DIST, double WAVE_NUMBER)
{   
    double z, y, y_max;

    unsigned int tid = threadIdx.x/* + blockIdx.x * blockDim.x*/;
    
    double Z = tid * SCREEN_STEP_SIZE, Y = 0;

    double temp = WAVE_NUMBER / SCREEN_DIST;


    // Make sure the per-thread accumulator is zero before we begin
    buf[tid] = 0;

    for (z = -APERTURE_RADIUS; z <= APERTURE_RADIUS; z += APERTURE_STEP_SIZE)
    {
        y_max = sqrt(APERTURE_RADIUS * APERTURE_RADIUS - z * z);

        for (y = -y_max; y <= y_max; y += APERTURE_STEP_SIZE)
        {
            buf[tid] += cos(temp * (Y * y + Z * z));
        }
    }   
}


int main(void)
{
    double *dev_mem;
    double *buf = NULL;
    cudaError_t cudaStatus;

    unsigned int screen_elems = 1000;


    if ((buf = (double*)malloc(screen_elems * sizeof(double))) == NULL)
    {
        printf("Could not allocate memory...");
        return -1;
    }

    memset(buf, 0, screen_elems * sizeof(double));


    if ((cudaStatus = cudaMalloc((void**)&dev_mem, screen_elems * sizeof(double))) != cudaSuccess)
    {
        printf("cudaMalloc failed with code %u", cudaStatus);
        return cudaStatus;
    }


    kernel<<<1, 1000>>>(dev_mem, 1e-3, 5e-5, 50e-9, 10.0, 2 * M_PI / 5e-7);

    cudaDeviceSynchronize();

    if ((cudaStatus = cudaMemcpy(buf, dev_mem, screen_elems * sizeof(double), cudaMemcpyDeviceToHost)) != cudaSuccess)
    {
        printf("cudaMemcpy failed with code %u", cudaStatus);
        return cudaStatus;
    }


    cudaFree(dev_mem);

    cudaDeviceReset();

    free(buf);

    return 0;
}

下面的内核使用共享内存，执行时间大约为10.6秒，在Visual中也是这样：

__shared__ double cache[1000];


__global__ void kernel(double *buf, double SCREEN_STEP_SIZE, double APERTURE_RADIUS,
    double APERTURE_STEP_SIZE, double SCREEN_DIST, double WAVE_NUMBER)
{   
    double z, y, y_max;

    unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
    
    double Z = tid * SCREEN_STEP_SIZE, Y = 0;

    double temp = WAVE_NUMBER / SCREEN_DIST;


    // Make sure the per-thread accumulator is zero before we begin
    cache[tid] = 0;

    for (z = -APERTURE_RADIUS; z <= APERTURE_RADIUS; z += APERTURE_STEP_SIZE)
    {
        y_max = sqrt(APERTURE_RADIUS * APERTURE_RADIUS - z * z);

        for (y = -y_max; y <= y_max; y += APERTURE_STEP_SIZE)
        {
            cache[tid] += cos(temp * (Y * y + Z * z));
        }
    }   

    buf[tid] = cache[tid];
} 

循环中最内部的行通常执行数百万次，这取决于传递给内核的五个常量。所以，我并没有破坏片外的全球内存，而是预期片上共享内存版本会快得多，但显然不是--我错过了什么？

Answer 1

比方说..。每个线程只会从(共享)内存中读取和写入相同的单个内存位置，这是我正在研究的算法所特有的。

在这种情况下，使用共享内存是没有意义的。共享记忆的全部意义是分享..。在一个块中的所有线程中。根据您的假设，您应该将元素保存在寄存器中，而不是共享内存中。事实上，在您的"MWE添加“内核--这可能是您应该做的。

如果您的线程要共享信息，那么这种共享的模式将决定如何最好地利用共享内存。

还请记住，如果您不反复读取数据或从多个线程读取数据，那么共享内存帮助您的可能性就小得多--因为您总是必须至少从全局内存读取一次数据，并至少写入共享内存一次才能将数据存储在共享内存中。