将数组/指针作为模板参数传递

Question

我正在尝试创建以下类型的模板函数：

template <bool isHorizontal, float* kernel>
__global__ void smoothFilterColumns(const TwImageCUDA_Device* source, TwImageCUDA_Device* 
destination)
{
// code...
}

（不用担心__global__；这是在 CUDA 上实现的。）

不幸的是，它不允许我像这样创建此函数的实例：

float ptrKernel[] = {1, 2, 1};
smoothFilterColumns<true, ptrKernel>(dxBuffer->cuda_image, dxOutput->cuda_image);

我尝试了各种float*s 和float[]s，有和没有const 修饰符。甚至可以创建这种模板吗？

提前致谢。

注意。内核作为模板参数而不是普通函数参数传递，因为这允许我通过展开循环在 CUDA 中创建更高效​​的代码。

更新 指向浮点的指针在标准 C++ 中用作模板参数，但显然没有办法让它们与 CUDA 设备函数一起使用，因为它们需要指向设备地址的指针，而无法在外部定义这些指针.如果有人让它工作，请告诉我。

Answer 1

我怀疑你会得到它的工作。正如其他人指出的那样，C++ 标准规定作为模板参数传递的任何对象或函数都必须具有外部链接（因此未在当前翻译单元的范围内定义）。问题是 CUDA 目前根本不支持外部链接 - 设备代码中使用的每个符号都必须具有内部链接（即在同一个翻译单元中定义）。这种限制的根本原因是 CUDA 目前没有设备代码的链接器。

Answer 2

我猜你作为模板参数传递的ptrKernel 变量是一个局部变量。无论如何，您可以作为非类型模板参数传递的内容是有限制的。根据 C++ 标准 (14.3.2)，允许以下内容：

• 整数或枚举类型的整数常量表达式
• 非类型模板参数的名称
• 具有外部链接的对象或函数的名称
• 具有外部链接的对象或函数的地址
• 指向成员的指针

确保ptrKernel 变量满足这些要求（同样，我的猜测是它不是具有外部链接的变量，即全局或静态类成员）。

Answer 3

请确保ptrKernel 有外部链接。

// static float ptrKernel[] = { ... };
// ^ won't work.

// const float ptrKernel[] = { ... };
// ^ won't work.

float ptrKernel[] = { ... };
// ^ ok.

void func() {
   // float ptrKernel[] = { ... };
   // ^ won't work (not global variable).
   ...
}

---

这是对非类型模板的限制，如 §[temp.arg.nontype]/1 中所述：

用于非类型、非模板的 模板参数 模板参数 应为以下之一：

• 整数或枚举类型的整数常量表达式；或
• 非类型模板参数的名称；或
• 对象或函数的名称具有外部链接，包括函数模板和函数template-id，但不包括非静态类成员，表示为id 表达式;或
• 具有外部链接的对象或函数的地址，包括函数模板和函数template-id，但不包括非静态类成员，表示为@987654324 @ id-expression 如果名称引用函数或数组，则 & 是可选的；或
• 指向成员的指针，如 5.3.1 中所述。

Answer 4

这可以在 CUDA 中工作吗？

template <bool isHorizontal, class Kernel>
__global__ void smoothFilterColumns(
    const TwImageCUDA_Device* source, TwImageCUDA_Device* destination)
{
    const float *kernel = Kernel::ptr();
    // code...
}

struct Kernel_1_2_1
{
    static const float *ptr()
    {
        static const float kernel[] = {1, 2, 1};
        return kernel;
    }
}

smoothFilterColumns<true, Kernel_1_2_1>(
    dxBuffer->cuda_image, dxOutput->cuda_image);

您也许可以使内核成为struct 的数据成员。并且您可能想要添加一种机制来传递内核大小。

Answer 5

不会工作。您正在尝试将 CPU-RAM 指针传递给 GPU-RAM 内核。

您可以采用不同的方式 1) 是使用多个模板嵌入所有常量值 取决于你的内核的不同长度，否则你 创建一个函子类来处理您要应用的转换的细节：

这是一个让您理解的工作示例。不要忘记设备说明符。

// with 3 int
template<int amount, int k0,int k1, int k2>
__global__ void apply_kernel(const float *input, float *output){


}

// with four int
template<int amount, int k0,int k1, int k2, int k3>
__global__ void apply_kernel(const float *input, float *output){


}

// with five int 
template<int amount, int k0,int k1, int k2, int k3, int k4>
__global__ void apply_kernel(const float *input, float *output){


}

class KernelOperator {
public:
      __host__ __device__ KernelOperator() {
      }
      __host__ __device__ int operator*(int value){
            return value * 2;
      }
};


// with KernelOperator
template<class T>
__global__ void apply_kernel(const float *input, float *output){
           T value;

}

int main(){
    apply_kernel<0, 1,2,1><<<10, 20>>>(NULL,NULL);

    apply_kernel< KernelOperator ><<<10, 20>>>(NULL,NULL);
}