将数据从全局加载到共享内存时如何避免银行冲突

Posted 2023-03-23

【中文标题】将数据从全局加载到共享内存时如何避免银行冲突

【英文标题】：How to avoid bank conflicts when loading data from global to shared memory

【发布时间】：2012-10-22 10:37:42

【问题描述】：

一个问题涉及对存储在计算能力 1.3 GPU 的全局内存中的无符号字符数组进行跨步访问。为了绕过全局内存的合并要求，线程顺序访问全局内存并将数组复制到共享内存，仅使用 2 个内存事务，例如：

#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>

__global__ void kernel ( unsigned char *d_text, unsigned char *d_out ) 

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

    extern __shared__ unsigned char s_array[];

    uint4 *uint4_text = ( uint4 * ) d_text;
    uint4 var;

    //memory transaction
    var = uint4_text[0];

    uchar4 c0 = *reinterpret_cast<uchar4 *>(&var.x);
    uchar4 c4 = *reinterpret_cast<uchar4 *>(&var.y);
    uchar4 c8 = *reinterpret_cast<uchar4 *>(&var.z);
    uchar4 c12 = *reinterpret_cast<uchar4 *>(&var.w);

    s_array[threadIdx.x*16 + 0] = c0.x;
    s_array[threadIdx.x*16 + 1] = c0.y;
    s_array[threadIdx.x*16 + 2] = c0.z;
    s_array[threadIdx.x*16 + 3] = c0.w;

    s_array[threadIdx.x*16 + 4] = c4.x;
    s_array[threadIdx.x*16 + 5] = c4.y;
    s_array[threadIdx.x*16 + 6] = c4.z;
    s_array[threadIdx.x*16 + 7] = c4.w;

    s_array[threadIdx.x*16 + 8] = c8.x;
    s_array[threadIdx.x*16 + 9] = c8.y;
    s_array[threadIdx.x*16 + 10] = c8.z;
    s_array[threadIdx.x*16 + 11] = c8.w;

    s_array[threadIdx.x*16 + 12] = c12.x;
    s_array[threadIdx.x*16 + 13] = c12.y;
    s_array[threadIdx.x*16 + 14] = c12.z;
    s_array[threadIdx.x*16 + 15] = c12.w;

    d_out[idx] = s_array[threadIdx.x*16];


int main ( void ) 

    unsigned char *d_text, *d_out;

    unsigned char *h_out = ( unsigned char * ) malloc ( 32 * sizeof ( unsigned char ) );
    unsigned char *h_text = ( unsigned char * ) malloc ( 32 * sizeof ( unsigned char ) );

    int i;

    for ( i = 0; i < 32; i++ )
        h_text[i] = 65 + i;

    cudaMalloc ( ( void** ) &d_text, 32 * sizeof ( unsigned char ) );
    cudaMalloc ( ( void** ) &d_out, 32 * sizeof ( unsigned char ) );

    cudaMemcpy ( d_text, h_text, 32 * sizeof ( unsigned char ), cudaMemcpyHostToDevice );

    kernel<<<1,32,16128>>>(d_text, d_out );

    cudaMemcpy ( h_out, d_out, 32 * sizeof ( unsigned char ), cudaMemcpyDeviceToHost );

    for ( i = 0; i < 32; i++ )
        printf("%c\n", h_out[i]);

    return 0;

问题是在将数据复制到共享内存时会发生存储库冲突（nvprof 报告的上述示例中存在 384 个冲突），从而导致线程的序列化访问。

共享内存被划分为 16 个（或在较新的设备架构上为 32 个）32 位组，以便同时为同一个半扭曲的 16 个线程提供服务。数据在 bank 之间交错，第 i 个 32 位字始终存储在 i % 16 - 1 共享内存 bank 中。

由于每个线程在一次内存事务中读取 16 个字节，因此字符将以跨步方式存储到共享内存中。这会导致线程 0、4、8、12 之间发生冲突； 1、5、9、13； 2、6、10、14； 3、7、11、15 个相同的半经线。消除银行冲突的一种简单方法是使用 if/else 分支以类似于以下的循环方式将数据存储到共享内存，但会导致一些严重的线程分歧：

int tid16 = threadIdx.x % 16;

