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）的 warp 调度程序。如果您的代码在线程中有 2 个（或更多）连续且独立的指令（或编译器以某种方式重新排序），您也可以从调度程序中利用此 ILP。

这是一个很好解释的问题，关于 2-wide warp 调度程序 + 流水线如何协同工作。 How do nVIDIA CC 2.1 GPU warp schedulers issue 2 instructions at a time for a warp?

还可以查看 Vasily Volkov 在 GTC 2010 上的演示。他通过实验发现了 ILP 如何提高 CUDA 代码性能。 http://www.cs.berkeley.edu/~volkov/volkov10-GTC.pdf

就 GPU 上的乱序执行而言，我不这么认为。如您所知，硬件指令重新排序、推测执行所有这些东西对于每个 SM 来说都太昂贵了。而线程级并行可以弥补乱序执行的不足。当遇到真正的依赖关系时，其他一些扭曲可以启动并填充管道。

Answer 2

下面的代码报告了指令级并行 (ILP) 的示例。

示例中的__global__ 函数只是在两个数组之间执行赋值。对于ILP=1 的情况，我们的线程数与数组元素N 的数量一样多，因此每个线程执行一次赋值。与此相反，对于ILP=2 的情况，我们有许多N/2 线程，每个线程处理2 元素。一般来说，对于ILP=k的情况，我们有多个N/k线程，每个线程处理k元素。

除了代码之外，我还报告了在NVIDIA GT920M（Kepler 架构）上针对N 和ILP 的不同值执行的时间。可以看出：

1. 对于较大的N 值，内存带宽接近GT920M 卡的最大值，即14.4GB/s；
2. 对于任何固定的N，更改ILP 的值不会改变性能。

关于第 2 点，我还在 Maxwell 上测试了相同的代码，并观察到了相同的行为（针对 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

是读或写的次数，

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