1

GPUにFDTD方程式を実装しようとしていました。私は当初、グローバル メモリを使用するカーネルを実装していました。記憶の合体はそれほど素晴らしいものではありませんでした。したがって、共有メモリを使用して値をロードする別のカーネルを実装しました。私はのグリッドに取り組んでい1024x1024ます。

コードは以下です

__global__ void update_Hx(float *Hx, float *Ez, float *coef1, float* coef2){
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;
    __shared__ float  Ez_shared[BLOCKSIZE_HX][BLOCKSIZE_HY + 1];
    /*int top = offset + x_index_dim;*/
    if(threadIdx.y == (blockDim.y - 1)){
        Ez_shared[threadIdx.x][threadIdx.y] = Ez[offset];
        Ez_shared[threadIdx.x][threadIdx.y + 1] = Ez[offset + x_index_dim];
   }
    else{
        Ez_shared[threadIdx.x][threadIdx.y] = Ez[offset];
    }
}

定数BLOCKSIZE_HX=16およびBLOCKSIZE_HY= 16。

ビジュアル プロファイラーを実行すると、メモリが結合されていないと表示されます。

編集: 2.1 の cuda 計算機能を備えた GT 520 グラフィック カードを使用しています。My Global L2 transactions / Access =7.5つまり、回線の実行の 245 760ための L2 トランザクション があります32768Ez_shared[threadIdx.x][threadIdx.y] = Ez[offset];

Global memory load efficiencyです50%。

Global memory load efficiency=100 * gld_requested_throughput/ gld_throughput

スレッドが 16 個の連続した値を参照しているにもかかわらず、メモリ アクセスが非常に多い理由がわかりません。誰かが私が間違っていることを指摘できますか?

編集：すべての助けをありがとう。

4

1 に答える 1

1

ここで問題になるのは、メモリ アクセス パターンです。16 フ​​ロートの連続した領域、つまり 64 バイトにアクセスしているため、(L1 と L2 の両方で) 50% の効率しか得られませんが、L1 トランザクション サイズは 128 バイトです。これは、要求された 64 バイトごとに、128 バイトを L1 にロードする必要があることを意味します (結果的に L2 にもロードする必要があります)。

また、共有メモリ バンクの競合の問題もありますが、現在、グローバル メモリのロード効率に悪影響はありません。

負荷効率の問題は、いくつかの方法で解決できます。最も簡単な方法は、x 次元のブロック サイズを 32 に変更することです。それができない場合は、グローバル メモリ データ レイアウトを変更して、2 つの連続する blockIdx.y ([0, 1]、[2,3] など) がそれぞれ. ) 値は、連続メモリ ブロックにマップされます。それでもオプションではなく、グローバル データを一度だけロードする必要がある場合は、キャッシュされていないグローバル メモリ ロードを使用して L1 をバイパスできます。L2 は 32 バイトのトランザクションを使用するため、64 バイトが 2 つの L2 にロードされるオーバーヘッドのないトランザクション。

于 2013-02-12T10:59:24.887 に答える