if ( tid16 < 4 ) 

    s_array[threadIdx.x * 16 + 0] = c0.x;
    s_array[threadIdx.x * 16 + 1] = c0.y;
    s_array[threadIdx.x * 16 + 2] = c0.z;
    s_array[threadIdx.x * 16 + 3] = c0.w;

    s_array[threadIdx.x * 16 + 4] = c4.x;
    s_array[threadIdx.x * 16 + 5] = c4.y;
    s_array[threadIdx.x * 16 + 6] = c4.z;
    s_array[threadIdx.x * 16 + 7] = c4.w;

    s_array[threadIdx.x * 16 + 8] = c8.x;
    s_array[threadIdx.x * 16 + 9] = c8.y;
    s_array[threadIdx.x * 16 + 10] = c8.z;
    s_array[threadIdx.x * 16 + 11] = c8.w;

    s_array[threadIdx.x * 16 + 12] = c12.x;
    s_array[threadIdx.x * 16 + 13] = c12.y;
    s_array[threadIdx.x * 16 + 14] = c12.z;
    s_array[threadIdx.x * 16 + 15] = c12.w;

 else if ( tid16 < 8 ) 

    s_array[threadIdx.x * 16 + 4] = c4.x;
    s_array[threadIdx.x * 16 + 5] = c4.y;
    s_array[threadIdx.x * 16 + 6] = c4.z;
    s_array[threadIdx.x * 16 + 7] = c4.w;

    s_array[threadIdx.x * 16 + 8] = c8.x;
    s_array[threadIdx.x * 16 + 9] = c8.y;
    s_array[threadIdx.x * 16 + 10] = c8.z;
    s_array[threadIdx.x * 16 + 11] = c8.w;

    s_array[threadIdx.x * 16 + 12] = c12.x;
    s_array[threadIdx.x * 16 + 13] = c12.y;
    s_array[threadIdx.x * 16 + 14] = c12.z;
    s_array[threadIdx.x * 16 + 15] = c12.w;

    s_array[threadIdx.x * 16 + 0] = c0.x;
    s_array[threadIdx.x * 16 + 1] = c0.y;
    s_array[threadIdx.x * 16 + 2] = c0.z;
    s_array[threadIdx.x * 16 + 3] = c0.w;

 else if ( tid16 < 12 ) 

    s_array[threadIdx.x * 16 + 8] = c8.x;
    s_array[threadIdx.x * 16 + 9] = c8.y;
    s_array[threadIdx.x * 16 + 10] = c8.z;
    s_array[threadIdx.x * 16 + 11] = c8.w;

    s_array[threadIdx.x * 16 + 12] = c12.x;
    s_array[threadIdx.x * 16 + 13] = c12.y;
    s_array[threadIdx.x * 16 + 14] = c12.z;
    s_array[threadIdx.x * 16 + 15] = c12.w;

    s_array[threadIdx.x * 16 + 0] = c0.x;
    s_array[threadIdx.x * 16 + 1] = c0.y;
    s_array[threadIdx.x * 16 + 2] = c0.z;
    s_array[threadIdx.x * 16 + 3] = c0.w;

    s_array[threadIdx.x * 16 + 4] = c4.x;
    s_array[threadIdx.x * 16 + 5] = c4.y;
    s_array[threadIdx.x * 16 + 6] = c4.z;
    s_array[threadIdx.x * 16 + 7] = c4.w;

 else 

    s_array[threadIdx.x * 16 + 12] = c12.x;
    s_array[threadIdx.x * 16 + 13] = c12.y;
    s_array[threadIdx.x * 16 + 14] = c12.z;
    s_array[threadIdx.x * 16 + 15] = c12.w;

    s_array[threadIdx.x * 16 + 0] = c0.x;
    s_array[threadIdx.x * 16 + 1] = c0.y;
    s_array[threadIdx.x * 16 + 2] = c0.z;
    s_array[threadIdx.x * 16 + 3] = c0.w;

    s_array[threadIdx.x * 16 + 4] = c4.x;
    s_array[threadIdx.x * 16 + 5] = c4.y;
    s_array[threadIdx.x * 16 + 6] = c4.z;
    s_array[threadIdx.x * 16 + 7] = c4.w;

    s_array[threadIdx.x * 16 + 8] = c8.x;
    s_array[threadIdx.x * 16 + 9] = c8.y;
    s_array[threadIdx.x * 16 + 10] = c8.z;
    s_array[threadIdx.x * 16 + 11] = c8.w;

任何人都可以提出更好的解决方案吗？我已经研究过 SDK 的缩减示例，但我不确定它是否适用于我的问题。

