0

nvidia gtx980カードで、トランザクション数を観察することにより、結合されたメモリアクセスをテストするための簡単なカーネルを作成しました。カーネルは、

__global__
void copy_coalesced(float * d_in, float * d_out)
{
    int tid = threadIdx.x + blockIdx.x*blockDim.x;

    d_out[tid] = d_in[tid];
}

次のカーネル構成でこれを実行すると

#define BLOCKSIZE 32   

int data_size  = 10240;                  //always a multiply of the BLOCKSIZE
int gridSize   = data_size / BLOCKSIZE;

copy_coalesced<<<gridSize, BLOCKSIZE>>>(d_in, d_out);

カーネル内のデータ アクセスは完全に統合されており、データ型は float (4 バイト) であるため、予想されるロード/ストア トランザクションの数は次のようになります。

ロード トランザクション サイズ = 32 バイト

トランザクションごとにロードできる float の数 = 32 バイト / 4 バイト = 8

10240 のデータをロードするのに必要なトランザクション数 = 10240/8 = 1280 トランザクション

データの書き込みにも同じ量のトランザクションが予想されます。

しかし、nvprof メトリックを観察すると、以下の結果が得られました

gld_transactions    2560
gst_transactions    1280

gld_transactions_per_request    8.0
gst_transactions_per_request    4.0

データのロードに必要なトランザクションの 2 倍が必要な理由がわかりません。しかし、ロード/ストアの効率に関しては、両方のメトリックが 100% を示します

ここで何が欠けていますか？

4

1 に答える 1