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

您实际上拥有一组结构 (AoS)。这对于 GPU 编程来说是非常糟糕的。您可以使用标准的 AoS->SoA 数据重组方法来修复访问，以便相邻线程访问相邻元素（这将防止库冲突）。
以更大的块加载数据，例如重新组织您的结构，使其可以表示 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];
};

而我们的 "load" 操作如下所示：

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];
};

而我们的 "load" 操作如下所示：

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 路存储体冲突 iteration/transaction 为一个新序列，每个循环涉及 4 路存储体冲突 iteration/transaction。我们有相同数量的 128 字节事务进入共享内存，但每个事务现在都承受着 4 路存储体冲突而不是 16 路存储体冲突。

CUDA：最小化大数据类型的银行冲突

CUDA: minimize bank conflict for large data type

c++

memory

cuda

gpu

nvidia