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」はグローバル メモリ バッファの幅 (float4 の数) に設定され、「d」は 8 (コピーされた float4 の数) に設定されます。

そのような構成とメモリのさらなる使用は、バンクの競合につながる可能性がありますか、またはブロードキャストが適用されますか? これは、スレッドが 8 つではなく 5 つの float4 のみをコピーする場合にも当てはまりますか?

MK

PS同じトピックはこちら

4

1 に答える 1

1

プリフェッチ フェーズ中に、バンクの競合が発生します。たとえば、ID が 0、4、8 threadIdx.x + threadIdx.y * blockDim.x、... 28 の最初のワープ内のスレッドは、同じバンクにアクセスします。スレッド (0,0) とスレッド (4,0) がi等しい 0 アクセスs_dst[0]でありs_dst[32]、同じバンクに属していることがわかります。

さらに使用中にバンクの競合が発生した場合は、アクセスするスキームによって異なりますs_dst。

ブロードキャスト メカニズムは、スレッドが同時に同じアドレスを読み取る場合にのみ適用されます。

バンク競合が発生する回数は、 の値によって異なりますd。d mod 32 == 1競合が発生しない場合。

編集: 私見プリフェッチ フェーズでバンクの競合を回避する最善の方法は、特にd変更されている場合は、作業をワープ間で均等に分割することです。n値を共有メモリにプリフェッチする必要があるとします。w_idワープの ID であり、ワープl_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 に答える