1

在我的程序中，我使用共享内存来预取数据。一个 2D 线程块，尺寸为 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'设置为全局内存缓冲区的宽度（以float4s的数量为单位），'d'设置为8（复制的float4s的数量）。

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

MK

PS 相同的主题在这里

4

1 回答 1

1

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

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

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

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

编辑： 恕我直言，在预取阶段避免银行冲突的最佳方法，特别是如果d正在改变，是在经纱之间平均分配工作。假设您需要将值预取n到共享内存，w_id是 warp 的 ID 和 warpl_id中的线程 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。

于 2013-02-27T11:58:32.707 回答