解释GPU NVIDIA执行时间

Question

我编写了简单的代码来使用 GPU 对数组求和。我使用一维网格和块。每个块计算数组的“2 倍块大小”元素并将这些元素减少为 1 个值。然后主机代码将对这些块的减少值求和。这是代码：

__global__ void reduceBlksKernel1(int * in, int n, int * out) {
    int i = blockIdx.x * blockDim.x * 2 + threadIdx.x * 2;

    for (int stride = 1; stride < 2 * blockDim.x; stride *= 2) {
        if (threadIdx.x % stride == 0) {
            if (i + stride < n) {
                in[i] += in[i + stride];
            }
        }
        __syncthreads();
    }

    // use 1 thread to write the reduced value to out array
    if (threadIdx.x == 0) {
        out[blockIdx.x] = in[blockIdx.x * blockDim.x * 2];
    }
}

我使用 Google colab，这里是 GPU 的信息：

我已经运行一些不同块大小的测试：

有人可以帮我解释为什么尽管所有情况下的占用率都是 100%，但执行时间会随着每个 SM 块的增加而减少吗？谢谢！

Answer 1

让我们考虑一下内核的这一部分，其中“大部分”工作正在完成：

for (int stride = 1; stride < 2 * blockDim.x; stride *= 2) {
    if (threadIdx.x % stride == 0) {
        if (i + stride < n) {
            in[i] += in[i + stride];
        }
    }
    __syncthreads();
}

每次通过for-loop，让我们先问一个问题，有多少个“活动”线程？这是第一个if-test给出的，也就是threadIdx.x%stride==0。 stride 一开始是 1，然后是 2，然后是 4，依此类推。所以第一遍有完整的线程，第二遍选择每第二个线程，第三遍选择每第四个线程，依此类推。

现在让我们问一个问题，对于给定的线程块大小，在所有迭代中我们有多少平均个活动线程for-loop。

最后，让我们针对上述问题构建一个 table，显示每个线程块配置的平均活动线程数，同时考虑驻留多个线程块的乘法因子。 table 中的第一个平均值将像这样计算：

 1024+512+256+128+64+32+16+8+4+2+1 = 2047/11 = ~186 active threads (average)

对于这种情况，只能驻留 1 个线程块，因为 Turing 每个 SM 最多可补充 1024 个线程。 table 在其他情况下会是什么样子？

threads/block    average per block   number of blocks resident  average per SM
1024             186                 1                          186
512              102                 2                          204
256              56.8                4                          227
128              31.9                8                          255

因此您可以看到，使用较小的线程块，但使用更多的线程块（至少对于这个范围和这部分分析）会导致每个 SM 的平均活动线程数更多。实际上，SM 平均消耗更高的内存流量，因此可以更快地完成此 memory-bound 工作负载。

那我们该怎么办呢？

我的建议是在缩减操作之前使用 grid-stride 循环。您可以通过查看缩减 #7 中的 canonical parallel reduction material. In that material, note the introduction of the grid-stride loop 了解这一点。这个想法是选择线程块的数量来填充您的 SM，并使用网格步幅循环来有效地加载“所有”数据。这将数据加载和部分缩减移动到非常有效的代码配置，同时将 less-efficient 扫描样式缩减保留到整个工作的希望较小的部分。这将最大限度地减少在扫描阶段花费的内核时间，从而最大限度地减少线程块大小对性能可变性的影响。

当然，即使有上述建议，如果您想要最后一点性能，您可能希望在各种线程块配置下评估您的代码以获得最佳性能。这种“ninja-tuning”的建议几乎适用于所有 GPU 代码。

顺便说一句，这里只使用全局内存的缩减方法不是我推荐的方法（请参阅链接的规范 material 以获得更好的方法），但这似乎不是你的问题的主旨。

解释GPU NVIDIA执行时间

Explain GPU NVIDIA execution time

c++

architecture

cuda

gpu

nvidia