c - グローバルメモリから共有メモリにデータをロードする際のバンク競合を回避する方法

Question

問題には、計算機能 1.3 GPU のグローバルメモリに格納されている unsigned char 配列へのストライドアクセスが含まれます。グローバルメモリの合体要件を回避するために、スレッドはグローバルメモリに順次アクセスし、次の例では 2 つのメモリトランザクションのみを使用して配列を共有メモリにコピーします。

#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>

__global__ void kernel ( unsigned char *d_text, unsigned char *d_out ) {

    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    extern __shared__ unsigned char s_array[];

    uint4 *uint4_text = ( uint4 * ) d_text;
    uint4 var;

    //memory transaction
    var = uint4_text[0];

    uchar4 c0 = *reinterpret_cast<uchar4 *>(&var.x);
    uchar4 c4 = *reinterpret_cast<uchar4 *>(&var.y);
    uchar4 c8 = *reinterpret_cast<uchar4 *>(&var.z);
    uchar4 c12 = *reinterpret_cast<uchar4 *>(&var.w);

    s_array[threadIdx.x*16 + 0] = c0.x;
    s_array[threadIdx.x*16 + 1] = c0.y;
    s_array[threadIdx.x*16 + 2] = c0.z;
    s_array[threadIdx.x*16 + 3] = c0.w;

    s_array[threadIdx.x*16 + 4] = c4.x;
    s_array[threadIdx.x*16 + 5] = c4.y;
    s_array[threadIdx.x*16 + 6] = c4.z;
    s_array[threadIdx.x*16 + 7] = c4.w;

    s_array[threadIdx.x*16 + 8] = c8.x;
    s_array[threadIdx.x*16 + 9] = c8.y;
    s_array[threadIdx.x*16 + 10] = c8.z;
    s_array[threadIdx.x*16 + 11] = c8.w;

    s_array[threadIdx.x*16 + 12] = c12.x;
    s_array[threadIdx.x*16 + 13] = c12.y;
    s_array[threadIdx.x*16 + 14] = c12.z;
    s_array[threadIdx.x*16 + 15] = c12.w;

    d_out[idx] = s_array[threadIdx.x*16];
}

int main ( void ) {

    unsigned char *d_text, *d_out;

    unsigned char *h_out = ( unsigned char * ) malloc ( 32 * sizeof ( unsigned char ) );
    unsigned char *h_text = ( unsigned char * ) malloc ( 32 * sizeof ( unsigned char ) );

    int i;

    for ( i = 0; i < 32; i++ )
        h_text[i] = 65 + i;

    cudaMalloc ( ( void** ) &d_text, 32 * sizeof ( unsigned char ) );
    cudaMalloc ( ( void** ) &d_out, 32 * sizeof ( unsigned char ) );

    cudaMemcpy ( d_text, h_text, 32 * sizeof ( unsigned char ), cudaMemcpyHostToDevice );

    kernel<<<1,32,16128>>>(d_text, d_out );

    cudaMemcpy ( h_out, d_out, 32 * sizeof ( unsigned char ), cudaMemcpyDeviceToHost );

    for ( i = 0; i < 32; i++ )
        printf("%c\n", h_out[i]);

    return 0;
}

問題は、データを共有メモリにコピーするときにバンクの競合が発生し (nvprof によって報告された上記の例では 384 の競合)、スレッドのシリアル化されたアクセスにつながることです。

共有メモリは、同じハーフワープの 16 のスレッドを同時に処理するために、16 (新しいデバイスアーキテクチャでは 32) の 32 ビットバンクに分割されます。データはバンク間でインターリーブされ、i 番目の 32 ビットワードは常に i % 16 - 1 共有メモリバンクに格納されます。

