解释 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 的信息：

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

有人可以帮我解释为什么执行时间会随着每个 SM 的更多块而减少，尽管事实上所有情况下的占用数都是 100%？谢谢！

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循环时，我们首先要问一个问题，有多少个“活动”线程？这是由第一个 if 测试给出的，即threadIdx.x%stride==0。 stride 从 1 开始，然后是 2，然后是 4，以此类推。所以第一遍有完整的线程，第二遍选择每个第二个线程，第三遍选择每个第四个线程，依此类推。

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

最后，让我们为上述问题建立一个表格，显示每个线程块配置的平均活动线程，并考虑多个线程块驻留的乘法因素。该表中的第一个平均值将按如下方式计算：

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

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

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 平均消耗更高的内存流量，因此可以更快地完成这种受内存限制的工作负载。

那么我们应该怎么做呢？

我的建议是在您的归约操作之前使用网格步长循环。您可以通过查看canonical parallel reduction material 来了解这一点。在该材料中，请注意在减少 #7 处引入了 grid-stride loop。这个想法是选择线程块的数量来填充你的 SM，并使用网格步幅循环来有效地加载“所有”数据。这会将数据加载和部分缩减转移到非常有效的代码配置中，同时将效率较低的扫描样式缩减保留到整个工作的较小部分。这将最大限度地减少在扫描阶段花费的内核时间，最大限度地减少线程块大小对性能可变性的影响。

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

顺便说一句，我不推荐您在此处仅使用全局内存的减少方法（请参阅链接的规范材料以获得更好的方法），但这似乎不是您问题的重点。