NVIDIA GPUで命令レベルの並列処理（ILP）とアウトオブオーダー実行

Question

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

まずは、高価なハードウェアを搭載していないと思います。ただし、CUDA progamming guideを読むときは、パフォーマンス向上のために命令レベル並列化（ILP）を使用することをお勧めします。

ILPは、アウトオブオーダー実行をサポートするハードウェアが利用できる機能ではありませんか？または、NVIDIAのILPは単にコンパイラレベルの命令の並べ替えを意味するため、実行時に命令の順序は変わりません。言い換えれば、コンパイラやプログラマだけが、実行時にインオーダー実行によってILPを実現できるような方法で命令の順序を並べ替える必要がありますか？

Answer 1

パイプライン処理は一般的なILP手法であり、確かにNVidiaのGPUに実装されています。あなたは、パイプライン処理がアウト・オブ・オーダー実行に依存していないことに同意すると思います。 さらに、NVidia GPUには、計算機能2.0以降の複数のワープスケジューラがあります（2または4）。あなたのコードが2つ（またはそれ以上）の連続した独立した命令をスレッド内に持つ場合（または何らかの形でコンパイラがそれを並べ替える）、このILPをスケジューラからも悪用します。

ここでは、2ワイドワープスケジューラ+パイプライニングがどのように連携するかについて説明しています。 How do nVIDIA CC 2.1 GPU warp schedulers issue 2 instructions at a time for a warp?

また、GTC 2010のVasily Volkov氏のプレゼンテーションでも、ILPがどのようにCUDAコードのパフォーマンスを向上させるかが実験的に分かりました。 http://www.cs.berkeley.edu/~volkov/volkov10-GTC.pdf

GPUのアウトオブオーダー実行に関して、私はそうは思わない。ハードウェアの命令の並べ替え、投機的な実行はすべてあなたが知っているように、SMごとに実装するには高価です。スレッドレベルの並列処理は、アウトオブオーダー実行の欠如のギャップを埋めることができます。真の依存関係に遭遇すると、いくつかの他のワープが蹴り込んでパイプを埋めることができます。

Answer 2

以下のコードは、命令レベル並列性（ILP）の例を示しています。

この例の__global__関数は、2つの配列間の代入を単に実行します。 ILP=1の場合、配列要素の数と同じ数のスレッドがあり、各スレッドは1つの割り当てを実行します。Nこれとは逆に、ケースILP=2の場合、要素を処理するスレッドがN/2個あります。一般に、ケースILP=kの場合、エレメントを処理するスレッドがN/k個あります。コードに加え

は、私はまた、タイミングを報告しています下に、NとILPの異なる値について、NVIDIA GT920M（ケプラーアーキテクチャ）上で行いました。それは見ることができるように：

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

私はマックスウェルで同じコードをテストし、同じ動作を観察しました（ILPに対してパフォーマンスに変化はありません）。ILPに対するパフォーマンスの変更については、The efficiency and performance of ILP for the NVIDIA Kepler architectureの回答を参照して、フェルミアーキテクチャのテストも報告してください。

メモリ速度は以下の式で計算された：

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

4.f * N * numITER 

がリードの数であり、OR書き込み、

2.f * 4.f * N * numITER 

がリードの数でありますAND書き込み、

timeTotal * 1e-3 

は、seconds（timeTotalはms）の時刻です。

CODE

// --- 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; 

} 

PERFORMANCE

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