各スレッドは 1 回のメモリトランザクションで 16 バイトを読み取るため、文字はストライド方式で共有メモリに格納されます。これにより、スレッド 0、4、8、12 の間で競合が発生します。1、5、9、13; 2、6、10、14; 同じハーフワープの 3、7、11、15。バンクの競合をなくす単純な方法は、if/else 分岐を使用して、次のようなラウンドロビン方式でデータを共有メモリに格納することですが、深刻なスレッドの分岐が発生します。

int tid16 = threadIdx.x % 16;

if ( tid16 < 4 ) {

    s_array[threadIdx.x * 16 + 0] = c0.x;
    s_array[threadIdx.x * 16 + 1] = c0.y;
    s_array[threadIdx.x * 16 + 2] = c0.z;
    s_array[threadIdx.x * 16 + 3] = c0.w;

    s_array[threadIdx.x * 16 + 4] = c4.x;
    s_array[threadIdx.x * 16 + 5] = c4.y;
    s_array[threadIdx.x * 16 + 6] = c4.z;
    s_array[threadIdx.x * 16 + 7] = c4.w;

    s_array[threadIdx.x * 16 + 8] = c8.x;
    s_array[threadIdx.x * 16 + 9] = c8.y;
    s_array[threadIdx.x * 16 + 10] = c8.z;
    s_array[threadIdx.x * 16 + 11] = c8.w;

    s_array[threadIdx.x * 16 + 12] = c12.x;
    s_array[threadIdx.x * 16 + 13] = c12.y;
    s_array[threadIdx.x * 16 + 14] = c12.z;
    s_array[threadIdx.x * 16 + 15] = c12.w;

} else if ( tid16 < 8 ) {

    s_array[threadIdx.x * 16 + 4] = c4.x;
    s_array[threadIdx.x * 16 + 5] = c4.y;
    s_array[threadIdx.x * 16 + 6] = c4.z;
    s_array[threadIdx.x * 16 + 7] = c4.w;

    s_array[threadIdx.x * 16 + 8] = c8.x;
    s_array[threadIdx.x * 16 + 9] = c8.y;
    s_array[threadIdx.x * 16 + 10] = c8.z;
    s_array[threadIdx.x * 16 + 11] = c8.w;

    s_array[threadIdx.x * 16 + 12] = c12.x;
    s_array[threadIdx.x * 16 + 13] = c12.y;
    s_array[threadIdx.x * 16 + 14] = c12.z;
    s_array[threadIdx.x * 16 + 15] = c12.w;

    s_array[threadIdx.x * 16 + 0] = c0.x;
    s_array[threadIdx.x * 16 + 1] = c0.y;
    s_array[threadIdx.x * 16 + 2] = c0.z;
    s_array[threadIdx.x * 16 + 3] = c0.w;

} else if ( tid16 < 12 ) {

    s_array[threadIdx.x * 16 + 8] = c8.x;
    s_array[threadIdx.x * 16 + 9] = c8.y;
    s_array[threadIdx.x * 16 + 10] = c8.z;
    s_array[threadIdx.x * 16 + 11] = c8.w;

    s_array[threadIdx.x * 16 + 12] = c12.x;
    s_array[threadIdx.x * 16 + 13] = c12.y;
    s_array[threadIdx.x * 16 + 14] = c12.z;
    s_array[threadIdx.x * 16 + 15] = c12.w;

    s_array[threadIdx.x * 16 + 0] = c0.x;
    s_array[threadIdx.x * 16 + 1] = c0.y;
    s_array[threadIdx.x * 16 + 2] = c0.z;
    s_array[threadIdx.x * 16 + 3] = c0.w;

    s_array[threadIdx.x * 16 + 4] = c4.x;
    s_array[threadIdx.x * 16 + 5] = c4.y;
    s_array[threadIdx.x * 16 + 6] = c4.z;
    s_array[threadIdx.x * 16 + 7] = c4.w;

} else {

    s_array[threadIdx.x * 16 + 12] = c12.x;
    s_array[threadIdx.x * 16 + 13] = c12.y;
    s_array[threadIdx.x * 16 + 14] = c12.z;
    s_array[threadIdx.x * 16 + 15] = c12.w;

    s_array[threadIdx.x * 16 + 0] = c0.x;
    s_array[threadIdx.x * 16 + 1] = c0.y;
    s_array[threadIdx.x * 16 + 2] = c0.z;
    s_array[threadIdx.x * 16 + 3] = c0.w;

    s_array[threadIdx.x * 16 + 4] = c4.x;
    s_array[threadIdx.x * 16 + 5] = c4.y;
    s_array[threadIdx.x * 16 + 6] = c4.z;
    s_array[threadIdx.x * 16 + 7] = c4.w;

    s_array[threadIdx.x * 16 + 8] = c8.x;
    s_array[threadIdx.x * 16 + 9] = c8.y;
    s_array[threadIdx.x * 16 + 10] = c8.z;
    s_array[threadIdx.x * 16 + 11] = c8.w;
}

