CUDA 中的编译时信息

Question

我正在优化一个时间紧迫的 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 的回答，我认为我的问题已解决。但是，我仍在寻找其他替代方法来向编译器提供编译时信息。

Answer 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 可能使用不同的开关）。