c++ - 内核参数的 CUDA C++ 模板化

Question

我正在尝试基于 bool 变量对 CUDA 内核进行模板化(如此处所示:Should I unify two similar kernels with an 'if' statement, risking performance loss?)，但我不断收到编译器错误，提示我的函数不是模板。我认为我只是遗漏了一些明显的东西，所以非常令人沮丧。

以下不起作用:

实用工具

#include "kernels.cuh"
//Utility functions

内核.cuh

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

内核.cu

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

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

主.cu

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

以下确实有效:

实用工具

#include "kernels.cuh"
//Utility functions

内核.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

主.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中包含的接口(interface)从kernels.cu中分离出来，这是你的第一个版本的代码没有的主要原因编译。因此，您必须在使用模板的任何文件中同时包含接口(interface)和实现，即，您必须在 main.cu 中同时包含 kernels.cuh 和 kernels .cu.

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

有几个关于在 C++ 中使用模板的教程。 An Idiot's Guide to C++ Templates - Part 1 ，除了令人恼火的标题外，还将为您提供有关该主题的逐步介绍。