11

我正在尝试基于布尔变量对 CUDA 内核进行模板化（如下所示：我是否应该使用 'if' 语句统一两个类似的内核，冒着性能损失的风险？），但我不断收到一个编译器错误，提示我的函数是不是模板。我认为我只是遗漏了一些明显的东西，所以这很令人沮丧。

以下不起作用：

实用工具.cuh

#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

主文件

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

以下确实有效：

实用工具.cuh

#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

主文件

#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
4

1 回答 1

23

所有显式特化声明必须在模板实例化时可见。您的显式特化声明仅在 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++ 中使用模板的教程。C++ 模板的白痴指南 - 第 1 部分，除了烦人的标题之外，将为您提供对该主题的逐步介绍。

于 2013-11-08T18:04:10.267 回答