c - 从全局加载数据到共享内存时如何避免bank冲突

Question

一个问题涉及跨步访问存储在计算能力 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 位存储体，以便同时为同一个 half-warp 的 16 个线程提供服务。数据在存储体之间交错，第 i 个 32 位字始终存储在第 i % 16 - 1 个共享存储体中。

由于每个线程通过一次内存事务读取 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 的缩减示例，但我不确定它是否适用于我的问题。

Answer 1

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

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

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

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

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

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