为什么这个简单的 openCl 内核的矢量化会使其变慢？

Question

我目前正在使用 OpenCl 在 GPU 上实现一个简单的并行索引总和。这是需要对大型数组进行连续流压缩的项目的一部分，因此我认为首先熟悉该算法，然后尝试构建一个更高级的版本（根据我有限的理解，所有这些都派生出来）是一个好主意来自 Blelloch 算法）。

目前，我有一个名为

step_naive_prefix_sum

的内核，它只是将

a

添加到自身上，并进行一定的移位，并将结果保存到

b

中。

kernel void step_naive_prefix_sum(global int* a, global int* b, int offset, int nels)
{
    int i = get_global_id(0);
    if (i >= nels) return;
    b[i] = a[i];
    if (i >= offset) {
        b[i] += a[i - offset];
    }
}

内核工作，算法的其余部分也依次工作。

在非常小的测试负载上，我得到了这些结果（作为每次完成时通过

clGetEventProfilingInfo

报告的内核时间的总和）：

prefix_sum: 68.112549 ms

我认为这种任务非常适合使用 OpenCl 向量类型进行向量化。所以我迅速编码了变体：

kernel void step_naive_prefix_sum_vectorized2(global int2* a, global int2* b, int offset, int ncouples)
{
    int i = get_global_id(0);
    if (i >= ncouples) return;
    b[i] = a[i];
    if (i >= offset) {
        b[i] += a[i - offset];
    }
}

kernel void step_naive_prefix_sum_vectorized4(global int4* a, global int4* b, int offset, int nquarts)
{
    int i = get_global_id(0);
    if (i >= nquarts) return;
    b[i] = a[i];
    if (i >= offset) {
        b[i] += a[i - offset];
    }
}

并修改了调用内核的循环，使其处理 offset = 1、offset = 2、offset >= 4 的情况：

• offset = 1：调用普通内核
• offset = 2：调用 int2 内核，全局大小减半（仍四舍五入为本地组大小的倍数）并传递 n Couples = nels/2
• offset >=4：调用 int4 内核，全局大小除以 4（仍四舍五入为本地组大小的倍数）并传递 n Couples = nels/4

请注意，根据设计，将数组大小填充到（远高于 4）2 的幂，因此无需检查数组末尾的边缘情况。

然而，时间变成了：

prefix_sum: 79.486565 ms

这让我很惊讶。在一些额外的测试负载中，结果保持一致。

有人可以帮助我了解出了什么问题，并可能为我指出正确的方向吗？ 当然，任何其他关于如何使这个（或类似的）运行得更快的提示都是可以接受的！

Answer 1

向量类型破坏了合并的内存访问。

您的内核受带宽限制。

int2

和

int4

是保存 2 或 4 个整数的结构，因此连续线程不再访问连续的内存位置，而是访问 2 或 4 步长的位置。

在旧的 GPU 架构上，跨步内存访问会影响性能。现代架构具有内置补偿机制，但它仍然会对内存带宽产生负面影响。