cuda - CUDA 中的每计数指令 (IPC) 和指令级并行 (ILP)

Question

我观察到 IPC 下降，因为 ILP 在尝试加速我的加密内核时针对 32-bit int 操作上升。内核由 ADD 和 XOR 操作的长序列相当展开的循环组成，每个 192 的吞吐量应为 160 ops Kepler 上每个周期的内核数(GTX Titan/780)。

我的内核的

IPC 达到了 3.28 的上限。使用 ILP 甚至会降低 IPC。显然 ILP 无法帮助实现我的目标——充分利用管道，所以我写了一些小实验。我把 ILP 4 的代码放在最后。

剖面测量

• 结果是在 GTX Titan 上测得的。
• cubin 检查输出以确保在优化过程中没有指令被删除。
• Executed IPC 和 issued IPC 几乎一样，所以我只列出其中之一。

ADD 指令(XOR 具有相同的行为)

             | ILP 1  | ILP 2   | ILP 4  | ILP 8
--------------------------------------------------
 IPC         |   4.00 |   3.32  |   2.72 |   3.44
--------------------------------------------------
 Issue Slot  | 99.17% | 59.34%  | 48.61% | 61.71%
 Utilization |        |         |        |

• 我希望 ILP 2、4 和 8 会提供更好的性能，但事实并非如此。
• 回想一下整数吞吐量是 160。每个 SM 的 4 warp 调度程序每个周期应该双重发出最多 5 指令，因此 IPC 应该向 5。我该如何解释我观察到的东西？当 IPC = 4 时，为什么 issue slot 的使用率为 99%？

Float/Int ADD 指令组合

如果我修改 ILP 4 的代码来执行两个 int ADD 和两个 float 添加:

IPC: 5.1
Issue slot utilization: 99.12%

奇怪的是，似乎 warp 调度程序在发出 float 操作方面做得更好。

讨论

• 现有文献表明使用 ILP 有助于达到浮点运算的最佳性能。为什么 ILP 不适用于整数？我如何为整数运算执行此操作？
• 我的内核理论上应该对每个候选对象执行 2.25 整数运算。这与我在 cuobjdump 中观察到的一致。有 2^48 个候选者，因此 GTX Titan 上的最小运行时间应该是 2.25 * 2^48/(2688 * 160/192)/876 MHz = 322.75s。这个估计合理吗？
• 我的内核的测量性能是 523s。这确实意味着整数吞吐量仅为 160 * 3.28(测量 IPC)/5(最大 IPC)。

ILP测试代码

__device__ int x[10];
__global__ void test(int flag = 0)
{
    int a = x[0], b = x[1], c = x[2], d = x[3];
    int _a = x[4], _b = x[5], _c = x[6], _d = x[7];

#pragma unroll 128
    for (int i = 0; i < 51200; ++i)
    {
        asm volatile("add.u32 %0, %0, %1;": "+r"(a): "r"(_a));
        asm volatile("add.u32 %0, %0, %1;": "+r"(b): "r"(_b));
        asm volatile("add.u32 %0, %0, %1;": "+r"(c): "r"(_c));
        asm volatile("add.u32 %0, %0, %1;": "+r"(d): "r"(_d));
    }

    int v = a + b + c + d;
    if (flag * v == 1)
        x[0] = v;
}

4个候选人的代码片段

每个候选人需要 9/4 = 2.25 操作。 Cuobjdump 也验证了这一点。

d ^= d2(1, 3);                 // d2 is located in constant memory
s ^= d;
t ^= d2(1, 16);
u ^= d2(1, 17);
v ^= some_const;
flag_s = min(flag_s, s);       // int min has throughput of 160
flag_t = flag_t || (s == t);   // setp.or should be the same
flag_u = flag_u || (s == u);
flag_v = flag_v || (s == v);

Answer 1

我正在提供从未回答列表中删除此问题的答案。

我没有观察到每计数执行的指令数 (IPC) 与指令级并行度 的变化。总体而言，在不知道任何进一步信息但由 OP 自己提供的信息(例如，启动配置)的情况下，很难争论 OP 观察到的效果的原因。

