CUDA 5.0 重播开销答案

【问题标题】：CUDA 5.0 Replay OverheadCUDA 5.0 重播开销
【发布时间】：2013-07-11 17:50:33
【问题描述】：

我是一个新手 CUDA 程序员。我最近了解了更多关于在较低占用率下实现更好性能的信息。这是一个代码 sn-p，我需要帮助来理解一些关于重放开销和指令级并行性的事情

__global__ void myKernel(double *d_dst, double *d_a1, double *d_a2, size_t SIZE) 
{

    int tId = threadIdx.x + blockDim.x * blockIdx.x;

    d_dst[tId]            = d_a1[tId] * d_a2[tId];
    d_dst[tId + SIZE]     = d_a1[tId + SIZE] * d_a2[tId + SIZE];
    d_dst[tId + SIZE * 2] = d_a1[tId + SIZE * 2] * d_a2[tId + SIZE * 2];
    d_dst[tId + SIZE * 3] = d_a1[tId + SIZE * 3] * d_a2[tId + SIZE * 3];
}

这是我的简单内核，它简单地将两个二维数组相乘以形成第三个二维数组（从逻辑角度来看），其中这些数组都作为平面一维数组放置在设备内存中。

下面我再展示一段代码sn-p：

void doCompute() {

    double *h_a1;
    double *h_a2;

    size_t SIZE = pow(31, 3) + 1;

   // Imagine h_a1, h_a2 as 2D arrays
   // with 4 rows and SIZE Columns
   // For convenience created as 1D arrays 

    h_a1 = (double *) malloc(SIZE * 4 * sizeof(double));
    h_a2 = (double *) malloc(SIZE * 4 * sizeof(double));

    memset(h_a1, 5.0, SIZE * 4 * sizeof(double));
    memset(h_a2, 5.0, SIZE * 4 * sizeof(double));

    double *d_dst;
    double *d_a1;
    double *d_a2;

    cudaMalloc(&d_dst, SIZE * 4 * sizeof(double));
    cudaMalloc(&d_a1,  SIZE * 4 * sizeof(double));
    cudaMalloc(&d_a2,  SIZE * 4 * sizeof(double));

    cudaMemcpy(d_a1, h_a1, SIZE * 4 * sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(d_a2, h_a2, SIZE * 4 * sizeof(double), cudaMemcpyHostToDevice);

    int BLOC_SIZE = 32;
    int GRID_SIZE = (SIZE + BLOC_SIZE - 1) / BLOC_SIZE;

    myKernel <<< GRID_SIZE, BLOC_SIZE >>> (d_dst, d_a1, d_a2, SIZE);
}

Q1) 我在这里打破了任何合并的内存访问模式吗？

Q2) 我可以说对内存的访问，它们在内核中的编码方式也是指令级并行的例子吗？如果是，我使用的是 ILP2 还是 ILP4？和为什么？

Q3) 如果我所做的一切都是正确的，那么为什么 nvvp 分析器会给我以下消息？

Total Replay Overhead: 4.6%
Global Cache Replay Overhead: 30.3%

如何减少或修复它们？

干杯，

【问题讨论】：

您选择 SIZE 作为非 2 的幂值不会提供最佳的合并可能性。但是缓存应该对此有所帮助。我认为，视觉分析器应该报告大约 50% - 60% 的全局内存使用效率数字。您的内核代码应该为 ILP 提供一些机会，因为每一行代码都不依赖于前一行。但 ILP 也将取决于您使用的特定 GPU。如果您将SIZE 设为 32 或 16 的倍数，您可能会得到不同/更好的结果。
@Robert 嗨，我要解决的问题的性质不允许我的 SIZE 正好是 2 的幂。所以我即兴发挥，我用额外的内存位置填充它们以使 SIZE 32 的倍数。就像在这种情况下一样，我添加了 1 分析器显示加载和存储效率为 100%。但这是否也意味着，它是一个合并的内存访问？ ILP 是否也适用于内存访问？ warp 中的线程正在跳跃，但它们没有破坏 128 字节缓存线？还是他们？我使用双打而不是花车？最后，如果一切都在这里，为什么还要开销？
是的，我的错误，SIZE 值可以被 32 整除，这应该足以提供良好的合并，因此 VP 报告了 100% 的全局内存负载效率。 ILP 和合并实际上是两个不同的概念。没有 ILP 机会的代码仍然可以有 100% 的带宽利用率、100% 的合并，并利用 100% 的可用内存带宽。您的代码应该 100% 合并。
顺便说一句，我不认为那些memset 电话正在按照您的想法进行。 memset 接受 byte 数量并设置 byte 位置。您正在获取 double 数量并将每个字节设置为 5。
感谢您的澄清。对于 ILP，只要在代码中找到可以独立执行的指令，ILP 就会发挥作用。任何依赖资源都会让warp进入睡眠状态，直到满足依赖关系并且warp被安排到下一个执行周期，这取决于负载或调度程序的工作方式。那是对的吗？最后一个问题仍然是什么是开销重放？为什么我会在这个例子中看到它？

标签： optimization cuda

【解决方案1】：

编译器对可能的 ILP 漏洞利用调度指令的能力有限。 GPU 本身也必须具有 ILP 能力，其程度因 GPU 代而异。是的，任何不可用的资源都可能导致扭曲停止，典型的资源是内存中需要的数据。您所询问的重播数量的定义在here 中给出。

因此，例如，全局缓存重放开销将由缓存未命中触发，并且您的代码将出现一些缓存未命中。即使您拥有 100% 的合并访问和（几乎）100% 的带宽利用效率，缓存未命中也是可能的。

【讨论】：