8

NVIDIA GPU は順不同での実行をサポートしていますか?

私の最初の推測では、そのような高価なハードウェアは含まれていません。ただし、ガイドを読むときはCUDA progamming guide、命令レベルの並列処理 (ILP) を使用してパフォーマンスを向上させることをお勧めします。

ILP は、アウトオブオーダー実行をサポートするハードウェアが利用できる機能ではないでしょうか? または、NVIDIA の ILP は単にコンパイラ レベルの命令の並べ替えを意味するため、その順序は実行時に固定されます。言い換えれば、コンパイラーおよび/またはプログラマーだけが、実行時に順序どおりに実行してILPを達成できるように命令の順序を調整する必要がありますか?

4

2 に答える 2

6

パイプライン処理は一般的な ILP 手法であり、NVidia の GPU に確実に実装されています。パイプライン処理が順不同の実行に依存していないことに同意していると思います。さらに、NVidia GPU には、コンピューティング機能 2.0 以降 (2 または 4) の複数のワープ スケジューラがあります。コードのスレッドに 2 つ (またはそれ以上) の連続した独立した命令がある場合 (またはコンパイラがそのように並べ替える場合)、スケジューラからもこの ILP を利用します。

これは、2 幅のワープ スケジューラとパイプライン処理がどのように連携するかについて、よく説明された質問です。 nVIDIA CC 2.1 GPU ワープ スケジューラは、ワープに対して一度に 2 つの命令をどのように発行しますか?

また、GTC 2010 に関する Vasily Volkov のプレゼンテーションもご覧ください。彼は、ILP が CUDA コードのパフォーマンスを向上させる方法を実験的に発見しました。http://www.cs.berkeley.edu/~volkov/volkov10-GTC.pdf

GPU でのアウトオブオーダー実行に関しては、そうは思いません。ハードウェア命令の並べ替え、投機的実行はすべて、SM ごとに実装するにはコストがかかりすぎることはご承知のとおりです。また、スレッド レベルの並列処理は、順不同で実行できないというギャップを埋めることができます。真の依存関係が発生すると、他のワープが発生してパイプを埋めることができます。

于 2013-07-26T21:59:47.483 に答える
1

次のコードは、命令レベルの並列処理 (ILP) の例を示しています。

この__global__例の関数は、2 つの配列間の割り当てを実行するだけです。の場合ILP=1、配列要素の数と同じ数のNスレッドがあるため、各スレッドは単一の割り当てを実行します。それとは反対に、 の場合、それぞれが要素を処理ILP=2する多数のN/2スレッドがあります。2一般に、 の場合、各要素を処理ILP=kする多数のN/kスレッドがあります。k

コードに加えて、以下では、 (ケプラー アーキテクチャ) で実行された と のさまざまな値のタイミングも報告していNVIDIA GT920Mます。ご覧のとおり：NILP

  1. の値が大きい場合、カードNの最大帯域幅、つまり に近いメモリ帯域幅に達します。GT920M14.4GB/s
  2. 固定されNた の場合、 の値を変更ILPしてもパフォーマンスは変わりません。

ポイント 2. に関しては、Maxwell で同じコードをテストしたところ、同じ動作が観察されました ( に対するパフォーマンスの変化はありませんILP)。に対するパフォーマンスの変化については、NVIDIA Kepler アーキテクチャレポートの ILP の効率とパフォーマンスILPへの回答を参照してください。また、Fermi アーキテクチャでもテストされます。

メモリ速度は、次の式で計算されています。

(2.f * 4.f * N * numITER) / (1e9 * timeTotal * 1e-3)

どこ

4.f * N * numITER

は読み取りまたは書き込みの数です。

2.f * 4.f * N * numITER

は読み取りと書き込みの数です。

timeTotal * 1e-3

seconds(timeTotalは)の時間msです。

コード

// --- GT920m - 14.4 GB/s
//     http://gpuboss.com/gpus/GeForce-GTX-280M-vs-GeForce-920M

#include<stdio.h>
#include<iostream>

#include "Utilities.cuh"
#include "TimingGPU.cuh"

#define BLOCKSIZE    32

#define DEBUG

/****************************************/
/* INSTRUCTION LEVEL PARALLELISM KERNEL */
/****************************************/
__global__ void ILPKernel(const int * __restrict__ d_a, int * __restrict__ d_b, const int ILP, const int N) {

    const int tid = threadIdx.x + blockIdx.x * blockDim.x * ILP;

    if (tid >= N) return;

    for (int j = 0; j < ILP; j++) d_b[tid + j * blockDim.x] = d_a[tid + j * blockDim.x];

}