在下面的代码中，我正在考虑使用 float 的示例，尽管我已经使用 int 测试了相同的代码，但没有改变概念结果。该代码使用ILP=1、ILP=2 和 实现循环Multiply Add (MAD) 操作ILP=4。

执行的IPC如下

ILP         IPC            FLOPs
1           3.924          67108864
2           4.323          67108864
4           4.016          67108864

对于 N=8192。代码已使用 CUDA 8.0 编译并在 NVIDIA GT920M 上运行。可以看出，对于 ILP 的不同考虑值，IPC 几乎保持不变。 浮点运算(FLOPs)由代码估计，假设每 MAD 2 FLOPs 与 Visual Profiler 测量的一致。

代码

#include<stdio.h>

#define N_ITERATIONS 8192

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

#define BLOCKSIZE   512

//#define DEBUG

/********************************************************/
/* KERNEL0 - NO INSTRUCTION LEVEL PARALLELISM (ILP = 0) */
/********************************************************/
__global__ void kernel0(float * __restrict__ d_a, const float * __restrict__ d_b, const float * __restrict__ d_c, const int N) {

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

    if (tid < N) {

        float a = d_a[tid];
        float b = d_b[tid];
        float c = d_c[tid];

        for (unsigned int i = 0; i < N_ITERATIONS; i++) {
            a = a * b + c;
        }

        d_a[tid] = a;
    }

}

/*****************************************************/
/* KERNEL1 - INSTRUCTION LEVEL PARALLELISM (ILP = 2) */
/*****************************************************/
__global__ void kernel1(float * __restrict__ d_a, const float * __restrict__ d_b, const float * __restrict__ d_c, const int N) {

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

    if (tid < N / 2) {

        float a1 = d_a[tid];
        float b1 = d_b[tid];
        float c1 = d_c[tid];

        float a2 = d_a[tid + N / 2];
        float b2 = d_b[tid + N / 2];
        float c2 = d_c[tid + N / 2];

        for (unsigned int i = 0; i < N_ITERATIONS; i++) {
            a1 = a1 * b1 + c1;
            a2 = a2 * b2 + c2;
        }

        d_a[tid] = a1;
        d_a[tid + N / 2] = a2;
    }

}

/*****************************************************/
/* KERNEL2 - INSTRUCTION LEVEL PARALLELISM (ILP = 4) */
/*****************************************************/
__global__ void kernel2(float * __restrict__ d_a, const float * __restrict__ d_b, const float * __restrict__ d_c, const int N) {

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

    if (tid < N / 4) {

        float a1 = d_a[tid];
        float b1 = d_b[tid];
        float c1 = d_c[tid];

        float a2 = d_a[tid + N / 4];
        float b2 = d_b[tid + N / 4];
        float c2 = d_c[tid + N / 4];

        float a3 = d_a[tid + N / 2];
        float b3 = d_b[tid + N / 2];
        float c3 = d_c[tid + N / 2];

        float a4 = d_a[tid + 3 * N / 4];
        float b4 = d_b[tid + 3 * N / 4];
        float c4 = d_c[tid + 3 * N / 4];

        for (unsigned int i = 0; i < N_ITERATIONS; i++) {
            a1 = a1 * b1 + c1;
            a2 = a2 * b2 + c2;
            a3 = a3 * b3 + c3;
            a4 = a4 * b4 + c4;
        }

        d_a[tid] = a1;
        d_a[tid + N / 4] = a2;
        d_a[tid + N / 2] = a3;
        d_a[tid + 3 * N / 4] = a4;
    }

}

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

    //const int N = 8192 * 64;
    const int N = 8192;
    //const int N = 1024;

    TimingGPU timerGPU;

    float *h_a = (float*)malloc(N*sizeof(float));
    float *h_a_result_host = (float*)malloc(N*sizeof(float));
    float *h_a_result_device = (float*)malloc(N*sizeof(float));
    float *h_b = (float*)malloc(N*sizeof(float));
    float *h_c = (float*)malloc(N*sizeof(float));

    for (int i = 0; i<N; i++) {
        h_a[i] = 2.;
        h_b[i] = 1.;
        h_c[i] = 2.;
        h_a_result_host[i] = h_a[i];
        for (unsigned int k = 0; k < N_ITERATIONS; k++) {
            h_a_result_host[i] = h_a_result_host[i] * h_b[i] + h_c[i];
        }
    }

    float *d_a; gpuErrchk(cudaMalloc((void**)&d_a, N*sizeof(float)));
    float *d_b; gpuErrchk(cudaMalloc((void**)&d_b, N*sizeof(float)));
    float *d_c; gpuErrchk(cudaMalloc((void**)&d_c, N*sizeof(float)));

    gpuErrchk(cudaMemcpy(d_a, h_a, N*sizeof(float), cudaMemcpyHostToDevice));
    gpuErrchk(cudaMemcpy(d_b, h_b, N*sizeof(float), cudaMemcpyHostToDevice));
    gpuErrchk(cudaMemcpy(d_c, h_c, N*sizeof(float), cudaMemcpyHostToDevice));

    /***********/
    /* KERNEL0 */
    /***********/
    timerGPU.StartCounter();
    kernel0 << <iDivUp(N, BLOCKSIZE), BLOCKSIZE >> >(d_a, d_b, d_c, N);
