10

我在 OpenCL 上运行 mandelbrot 生成器（来自静态参数的 2D 图像）。该程序很简单：

__kernel
void mandelbrot(__global uchar * output, 
                const float xstep,
                const float xoffset,
                const float ystep,
                const float yoffset,
                const int maxiter)
{
    int gid_y = get_global_id(1);
    int gid_x = get_global_id(0);

    //calculate x and y on the fly for every pixel. 
    //This is just as fast as reading precalculated rulers from global memory.
    float x = gid_x * xstep + xoffset;
    float y = gid_y * ystep + yoffset;

    float real = 0;
    float imag = 0;

    int out = 0;

    for(int curiter = 0; curiter < maxiter; curiter++) {
        float nreal = real*real - imag*imag + x;
        imag = 2* real*imag + y;
        real = nreal;

        if (real*real + imag*imag > 4.0f) {
            out = curiter;
            break;
        }
    }

    //normalize output
    out *= 256.0 / (float)maxiter;
    output[gid_y * get_global_size(0) + gid_x] = out;

}

[编辑] [发布完整内核，并按照建议交换行和列。通过这种方式，我在 AMD 上获得了 18% 的性能，但在 NVidia 上获得了 0% 的性能。原始代码是

output[get_global_id(0) * get_global_size(1) + get_global_id(1)] = out;

[/编辑]

我在我的 Nvidia Quadro 1000M 上运行它，它有 2 个计算单元和 96 个 CUDA 内核（每个计算单元 48 个内核）。

在将内核排队时，我正在通过更改本地组大小来玩弄。这些是我在生成 400Mpixel 图像时得到的不同尺寸的性能结果。所有数字都来自 OpenCL 分析器，不包括返回到操作系统的最终内存副本。图像为 40992x10272 - 高度和宽度都可以被 48 整除。

rows x columns
8x8: 397 MPixel/s
8x12: 505 MPixel/s
8x16: 523 MPixel/s
8x24: 521 MPixel/s
8x32: 520 MPixel/s
8x48: 520 MPixel/s

1x48: 321 MPixel/s
2x32: 424 MPixel/s
2x48: 523 MPixel/s
4x24: 519 MPixel/s
3x32: 525 MPixel/s
4x32: 525 MPixel/s
4x48: 525 MPixel/s

12x8: 490 MPixel/s
12x12:464 MPixel/s
12x24:505 MPixel/s
12x32:508 MPixel/s
12x48:433 MPixel/s

16x8: 499 MPixel/s
16x12:499 MPixel/s
16x16:472 MPixel/s
16x24:450 MPixel/s
16x32:440 MPixel/s
16x48:418 MPixel/s

其中一些数字让我感到困惑。虽然很清楚为什么我会在 48 列中获得最佳结果（感谢 SIMD 操作的工作方式），但我不明白：

  1. 为什么当我每组使用 16 行时性能会急剧下降？
  2. 为什么我的 1x48 性能很差？
  3. 为什么在天堂我能用 3x32、4x32 和 8x32 获得最佳性能？！？我原本预计 33% 的 SIMD 处理器处于空闲状态，但看起来工作组位于两个计算单元之间？！？
  4. 为什么 PREFERRED_WORK_GROUP_SIZE_MULTIPLE 返回 32 而不是 48？
  5. 仅考虑我从 OpenCL 信息结构中获得的信息，是否有一种非经验方法可以计算出任何 GPU（ATI/Nvidia/Intel HD）上最高性能的几何形状？

提前致谢

4

3 回答 3

22

我在这里回答了一个类似的问题，在阅读以下内容之前您可能会觉得有趣。

为什么当我每组使用 16 行时性能会急剧下降？

实际上，当您使用 12 行时，它已经降级了。内存访问通过事务进行。事务将一次获取一定数量的字节。现在，如果多个工作项尝试访问数组中的多个连续元素，则意味着一个事务可能足以为它们提供服务。

因为您以这种方式访问​​内存：

output[get_global_id(0) * get_global_size(1) + get_global_id(1)] = out;

这意味着本地大小在维度 0 中越大，事务的数量就越大，因为您必须访问非连续元素（由 get_global_size(1) 元素分隔）。并且全局内存访问是昂贵的。

因此，对于 12/16 行，您至少需要 12/16 个事务。这导致你的第二个问题：

为什么我的 1x48 性能很差？

根据我之前所说的，似乎性能应该很好，因为事务的数量很少。

但是这里出现了空闲线程的问题。正如其他人已经指出的那样，您获得的有关每个 SM 48 个内核的信息是错误的。线程在 NVIDIA 硬件上以 32 个组（对于 NVIDIA 称为 warp）执行。请注意，这些组称为波前，对于 AMD，最多可以有 64 个线程。由于在这种情况下您有一个由 48 个线程（1 x 48）组成的工作组，这意味着调度了 64 个线程。调度的线程数始终是 32 的倍数，因为您无法执行扭曲的一小部分。

因此，在这种情况下，您有四分之一的线程什么都不做。实际上，当您与 2x32（仍然是 64 个线程 - 2 个扭曲，但已充分利用）获得的结果进行比较时，321 MPixel/s 几乎是 424 MPixel/s 的 3/4。

值得注意的是这个结果：2x48: 523 MPixel/s。在这种情况下，您的工作组大小是 96 是 32 的倍数。所以没有空闲线程。

为什么在天堂我能用 3x32、4x32 和 8x32 获得最佳性能？！？

嗯，答案来自前两个：你使用了 32 的倍数，并且你保持维度 0 中的线程数相对较少。但让我们仔细看看你的结果：

2x32:  424 MPixel/s
3x32:  525 MPixel/s
4x32:  525 MPixel/s
8x32:  520 MPixel/s
16x32: 440 MPixel/s

最后两行的性能下降很容易用所说的来解释。但是，第一行和第二行之间的性能提升却没有。

在这种情况下，性能的提高来自其他地方。在第二种情况下，在同一个 SM上运行了足够多的 warp以隐藏访问内存延迟。您看到 REFERRED_WORK_GROUP_SIZE_MULTIPLE 值仅表明您应该尝试使用该值的 MULTIPLE 以获得最佳性能。可以同时在同一个 SM 上安排多个 warp。

那么它是怎样工作的？让我们以 3x32 为例。您有一个由 3 个 warp 组成的工作组。因为它们属于同一个工作组，所以按照 OpenCL 标准的要求，它们被安排在同一个 SM 上（如果不是这样，工作组内的线程之间就不可能同步）。

第一个 warp 开始运行，直到因为需要访问内存而停止。同时，warp 1 等待内存事务完成，warp 2 可以开始运行。由于 SM 上有很多寄存器，因此 SM 可以轻松快速地切换上下文以运行其他 warp。经线 1 的所有变量都保留在分配给经线 1 的寄存器上。然后经线 2 到达需要访问内存的行并停止。这时，下一个准备运行的warp就可以开始运行了。如果它的内存访问完成，它可能是warp 3，也可能是warp 1。在您的情况下，似乎运行的是经线 3，因为您在 2x32 和 3x32 之间存在差异。在第一种情况下，没有足够的扭曲计划来隐藏内存访问，尽管在第二种情况下有。

事实上，这也影响了问题 2 中 1x48 尺寸的不良性能。

为什么 PREFERRED_WORK_GROUP_SIZE_MULTIPLE 返回 32 而不是 48？

已经回答了。

仅考虑我从 OpenCL 信息结构中获得的信息，是否有一种非经验方法可以计算出任何 GPU（ATI/Nvidia/Intel HD）上最高性能的几何形状？

就像任何其他语言一样。当您了解它的底层工作原理时，它可以帮助您生成良好的第一个代码。但是您仍然必须对其进行基准测试，并通过反复试验来调整它。记住我刚刚写的只是影响性能的一小部分。从 OpenCL 查询一些信息并结合对 CPU/GPU 的良好理解显然会有所帮助……但仅此而已。

因为很多影响性能的参数都是对立的，你在一方得到的，在另一方就会失去。

因此，请继续进行基准测试；）。

于 2013-08-08T12:36:55.023 回答
1

这完全取决于您未显示的代码。这就是关键。如果您的代码非常简单，即：out = 8;那么您的假设可能是正确的。

但是，正如您所说，值 REFERRED_WORK_GROUP_SIZE_MULTIPLE 返回 32。这意味着，32 是计算单元可以在不影响性能的情况下并行启动的最大并发线程数。例如，启动超过 32 个是没有意义的。如果使用 32 个，您已经耗尽了本地内存存储，您需要递归到全局内存（这非常慢）。

如果您尝试超过建议的限制，您将获得 -> 性能下降。不是32更好，是相反。48不好。

我向你推荐：

  1. 如果可能，使用自动大小（将 null 作为本地大小传递给内核）。如果您不担心本地 worksize 形状，这将导致最大性能。
  2. 如果您需要手动设置本地大小，请使用 REFERRED_WORK_GROUP_SIZE_MULTIPLE 作为参考。
于 2013-08-07T16:58:24.353 回答
0

内核访问全局内存的方式至关重要，由工作组和全局维度决定：

  • 同一工作组中的连续工作项将写入哪些地址？这里的步幅是 get_global_size(1)，您可能想要交换 X 和 Y。在连续工作项中处理连续元素通常更快。这是最重要的因素。

  • 连续的工作组会写什么地址？连续的工作组将经常同时安排在不同的计算单元上。他们最终可能会竞争同一个频道/银行，从而导致性能损失。

  • 通常最好写 32 位整数而不是字节。

为了最大限度地提高性能，我建议您引入更多按钮来转动：编写内核计算单个工作项内的几个像素块（例如 4x2），然后对（块大小）x（工作组大小）的所有组合进行基准测试x（XY 交换）x（图像大小）。然后选择最适合您的 GPU。

于 2013-08-07T17:50:51.087 回答