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

CUDA C++ Templating of Kernel Parameter

本文关键字：参数 C++ 内核 CUDA      更新时间：2023-10-16

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

以下方法不起作用：

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...);

行在内核的末尾，它也可以工作。

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

kernel is not a template
invalid explicit instantiation declaration

如果它有所作为，我会在一行中编译我所有的 .cu 文件，例如：

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

所有显式专用化声明在模板实例化时必须可见。您的明确专业化声明仅在 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 branchn");
    }
    else
    {
        printf("False branchn");
    }
}
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 部分，除了令人恼火的标题外，还将为您提供该主题的分步介绍。

相关文章：