【参考方案1】：

虽然代码会导致银行冲突，但这并不意味着它会更慢。

在您的计算能力 1.3 GPU 上，具有 2 路银行冲突的共享内存事务只比没有银行冲突的一个多两个周期。在两个周期内，您甚至无法执行一条指令来解决银行冲突。与无冲突访问相比，4 路存储库冲突多使用六个周期，这足以执行一次额外的无冲突共享内存访问。

在您的情况下，代码很可能受到全局内存带宽（和延迟，即数百个周期，即比我们在这里讨论的 2..6 个周期大两个数量级）的限制​​。因此，您可能会有大量空闲周期可用，其中 SM 只是空闲等待来自全局内存的数据。然后银行冲突可以使用这些周期而不会减慢您的代码完全。

确保编译器将 .x、.y、.z 和 .w 的四个字节存储合并到一个 32 位访问中更为重要。使用cuobjdump -sass 查看编译后的代码，看看是否是这种情况。如果不是，请按照 Otter 的建议改用单词传输。

如果您只是从d_text 读取而不是从内核中写入，您也可以为其使用纹理，它仍然会比具有银行冲突的内核慢，但可能会提供其他优势来提高速度整体（例如，如果您不能保证全局内存中的数据正确对齐）。

另一方面，您的替代银行无冲突代码将快速的 256 字节全局内存拆分为四个 64 位事务，这些事务的效率要低得多，并且可能会溢出正在运行的内存事务的最大数量，因此你会招致四百到几千个全局内存延迟周期。 为避免这种情况，您需要首先使用 256 字节宽读取将数据传输到寄存器，然后以无存储冲突的方式将数据从寄存器移动到共享内存。尽管如此，仅 register->shmem 移动的代码将比我们试图解决的六个周期多得多。

【讨论】：

从内存中取出单词后发生银行冲突。我不确定全局内存延迟是否可以有效地隐藏它们。每字节复制仅用于测试目的。我实际上是直接将 uint4 字复制到共享内存中。 d_text 的每个字节只读取一次，因此 AFAIK 纹理缓存不会有任何好处。对于大小为 116.234.496 字节的 d_text 数组，将它们存储在共享内存中，然后从共享内存中读取以处理它们会导致 3736445 冲突。 SDK 中的缩减示例声称在避免银行冲突时提高了 2 倍以上

所以银行冲突在 GTX 260 上会消耗大约 0.25 毫秒的实时时间，而读取 116.234.496 字节大约需要 1 毫秒。你的内核需要多长时间才能执行？我的主要观点是不同的：硬件已经采取了解决银行冲突的最佳方法，你不能编写代码来执行相同的任务，即每个线程将 16 个连续字节传输到共享内存的速度更快。您唯一能做的就是重新排列数据或算法的布局，但为了帮助您实现这一点，我们需要更多关于您想要实现的目标的信息。

例如，如果您可以在写入共享内存之前将每个线程的 16 个字节处理成寄存器中的四个字节，这将完全消除银行冲突（但再次仅在您必须这样做时提供加速无论如何处理）。归约情况是一个理想的例子。

该算法用于字符串匹配。线程必须按顺序处理一组字节，以便可以使用来自先前扫描字符的信息来潜在地跳过扫描下一个字符。我不认为对字符进行任何进一步的处理以例如将它们从 16 减少到 4 是可行的，尽管我知道这将是一个理想的解决方案

可能读取 int 而不是 uint4 也不会有任何好处。它将消除存储库冲突，但仅使用 1.3 计算能力 GPU 上全局内存的一半带宽

【参考方案2】：

我认为 DWORD 复制无论如何都比按字节复制快。 试试这个而不是你的例子：

for(int i = 0; i < 4; i++)

    ((int*)s_array)[4 * threadIdx.x + i] = ((int*)d_text)[i];

【讨论】：

我实际上将 uint4 个字（128 位）直接存储到共享内存中。每字节处理是尝试找到解决银行冲突问题的临时方法

【参考方案3】：

为了避免存储库冲突，人们经常在共享内存中的虚拟矩阵中添加备用列。所以你也许可以将你的共享数组大小增加 1/16 并替换为

threadIdx.x * 17 + 0
threadIdx.x * 17 + 1
...
threadIdx.x * 17 + 15

对于一维数组，它可以是

s_array[idx + idx / 16] = source[idx];