#ifdef DEBUG
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());
#endif
    // --- Remember: timing is in ms
    printf("Number of operations = %f; GFlops = %f\n", (float)N*(float)N_ITERATIONS, (1.e-6)*((float)N*(float)N_ITERATIONS) / timerGPU.GetCounter());
    gpuErrchk(cudaMemcpy(h_a_result_device, d_a, N*sizeof(float), cudaMemcpyDeviceToHost));
    for (int i = 0; i<N; i++) if (h_a_result_device[i] != h_a_result_host[i]) { printf("Error at i=%i! Host = %f; Device = %f\n", i, h_a_result_host[i], h_a_result_device[i]); return 1; }

    /***********/
    /* KERNEL1 */
    /***********/
    gpuErrchk(cudaMemcpy(d_a, h_a, N*sizeof(float), cudaMemcpyHostToDevice));
    timerGPU.StartCounter();
    kernel1 << <iDivUp(N / 2, BLOCKSIZE), BLOCKSIZE >> >(d_a, d_b, d_c, N);
#ifdef DEBUG
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());
#endif
    // --- Remember: timing is in ms
    printf("Number of operations = %f; GFlops = %f\n", (float)N*(float)N_ITERATIONS, (1.e-6)*((float)N*(float)N_ITERATIONS) / timerGPU.GetCounter());
    gpuErrchk(cudaMemcpy(h_a_result_device, d_a, N*sizeof(float), cudaMemcpyDeviceToHost));
    for (int i = 0; i<N; i++) if (h_a_result_device[i] != h_a_result_host[i]) { printf("Error at i=%i! Host = %f; Device = %f\n", i, h_a_result_host[i], h_a_result_device[i]); return 1; }

    /***********/
    /* KERNEL2 */
    /***********/
    gpuErrchk(cudaMemcpy(d_a, h_a, N*sizeof(float), cudaMemcpyHostToDevice));
    timerGPU.StartCounter();
    kernel2 << <iDivUp(N / 4, BLOCKSIZE), BLOCKSIZE >> >(d_a, d_b, d_c, N);
#ifdef DEBUG
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());
#endif
    // --- Remember: timing is in ms
    printf("Number of operations = %f; GFlops = %f\n", (float)N*(float)N_ITERATIONS, (1.e-6)*((float)N*(float)N_ITERATIONS) / timerGPU.GetCounter());
    gpuErrchk(cudaMemcpy(h_a_result_device, d_a, N*sizeof(float), cudaMemcpyDeviceToHost));
    for (int i = 0; i<N; i++) if (h_a_result_device[i] != h_a_result_host[i]) { printf("Error at i=%i! Host = %f; Device = %f\n", i, h_a_result_host[i], h_a_result_device[i]); return 1; }

    cudaDeviceReset();

    return 0;

}