用于预取的共享内存配置

Question

在我的程序中，我使用共享内存来预取数据。一个二维线程块，尺寸为 8 x 4 (32)，获得 8 * 4 * 8 * sizeof(float4) 字节的共享内存。每个线程循环复制8个float4：

inline __device__ void pack(const float4 *g_src, float4 *s_dst, const unsigned int w, const unsigned int d) {
    uint2 indx = { blockIdx.x * blockDim.x + threadIdx.x, blockIdx.y * blockDim.y + threadIdx.y };
    uint2 sindx = { threadIdx.x, threadIdx.y };
    int i;

    for (i = 0; i < d; ++i) s_dst[(sindx.y * blockDim.x + sindx.x) * d + i] = g_src[(w * indx.y + indx.x) * d + i];
} 

其中“w”设置为全局内存缓冲区的宽度（以 float4 为单位），“d”设置为 8（复制的 float4 数量）。

这样的配置和内存的进一步使用会导致bank冲突，或者广播会被应用吗？当线程只复制时也会出现这种情况，比如 5 个 float4s，而不是 8 个？

MK

附： 同题here

Answer 1

在预取阶段会发生银行冲突。例如。第一个经线中的线程，ID 为（计算为threadIdx.x + threadIdx.y * blockDim.x）0、4、8、... 28 访问同一个银行。您可以将其视为线程 (0,0) 和线程 (4,0) for i 等于 0 访问 s_dst[0] 和 s_dst[32] 属于同一银行。

如果在进一步使用过程中发生银行冲突，取决于您将访问s_dst 的方案。

广播机制仅在线程同时读取同一地址时应用。

发生多少银行冲突取决于d 的值。如果d mod 32 == 1 不会有任何冲突。

编辑： 恕我直言，在预取阶段避免银行冲突的最佳方法，特别是如果d 正在改变，是在经纱之间平均分配工作。假设您需要将 n 值预取到共享内存，w_id 是 warp 的 ID，l_id 是 warp 内的线程 ID（从 0 到 31）。预取应该是这样的：

for(int i = l_id + w_id*WARP_SIZE; i < n; i += WARP_SIZE*COUNT_OF_WARPS_IN_BLOCK)
{
    s_dst[i] = ...;
}

但这仅有助于避免预取期间的银行冲突。正如我已经说过的，在进一步使用过程中避免冲突取决于您将访问s_dst 的方案。