CUDA:最大限度地减少大型数据类型的银行冲突

Question

32个银行被映射为32个连续的单词。我想知道数据类型是否很大，比如说

struct foo{
    float data[n];
};

__global__
void kernel(foo* d_ptr){
    __shared__ foo sh_data[number_threads_block];
    int tid = threadIdx.x;
    sh_data[tid] = d_ptr[tid + blockDim.x * blockIdx.x];
    __syncthreads();

    sh_data[tid] = ...
}

其中n被选择为8(或16，32)。然后，当我们访问sh_datatid时，如果我做对了，就会有8(或16，32)个银行冲突。

如果是这样的话，是否有任何技术可以尽量减少银行间的冲突？

谢谢

Answer 1

1. 实际上，您有一个结构数组(AoS)。这对GPU编程来说是典型的坏处。您可以使用标准的AoS->SoA数据重组方法来修复访问，以便相邻线程访问相邻的元素(这将防止银行冲突)。
2. 将数据加载到更大的块中，例如重组您的结构，以便它可以表示4个float4数量，而不是16 float数量。编译器可能能够将负载组织到float4负载中，这将减少银行冲突。如果确实需要这两种访问方法，甚至可以在结构中使用联合。

根据注释中的问题，让我们把第二个例子画出来。对于每个结构的16个float数量，无论是存储为float数组还是float4数组，存储模式如下：

(key: SXY  = float[Y] in Structure S[X], BX == Bank X)
B00 B01 B02 B03 B04 B05 B06 B07 B08 B09 B10 B11 B12 B13 B14 B15 B16 B17 B18 B19 B20 B21 B22 B23 B24 B25 B26 B27 B28 B29 B30 B31
S00 S01 S02 S03 S04 S05 S06 S07 S08 S09 S0A S0B S0C S0D S0E S0F S10 S11 S12 S13 S14 S15 S16 S17 S18 S19 S1A S1B S1C S1D S1E S1F
S20 S21 S22 S23 S24 S25 S26 S27 S28 S29 S2A S2B S2C S2D S2E S2F S30 S31 S32 S33 S34 S35 S36 S37 S38 S39 S3A S3B S3C S3D S3E S3F
S40 S41 S24 S43 S44 S45 S46 S47 S48 S49 S4A S4B S4C S4D S4E S4F S50 S51 S52 S53 S54 S55 S56 S57 S58 S59 S5A S5B S5C S5D S5E S5F
...

现在，假设每个线程的存储如下所示：

const int n = 16;
struct foo{
    float data[n];
};

我们的“装载”操作如下：

sh_data[tid] = ...

编译器无法在每个线程的单个指令中加载16x4字节，因此它将将上述加载操作分解为一系列请求。我认为这个序列要么是一个循环加载float数量，要么是一个循环加载字节(即一个memcpy)。假设它加载了float数量。因此，该循环的第一次迭代将请求S00、S10、S20、S30 .穿过经线。这仅仅是每个线程一个float，所以在整个过程中它是128个字节，所以理论上它可以在一个事务中被服务。但是S00，S20，S40 .在同一家银行，S10，S30，S50 .都在同一家银行，所以我们将有一个16路银行冲突，正如你在你的问题中所预测的那样。

现在，假设每个线程的存储如下所示：

const int n = 16;
struct foo{
    float4 data[n/4];
};

我们的“装载”操作如下：

sh_data[tid] = ...

同样，编译器不能在每个线程的单个指令中加载16x4字节。所以它必须打破转移。如果在本例中，我们可以引导编译器加载每个线程的float4，那么第一个循环迭代将尝试加载线程0的S00-S03，线程1的S10-S13等等。这个加载现在是512个字节，而不是128个字节。因此，该循环迭代的单个翘曲读取指令将被分解为4个共享内存事务。第一笔交易将包括S00-S03、S10-S13、S20-S23、S30-S33、S40-S43、S50-S53、S60-S63、S70-S73的负载请求。如果我们在上面的图表上查看这一点，我们就会发现，我们已经将涉及16路每个循环迭代/事务之间的银行冲突的前一个序列转换为一个新的序列，该序列涉及每个循环迭代/事务中的4路银行冲突。我们有相同数量的128字节事务要共享内存，但是现在每个事务都承受着4路银行冲突的负担，而不是16路银行冲突。