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

Ada*_*27X 9 c++ templates cuda

我试图基于一个布尔变量来模拟一个CUDA内核(如下所示:我应该用'if'语句统一两个类似的内核,冒着性能损失的风险吗?)但是我一直得到一个编译错误,说我的函数是不是模板.我认为我只是遗漏了一些显而易见的东西,所以非常令人沮丧.

以下不起作用:

util.cuh

#include "kernels.cuh"
//Utility functions
Run Code Online (Sandbox Code Playgroud)

kernels.cuh

    #ifndef KERNELS
    #define KERNELS
    template<bool approx>
    __global__ void kernel(...params...);
    #endif
Run Code Online (Sandbox Code Playgroud)

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
Run Code Online (Sandbox Code Playgroud)

main.cu

#include "kernels.cuh"
kernel<false><<<dimGrid,dimBlock>>>(...params...);
Run Code Online (Sandbox Code Playgroud)

以下DOES工作:

util.cuh

#include "kernels.cuh"
//Utility functions
Run Code Online (Sandbox Code Playgroud)

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
Run Code Online (Sandbox Code Playgroud)

main.cu

#include "kernels.cuh"
kernel<false><<<dimGrid,dimBlock>>>(...params...);
Run Code Online (Sandbox Code Playgroud)

如果我扔进去

template __global__ void kernel<false>(...params...);
Run Code Online (Sandbox Code Playgroud)

在kernels.cuh结束时它也有效.

我收到以下错误(均指上面标记的行):

kernel is not a template
invalid explicit instantiation declaration
Run Code Online (Sandbox Code Playgroud)

如果它有所不同,我会在一行中编译所有.cu文件,例如:

nvcc -O3 -arch=sm_21 -I. main.cu kernels.cu -o program
Run Code Online (Sandbox Code Playgroud)

Jac*_*ern 19

在模板实例化时,所有显式特化声明必须是可见的.您的显式专业化声明仅在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;
}
Run Code Online (Sandbox Code Playgroud)

编辑

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

举一个具体的例子,当编译器找到类似的东西时

template<class T>
__global__ void kernel(...params...)
{
    ...
    T a;
    ...
}
Run Code Online (Sandbox Code Playgroud)

它只检查函数语法,但不生成任何对象代码.因此,如果您使用上面的单个模板化函数编译文件,您将拥有一个"空"对象文件.这是合理的,因为编译器不知道分配哪种类型a.

只有在遇到函数模板的显式实例化时,编译器才会生成对象代码.在那一刻,这是模板化函数的编译如何工作,这种行为引入了对多文件项目的限制:模板化函数的实现(定义)必须与其声明在同一个文件中.因此,您不能将分隔的kernels.cuh头文件中包含的接口分开kernels.cu,这是代码的第一个版本无法编译的主要原因.因此,您必须包括接口和实现在使用模板,即你必须在任何文件中main.cu都,kernels.cuhkernels.cu.

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

有几个关于在C++中使用模板的教程.白痴的C++模板指南 - 第1部分,除了令人讨厌的标题之外,还将为您提供该主题的逐步介绍.