尝试了解 CUDA 中的内核优化。这些变化的影响可以忽略不计

Question

我有一个 CUDA mandelbrot 内核，它以标准方式通过：

__global__ void mandelbrot_worker(uint32_t* count_arr, uint32_t w, uint32_t h,
                                  const uint32_t max_count, const float xmin,
                                  const float xmax, const float ymin, const float ymax) {
    // Calculate global thread coordinates in the 2D Mandelbrot set grid
    int j = blockIdx.x * blockDim.x + threadIdx.x;  // x-coordinate
    int i = blockIdx.y * blockDim.y + threadIdx.y;  // y-coordinate

    // if (i >= h || j >= w) {
    //     printf("ERROR: OUT OF BOUNDS");
    //     return;
    // } // Bounds check

    // Map pixel coordinate (i, j) to complex plane coordinates (x, y)
    float x0 = xmin + (xmax - xmin) * j / w;
    float y0 = ymin + (ymax - ymin) * i / h;

注释掉 if 语句来检查是否有越界线程没有任何改变。在这种情况下，传入

w = 3840*8
h = 2160*8

条件永远不会被执行。

但我认为删除 if 语句会稍微改变时间。可能有，但显然 Mandelbrot 中的循环是最大的部分。

我也认为计算x0和y0效率低下，传入：

__global__ void mandelbrot_worker2(uint32_t* count_arr, uint32_t w, uint32_t h,
                                  const uint32_t max_count, const float xmin,
                                  const float dx_per_pixel, const float ymin, const float dy_per_pixel) {
    // Calculate global thread coordinates in the 2D Mandelbrot set grid
    int j = blockIdx.x * blockDim.x + threadIdx.x;  // x-coordinate
    int i = blockIdx.y * blockDim.y + threadIdx.y;  // y-coordinate

    // Map pixel coordinate (i, j) to complex plane coordinates (x, y)
    float x0 = xmin + dx_per_pixel * j ;
    float y0 = ymin + dy_per_pixel * i ;

最大迭代次数是 64，所以也许这只是一个很小的百分比。 最后，CUDA 的优化好像一直在进行？ -O3 标志对主机端影响很大，但似乎根本不影响 CUDA 上的代码速度。

所以问题是：

1. 与 C 不同，浮点项在 CUDA 中会重新排列吗？
2. 将常量从内核中取出有什么作用吗？忘记这个例子，循环花费了大部分时间。
3. 是否有其他编译器选项可以生成更好的代码？ 我们刚刚使用 nvcc -O3 mandel.cu

最后，这段代码是否按顺序写入内存，他们是否尝试以这种方式安排它，或者是随机顺序，因为线程执行他们想要的操作？我想到的一个想法是让一个线程将值写入 8 个寄存器，然后一次将它们全部写入。

Answer 1

条件永远不会被执行。

这取决于内核配置（即每个维度的线程和块数）。一般来说，人们使用的线程数可以被扭曲的大小整除，即 32（到目前为止在所有设备上）。由于

w

和

h

可以被32整除，所以没有问题。另请注意，除非某些线程由于预测而处于非活动状态，否则条件通常很便宜。最后但并非最不重要的一点是，无论如何，主计算循环应该占据大部分时间，所以你不应该关心这部分。

注意迭代次数取决于 Mandelbrot 中的位置。这具有巨大的影响。对于集合的很大一部分，

z

永远不会在合理的时间内收敛（也不会发散）。在某些区域，这会产生噪音并且扭曲发散是一个问题。重新映射是增加活动线程数量的关键。您需要 profile 您的代码来查看代码的哪一部分有效或无效（Nshight 对此非常有用）。

与 C 不同，浮点项在 CUDA 中会重新排列吗？

不可以，除非您手动请求。在 C 中也是如此（至少在主流编译器上）。 Clang 和 GCC 上的

-ffast-math

可以在没有 NaN 或 Inf 的情况下启用此功能以及许多其他优化（这对于您的目标应用程序可能存在潜在危险）。 AFAIK 默认情况下保证 IEEE-754 合规性。

将常量从内核中取出有什么作用吗？忘记这个例子，循环花费了大部分时间。

我不确定是否理解，但使用

w

和

h

常量可以显着提高除法的性能。但使用常量并不总是更快：它取决于实际完成的计算。乘法和加法的好处即使有的话也是很小的（除了一些非常特殊的情况）。对于迭代次数，它可以帮助编译器完全展开（当次数相当小时）。

是否有其他编译器选项可以生成更好的代码？我们只是使用 nvcc -O3 mandel.cu

快速数学可以提供帮助。话虽这么说，您可以自己优化代码，因此甚至不需要它。这个想法是尽可能多地使用 FMA 运算，而不仅仅是单独的乘法或加法。

这段代码是否按顺序写入内存

线程由 warp（32 个线程）执行，并且在合并内存访问时进行打包加载。这对于性能至关重要。 Warp 调度取决于硬件。块的大小有时会影响性能。

对于性能来说，最重要的是进行合并加载/存储，如果可能的话，所有线程都处于活动状态（否则带宽就会被浪费）。缓存也可以提供帮助，但是在这里，输出大小太大，您无法从中受益。

---

Mandelbrot 通常是计算密集型，因此您应该关心最小化主循环中的操作数量。 循环展开可以帮助实现这一目标。也将 MUL+ADD 合并到 FMA 操作。此外，warp divergence 是 Mandelbrot 上的一个关键问题，您应该尽可能减少以确保有效使用所有 SIMD 单元。如前所述，remapping 是一个有效的解决方案。您也可以检测一些模式来避免大量迭代（实际上相当复杂）。最重要的是，与所有 CUDA 内核一样，您应该检查 占用率 以最大限度地提高内核的性能。