相同类型的多个模板声明的目的

Question

我遇到过一个具有多个声明的代码库，并且只有一个 CUDA 内核

density(...)

template __global__ void density<5>(int width, int height);
template __global__ void density<4>(int width, int height);
template __global__ void density<3>(int width, int height);
template __global__ void density<2>(int width, int height);
template __global__ void density<1>(int width, int height);

这些是根据给定值通过

switch-case

机制调用的。内核本身定义为

template<int N>
__global__ void density(int width, int height){...}

我们可以简单地将变量

int N

传递给内核本身，而不是显式定义 5 个不同的模板参数。这样定义的目的是什么？

P.S：

int N

的值只有在运行时才知道。

Answer 1

这些是指示编译器使用该模板配置实例化该函数（生成和编译实际代码，并提供用于链接的函数）的指令。通常的原因是启用外部链接，以便变体可用并编译，以便可以从另一个编译单元调用它们。在编译单元 A 中定义的模板函数，无法在编译单元 B 中配置和编译。如果编译单元 B 要使用该函数，则必须在编译单元 A 中编译出合适的变体。

模板参数本身的一个典型原因是，代码设计者得出的结论是，如果模板参数在编译时已知（对于模板参数而言），而不是传递模板参数，则编译器可能能够更好地优化代码通过普通函数参数通过运行时参数获得的数字。

之所以专门模板实例化这5种类型，是因为代码设计者确定这些是实践中需要的，或者基于对外部链接需求的研究。

即使参数 N 仅在运行时已知，如果范围有限，也可以实例化每个变体，并获得编译时优化优势（如果有）。

在这种情况下，可以根据运行时参数使用 switch-case 来分派各种变体：

switch(runtime_N) {
  case 1:
    density<1><<<...>>>(width, height);
    break;
  case 2:
    density<2><<<...>>>(width, height);
    break;
  case 3:
    density<3><<<...>>>(width, height);
    break;
  case 4:
    density<4><<<...>>>(width, height);
    break;
  case 5:
    density<5><<<...>>>(width, height);
    break;
  default:
    density_unoptimized_for_N(width, height, runtime_N);
    break;
}

即使 N 的运行时分布超出范围 1..5，如果代码设计者确定最频繁出现的用法将在 1..5 范围内，则上述 switch-case 结构允许程序提取将 N 视为编译时常量的“大部分”好处，即使它不是。

这些都不是 CUDA 所独有或特定的，而是代表通用 C++。