8

NVIDIA GPU 是否支持乱序执行？

我的第一个猜测是它们不包含如此昂贵的硬件。但是，在阅读 时CUDA progamming guide，该指南建议使用指令级并行 (ILP) 来提高性能。

ILP 不是支持乱序执行的硬件可以利用的功能吗？或者 NVIDIA 的 ILP 仅仅意味着编译器级别的指令重新排序，因此它的顺序在运行时仍然是固定的。换句话说，只有编译器和/或程序员必须以这样一种方式安排指令的顺序，以便可以在运行时通过按顺序执行来实现 ILP？

4

2 回答 2

6

流水线是一种常见的 ILP 技术，并且肯定会在 NVidia 的 GPU 上实现。我猜你同意流水线不依赖于乱序执行。此外，NVidia GPU 具有多个计算能力 2.0 及更高版本（2 或 4）的 warp 调度程序。如果您的代码在线程中有 2 个（或更多）连续且独立的指令（或编译器以某种方式重新排序），您也可以从调度程序中利用此 ILP。

这是一个关于 2-wide warp scheduler + pipelining 如何一起工作的很好解释的问题。 nVIDIA CC 2.1 GPU warp 调度程序如何一次发出 2 条指令进行 warp？

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

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

于 2013-07-26T21:59:47.483 回答
1

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

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

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

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

关于第 2 点，我还在 Maxwell 上测试了相同的代码，并观察到了相同的行为（性能没有变化ILP）。有关性能的变化ILP，请参阅NVIDIA Kepler 架构的 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 回答