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

CUDA: minimize bank conflict for large data type

本文关键字：数据类型 冲突 大型 最大限度 CUDA      更新时间：2023-10-16

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_data[tid]时，如果我做对了，就会有8个（或16个，32个）银行冲突。

如果是这样的话，有什么技术可以最大限度地减少银行冲突吗？

感谢

  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。。。穿过经线。每个线程只有1个float，因此在整个warp中它是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字节。因此，该循环迭代的单个warp读取指令将被分解为4个共享内存事务。第一个事务将包括S00-S03、S10-S13、S20-S23、S30-S33、S40-S43、S50-S53、S60-S63、S70-S73的加载请求。如果我们在上图中对此进行检查，我们会发现我们已经将之前涉及每个循环迭代/事务的16路银行冲突的序列转换为涉及每个循环循环迭代/交易的4路银行冲突。我们有相同数量的128字节事务进入共享内存，但每个事务现在都有一个4路银行冲突，而不是16路银行冲突。

相关文章：