4

私は自分のCUDAアプリケーションでビジュアルプロファイラーを実行しました。データが大きすぎる場合、アプリケーションは単一のカーネルを複数回呼び出します。このカーネルには分岐がありません。

プロファイラーは、83.6 ％の高い命令再生オーバーヘッドと83.5％の高いグローバルメモリ命令再生オーバーヘッドを報告します。

カーネルの一般的な外観は次のとおりです。

// Decryption kernel
__global__ void dev_decrypt(uint8_t *in_blk, uint8_t *out_blk){

    __shared__ volatile word sdata[256];
    register uint32_t data;

    // Thread ID
#define xID (threadIdx.x + blockIdx.x * blockDim.x)
#define yID (threadIdx.y + blockIdx.y * blockDim.y)
    uint32_t tid = xID + yID * blockDim.x * gridDim.x;
#undef xID
#undef yID

    register uint32_t pos4 = tid%4;
    register uint32_t pos256 = tid%256;
    uint32_t blk = pos256&0xFC;

    // Indices
    register uint32_t index0 = blk + (pos4+3)%4;
    register uint32_t index1 = blk + (pos4+2)%4;

    // Read From Global Memory
    b0[pos256] = ((word*)in_blk)[tid+4] ^ dev_key[pos4];

    data  = tab(0,sdata[index0]);
    data ^= tab(1,sdata[index1]);
    sdata[pos256] = data ^ tab2[pos4];

    data  = tab(0,sdata[index0]);
    data ^= tab(1,sdata[index1]);
    sdata[pos256] = data ^ tab2[2*pos4];

    data  = tab(0,sdata[index0]);
    data ^= tab(1,sdata[index1]);
    data ^= tab2[3*pos4];

    ((uint32_t*)out_blk)[tid] = data + ((uint32_t*)in_blk)[tid];
}

ご覧のとおり、ブランチはありません。スレッドは、最初にスレッドID+16バイトに基づいてグローバルメモリから読み取ります。次に、スレッドIDに基づいてグローバルメモリからのデータを使用して操作を実行した後、出力バッファに書き込みます。

このカーネルにこれほど多くのオーバーヘッドがある理由はありますか？

4

1 に答える 1

3

この場合の命令再生のソースは、ワープ内の不均一な一定のメモリアクセスです。コードでtabは、は定数メモリに格納され、スレッドインデックスとデータ格納された共有メモリの組み合わせに従ってインデックスが付けられます。その結果、同じワープ内に不均一なアクセススレッドが表示されます。コンスタントメモリは、ワープ内のすべてのスレッドが同じワードにアクセスする場合を対象としています。その場合、値は1回の操作でコンスタントメモリキャッシュからブロードキャストできます。そうしないと、ワープのシリアル化が発生します。

小さな読み取り専用データセットへの不均一なアクセスが必要な場合は、データを一定のメモリとして保存するよりも、データをテクスチャにバインドする方がよいでしょう。

于 2012-07-23T05:42:02.817 に答える