1

一个问题涉及对存储在计算能力 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 的缩减示例，但我不确定它是否适用于我的问题。

4

3 回答 3

2

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

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

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

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

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

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

于 2012-11-02T02:24:07.023 回答
1

I think a DWORD copying is anyway faster than per-byte copying. Try this instead of your example:

for(int i = 0; i < 4; i++)
{
    ((int*)s_array)[4 * threadIdx.x + i] = ((int*)d_text)[i];
}
于 2012-11-01T23:51:56.147 回答
0

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

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

对于一维数组，它可以是

s_array[idx + idx / 16] = source[idx];
于 2021-12-03T16:25:33.730 回答