内核参数的CUDA C++模板

Question

我正在尝试基于布尔变量模板化 CUDA 内核（如下所示：Should I unify two similar kernels with an 'if' statement, risking performance loss?），但我不断收到编译器错误，提示我的函数不是模板。我认为我只是遗漏了一些明显的东西，所以这很令人沮丧。

以下方法不起作用：

util.cuh

#include "kernels.cuh"
//Utility functions

kernels.cuh

    #ifndef KERNELS
    #define KERNELS
    template<bool approx>
    __global__ void kernel(...params...);
    #endif

kernels.cu

template<bool approx>
__global__ void kernel(...params...)
{
    if(approx)
    {
       //Approximate calculation
    }
    else
    {
      //Exact calculation
    }
}

template __global__ void kernel<false>(...params...); //Error occurs here

main.cu

#include "kernels.cuh"
kernel<false><<<dimGrid,dimBlock>>>(...params...);

以下确实有效：

util.cuh

#include "kernels.cuh"
//Utility functions

kernels.cuh

#ifndef KERNELS
#define KERNELS
template<bool approx>
__global__ void kernel(...params...);
template<bool approx>
__global__ void kernel(...params...)
{
    if(approx)
    {
       //Approximate calculation
    }
    else
    {
      //Exact calculation
    }
}
#endif

main.cu

#include "kernels.cuh"
kernel<false><<<dimGrid,dimBlock>>>(...params...);

如果我扔了

template __global__ void kernel<false>(...params...);

kernels.cuh 末尾的一行也可以。

我收到以下错误（均指上面的标记行）：

kernel is not a template
invalid explicit instantiation declaration

如果有什么不同，我会在一行中编译所有 .cu 文件，例如：

nvcc -O3 -arch=sm_21 -I. main.cu kernels.cu -o program

Answer 1

所有显式特化声明必须在模板实例化时可见。您的显式特化声明仅在 kernels.cu 翻译单元中可见，而在 main.cu 中不可见。

以下代码确实可以正常工作（除了在显式实例化指令中添加__global__ 限定符）。

#include<cuda.h>
#include<cuda_runtime.h>
#include<stdio.h>
#include<conio.h>

template<bool approx>
__global__ void kernel()
{
    if(approx)
    {
        printf("True branch\n");
    }
    else
    {
        printf("False branch\n");
    }
}

template __global__ void kernel<false>();

int main(void) {
    kernel<false><<<1,1>>>();
    getch();
    return 0;
}

编辑

在 C++ 中，模板函数在遇到函数的显式实例化之前不会被编译。从这个角度来看，现在完全支持模板的 CUDA 的行为方式与 C++ 完全相同。

举个具体的例子，当编译器发现类似的东西时

template<class T>
__global__ void kernel(...params...)
{
    ...
    T a;
    ...
}

它只检查函数语法，但不产生目标代码。因此，如果您要编译具有上述单个模板函数的文件，您将拥有一个“空”目标文件。这是合理的，因为编译器不知道分配给a 的类型。

编译器仅在遇到函数模板的显式实例化时才会生成目标代码。这就是在那个时候，模板化函数的编译是如何工作的，并且这种行为引入了对多文件项目的限制：模板化函数的实现（定义）必须与其声明在同一个文件中。所以，你不能将kernels.cuh中包含的接口分离到一个与kernels.cu分离的头文件中，这就是你的代码的第一个版本无法编译的主要原因。因此，您必须在使用模板的任何文件中同时包含接口和实现，即，您必须在main.cu 中同时包含kernels.cuh 和kernels.cu。

由于没有显式实例化就不会生成代码，因此编译器允许在项目中多次包含具有声明和定义的同一模板文件，而不会产生链接错误。

有几个关于在 C++ 中使用模板的教程。 An Idiot's Guide to C++ Templates - Part 1，除了烦人的标题，还会为你一步步介绍主题。