1

次のようにストライド メモリ アクセスを実行するカーネルがあるとします。

__global__ void strideExample (float *outputData, float *inputData, int stride=2) 
{
        int index = (blockIdx.x * blockDim.x + threadIdx.x) * stride;
        outputData[index] = inputData[index]; 
}

トランザクションに含まれる要素の半分が使用されない (無駄な帯域幅になる) ため、ストライド サイズ 2 でアクセスすると、ロード/ストア効率が 50% になることを理解しています。より大きなストライド サイズのロード/ストア効率を計算するにはどうすればよいでしょうか? 前もって感謝します！

4

1 に答える 1

4

一般に：

load efficiency = requested loads / effective loads

requested loadsはソフトウェアが読み取りを要求したeffective loadsバイト数で、 はハードウェアが実際に読み取らなければならなかったバイト数です。店舗にも同じ式が適用されます。

完全に結合されたアクセスの効率は 1 です。

あなたのコードは正確に(blockIdx.x * blockDim.x + threadIdx.x) * sizeof(float)バイトを要求します。が正しくアラインされていると仮定するとoutputData( によって返されるポインタと同様)、ハードウェアは、トランザクション サイズ (SM/L1 の場合は 128 バイト、L1/L2 の場合は 32 バイト) に切り上げられたバイトcudaMallocを読み取る必要があります。(blockIdx.x * blockDim.x + threadIdx.x) * sizeof(float) * stride

ブロック サイズが十分に大きいと仮定すると、トランザクション サイズへの丸めは無視できるようになり、式を に単純化できます1 / stride。この場合、約 16.7% の負荷効率が得られます。

于 2016-11-17T06:18:08.487 に答える