誰でもこれに対するより良い解決策を思いつくことができますか? SDK のリダクションの例については既に調査しましたが、それが私の問題に適用できるかどうかはわかりません。

score 2 · Accepted Answer

コードが許可されると銀行の競合が発生しますが、それが遅いという意味ではありません。

コンピューティング機能1.3GPUでは、双方向のバンク競合がある共有メモリトランザクションは、バンク競合がない場合よりも2サイクル多くかかります。2つのサイクルでは、銀行の競合を回避するために1つの命令を実行することさえできません。4ウェイバンクコンフリクトは、コンフリクトフリーアクセスと比較してさらに6サイクルを使用します。これは、1つの追加のコンフリクトフリー共有メモリアクセスを実行するのに十分です。

あなたの場合、コードはグローバルメモリ帯域幅（およびレイテンシーは数百サイクル、つまりここで説明している2..6サイクルより2桁大きい）によって制限される可能性が非常に高くなります。したがって、SMがグローバルメモリからのデータを待機しているだけの場合は、おそらく十分なスペアサイクルが利用可能になります。銀行の競合は、コードの速度をまったく低下させることなく、これらのサイクルを使用できます。

コンパイラが.x、.y、.z、および.wの4つのバイト単位のストアを単一の32ビットアクセスにマージしていることを確認することがはるかに重要になります。を使用してコンパイルされたコードをcuobjdump -sass調べて、それが当てはまるかどうかを確認します。そうでない場合は、代わりに単語転送を使用するようにOtterのアドバイスに従ってください。

カーネル内からの読み取りのみをd_text行い、カーネル内からの書き込みを行わない場合は、テクスチャを使用することもできます。これは、バンクの競合があるカーネルよりも低速ですが、全体的な速度を向上させる他の利点を提供する可能性があります（たとえば、グローバルメモリ内のデータの適切な配置を保証するものではありません）。

一方、代替の銀行競合のないコードは、高速の256バイトのグローバルメモリを4つの64ビットトランザクションに分割します。これは、効果がはるかに低く、実行中のメモリトランザクションの最大数をオーバーフローする可能性があるため、グローバルメモリレイテンシの完全な400から数千サイクル。
これを回避するには、最初に256バイト幅の読み取りを使用してレジスタに転送し、次にデータをレジスタから共有メモリにバンク競合のない方法で移動する必要があります。それでも、register-> shmem moveのコードだけで、回避しようとしていた6サイクルよりもはるかに多くの時間がかかります。

score 1 · Accepted Answer

とにかく、DWORD コピーはバイトごとのコピーよりも高速だと思います。あなたの例の代わりにこれを試してください：

for(int i = 0; i < 4; i++)
{
    ((int*)s_array)[4 * threadIdx.x + i] = ((int*)d_text)[i];
}

c - グローバルメモリから共有メモリにデータをロードする際のバンク競合を回避する方法

3 に答える 3

Related

Reference