CUDA 中的编译时信息

Posted 2023-02-16

【中文标题】CUDA 中的编译时信息

【英文标题】：Compile-time information in CUDA

【发布时间】：2014-03-10 09:51:41

【问题描述】：

我正在优化一个时间紧迫的 CUDA 内核。我的应用程序接受各种影响行为的开关（例如，是否使用 3 阶或 5 阶导数）。考虑一组 50 个开关的近似值，其中每个开关都是一个整数变量（有时是布尔值，或浮点数，但这种情况与这个问题不太相关）。

所有这些开关在应用程序执行期间都是不变的。大多数这些开关都是运行时的，我将它们存储在常量内存中，以便利用缓存机制。其他一些开关可以是编译时的，如果客户想要更改开关中的值，他可以重新编译应用程序。一个非常简单的例子可能是：

__global__ void mykernel(const float* in, float *out)

    for ( /* many many times */ )
        if (compile_time_switch)
            do_this(in, out);
        else
            do_that(in, out);

假设do_this 和do_that 是计算密集型的并且非常便宜，我优化for 循环使其开销可以忽略不计，我必须将if 放在迭代中。如果编译器识别出compile_time_switch 是静态信息，它可以优化对“错误”函数的调用，并创建与 if 不存在一样优化的代码。现在真正的问题是：

我可以通过哪些方式向编译器提供此开关的静态值？我看到了下面列出的两种这样的方法，但它们都不适合我。还有哪些其他可能性？

---

模板参数

提供模板参数可以实现这种静态优化。

template<int compile_time_switch>
__global__ void mykernel(const float* in, float *out)

    for ( /* many many times */ )
        if (compile_time_switch)
            do_this(in, out);
        else
            do_that(in, out);

这个简单的解决方案对我不起作用，因为我无法直接访问调用内核的代码。

静态成员

考虑以下结构：

struct GlobalParameters

    static const bool compile_time_switch = true;
;

现在GlobalParameters::compile_time_switch 包含我想要的静态信息，并且该编译器将能够优化内核。不幸的是，CUDA 不支持这样的静态成员。

编辑：最后一条语句显然是错误的。结构的定义当然是合法的，您可以在设备代码中使用静态成员GlobalParameters::compile_time_switch。编译器内联变量，因此最终代码将直接包含该值，而不是运行时变量访问，这是您期望从优化器编译器获得的行为。所以，第二个选项实际上是合适的。

由于这个事实和 kronos 的回答，我认为我的问题已解决。但是，我仍在寻找其他替代方法来向编译器提供编译时信息。

【参考方案1】：

你的第三个选项是预处理器定义：

#define compile_time_switch 1

__global__ void mykernel(const float* in, float *out)

    for ( /* many many times */ )
        if (compile_time_switch)
            do_this(in, out);
        else
            do_that(in, out);

预处理器将完全丢弃 else 情况，编译器在他的死代码消除过程中没有什么可优化的，因为没有死代码。

此外，您可以使用 -D 命令行开关指定定义，并且（我认为）nvidia 支持的任何编译器都将接受 -D（msvc 可能使用不同的开关）。

【讨论】：

好的，好的，谢谢分享。实际上，这是我们目前的解决方案。我们试图避免过多地利用预处理器，因为它会降低代码的可读性，此外，我们希望对编译时和运行时切换都有类似的机制，即切换是编译时还是运行时-从用户的角度来看，时间应该尽可能透明（用户=读取开关的程序员）。

@Spiros Yves Daoust 对我的帖子进行了一些编辑。这些更改增加了代码的可读性。

哦，是的，当然你可以这样使用宏。因此，总而言之，使用预处理器是一种有效的选择。我仍在寻找其他可能的方法来指定编译时信息，以便编译器可以优化死代码。