0

私は初心者のCUDAプログラマーです。私は最近、より少ない占有率でより良いパフォーマンスを達成することについてさらに学びました. コード スニペットを次に示します。リプレイ オーバーヘッドと命令レベルの並列処理に関するいくつかのことを理解するための助けが必要です

__global__ void myKernel(double *d_dst, double *d_a1, double *d_a2, size_t SIZE) 
{

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

    d_dst[tId]            = d_a1[tId] * d_a2[tId];
    d_dst[tId + SIZE]     = d_a1[tId + SIZE] * d_a2[tId + SIZE];
    d_dst[tId + SIZE * 2] = d_a1[tId + SIZE * 2] * d_a2[tId + SIZE * 2];
    d_dst[tId + SIZE * 3] = d_a1[tId + SIZE * 3] * d_a2[tId + SIZE * 3];
}

これは単純なカーネルで、単純に 2 つの 2D 配列を乗算して 3 つ目の 2D 配列を形成し (論理的な観点から)、これらの配列はすべてフラットな 1D 配列としてデバイス メモリに配置されます。

以下に、別のコード スニペットを示します。

void doCompute() {

    double *h_a1;
    double *h_a2;

    size_t SIZE = pow(31, 3) + 1;

   // Imagine h_a1, h_a2 as 2D arrays
   // with 4 rows and SIZE Columns
   // For convenience created as 1D arrays 

    h_a1 = (double *) malloc(SIZE * 4 * sizeof(double));
    h_a2 = (double *) malloc(SIZE * 4 * sizeof(double));

    memset(h_a1, 5.0, SIZE * 4 * sizeof(double));
    memset(h_a2, 5.0, SIZE * 4 * sizeof(double));

    double *d_dst;
    double *d_a1;
    double *d_a2;

    cudaMalloc(&d_dst, SIZE * 4 * sizeof(double));
    cudaMalloc(&d_a1,  SIZE * 4 * sizeof(double));
    cudaMalloc(&d_a2,  SIZE * 4 * sizeof(double));

    cudaMemcpy(d_a1, h_a1, SIZE * 4 * sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(d_a2, h_a2, SIZE * 4 * sizeof(double), cudaMemcpyHostToDevice);

    int BLOC_SIZE = 32;
    int GRID_SIZE = (SIZE + BLOC_SIZE - 1) / BLOC_SIZE;

    myKernel <<< GRID_SIZE, BLOC_SIZE >>> (d_dst, d_a1, d_a2, SIZE);
}

Q1) 合体メモリ アクセス パターンを破っていますか?

Q2) メモリへのアクセス、カーネル内でのコーディング方法も、命令レベルの並列処理の例であると言えますか? はいの場合、ILP2 または ILP4 を使用していますか? なぜ？

Q3) 私がしていることはすべて正しいのに、nvvp プロファイラーが次のメッセージを表示するのはなぜですか?

Total Replay Overhead: 4.6%
Global Cache Replay Overhead: 30.3%

どうすればそれらを減らしたり修正したりできますか?

乾杯、

4

1 に答える 1

1

コンパイラには、ILP の悪用の可能性がある命令をスケジュールする機能が制限されています。GPU 自体にも ILP 機能が必要であり、その範囲は GPU の世代によって異なります。はい、利用できないリソースはワープを停止させる可能性があります。典型的なものは、メモリから必要なデータです。あなたが求めているリプレイ量の定義はここにあります。

したがって、たとえば、グローバル キャッシュ リプレイ オーバーヘッドはキャッシュ ミスによってトリガーされ、コードにはいくつかのキャッシュ ミスが発生します。100% の合体アクセスと (ほぼ) 100% の帯域幅使用効率がある場合でも、キャッシュ ミスが発生する可能性があります。

于 2013-07-11T19:25:50.077 に答える