用于预取的共享内存配置

Question

在我的程序中，我使用共享内存来预取数据。一个二维线程块，维度为8乘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];
} 

其中CCD_ 3被设置为全局存储器缓冲器的宽度(以float4s的数量为单位(并且CCD_ 5被设置为8(复制的float4s的数量(。

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

MK-

第页。SNvidia论坛中的相同主题

Answer 1

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

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

只有当线程同时读取相同地址时，才应用广播机制。

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

编辑：IMHO在预取阶段避免存储体冲突的最佳方法，特别是在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的方案。