/********/
/* MAIN */
/********/
int main() {

    //const int N = 8192;
    const int N = 524288 * 32;
    //const int N = 1048576;
    //const int N = 262144;
    //const int N = 2048;

    const int numITER = 100;

    const int ILP = 16;

    TimingGPU timerGPU;

    int *h_a = (int *)malloc(N * sizeof(int));
    int *h_b = (int *)malloc(N * sizeof(int));

    for (int i = 0; i<N; i++) {
        h_a[i] = 2;
        h_b[i] = 1;
    }

    int *d_a; gpuErrchk(cudaMalloc(&d_a, N * sizeof(int)));
    int *d_b; gpuErrchk(cudaMalloc(&d_b, N * sizeof(int)));

    gpuErrchk(cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice));
    gpuErrchk(cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice));

    /**************/
    /* ILP KERNEL */
    /**************/
    float timeTotal = 0.f;
    for (int k = 0; k < numITER; k++) {
        timerGPU.StartCounter();
        ILPKernel << <iDivUp(N / ILP, BLOCKSIZE), BLOCKSIZE >> >(d_a, d_b, ILP, N);
#ifdef DEBUG
        gpuErrchk(cudaPeekAtLastError());
        gpuErrchk(cudaDeviceSynchronize());
#endif
        timeTotal = timeTotal + timerGPU.GetCounter();
    }

    printf("Bandwidth = %f GB / s; Num blocks = %d\n", (2.f * 4.f * N * numITER) / (1e6 * timeTotal), iDivUp(N / ILP, BLOCKSIZE));
    gpuErrchk(cudaMemcpy(h_b, d_b, N * sizeof(int), cudaMemcpyDeviceToHost));
    for (int i = 0; i < N; i++) if (h_a[i] != h_b[i]) { printf("Error at i = %i for kernel0! Host = %i; Device = %i\n", i, h_a[i], h_b[i]); return 1; }

    return 0;

}

パフォーマンス

GT 920M
N = 512  - ILP = 1  - BLOCKSIZE = 512 (1 block  - each block processes 512 elements)  - Bandwidth = 0.092 GB / s

N = 1024 - ILP = 1  - BLOCKSIZE = 512 (2 blocks - each block processes 512 elements)  - Bandwidth = 0.15  GB / s

N = 2048 - ILP = 1  - BLOCKSIZE = 512 (4 blocks - each block processes 512 elements)  - Bandwidth = 0.37  GB / s
N = 2048 - ILP = 2  - BLOCKSIZE = 256 (4 blocks - each block processes 512 elements)  - Bandwidth = 0.36  GB / s
N = 2048 - ILP = 4  - BLOCKSIZE = 128 (4 blocks - each block processes 512 elements)  - Bandwidth = 0.35  GB / s
N = 2048 - ILP = 8  - BLOCKSIZE =  64 (4 blocks - each block processes 512 elements)  - Bandwidth = 0.26  GB / s
N = 2048 - ILP = 16 - BLOCKSIZE =  32 (4 blocks - each block processes 512 elements)  - Bandwidth = 0.31  GB / s

N = 4096 - ILP = 1  - BLOCKSIZE = 512 (8 blocks - each block processes 512 elements)  - Bandwidth = 0.53  GB / s
N = 4096 - ILP = 2  - BLOCKSIZE = 256 (8 blocks - each block processes 512 elements)  - Bandwidth = 0.61  GB / s
N = 4096 - ILP = 4  - BLOCKSIZE = 128 (8 blocks - each block processes 512 elements)  - Bandwidth = 0.74  GB / s
N = 4096 - ILP = 8  - BLOCKSIZE =  64 (8 blocks - each block processes 512 elements)  - Bandwidth = 0.74  GB / s
N = 4096 - ILP = 16 - BLOCKSIZE =  32 (8 blocks - each block processes 512 elements)  - Bandwidth = 0.56  GB / s

N = 8192 - ILP = 1  - BLOCKSIZE = 512 (16 blocks - each block processes 512 elements) - Bandwidth = 1.4  GB / s
N = 8192 - ILP = 2  - BLOCKSIZE = 256 (16 blocks - each block processes 512 elements) - Bandwidth = 1.1  GB / s
N = 8192 - ILP = 4  - BLOCKSIZE = 128 (16 blocks - each block processes 512 elements) - Bandwidth = 1.5  GB / s
N = 8192 - ILP = 8  - BLOCKSIZE =  64 (16 blocks - each block processes 512 elements) - Bandwidth = 1.4  GB / s
N = 8192 - ILP = 16 - BLOCKSIZE =  32 (16 blocks - each block processes 512 elements) - Bandwidth = 1.3  GB / s

...

N = 16777216 - ILP = 1  - BLOCKSIZE = 512 (32768 blocks - each block processes 512 elements) - Bandwidth = 12.9  GB / s
N = 16777216 - ILP = 2  - BLOCKSIZE = 256 (32768 blocks - each block processes 512 elements) - Bandwidth = 12.8  GB / s
N = 16777216 - ILP = 4  - BLOCKSIZE = 128 (32768 blocks - each block processes 512 elements) - Bandwidth = 12.8  GB / s
N = 16777216 - ILP = 8  - BLOCKSIZE =  64 (32768 blocks - each block processes 512 elements) - Bandwidth = 12.7  GB / s
N = 16777216 - ILP = 16 - BLOCKSIZE =  32 (32768 blocks - each block processes 512 elements) - Bandwidth = 12.6  GB / s
于 2017-06-26T21:10:49.070 に答える