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

如果是这种情况，是否有任何技术可以最大限度地减少银行冲突？

谢谢

Answer 1

1. 您实际上拥有一个结构数组 (AoS)。这对于 GPU 编程来说无疑是不利的。您可以使用标准 AoS->SoA 数据重组方法来修复访问，以便相邻线程访问相邻元素（这将防止存储库冲突）。

2. 以更大的块加载数据，例如重新组织您的结构，使其可以表示 4 个 float4 数量而不是 16 float 数量。编译器可能能够将加载组织成float4 加载，这将减少存储库冲突。如果你真的需要这两种访问方法，你甚至可以在结构中使用联合。

根据 cmets 中的问题，让我们用图表说明第二种情况。对于每个结构的 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] = ...

编译器cannot load 16x4 bytes in a single instruction per thread 因此它将上述加载操作分解为一系列请求。 我认为这个序列要么是循环加载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 路银行冲突。