cuda 设备功能和模板答案

【问题标题】：cuda device function and templatescuda 设备功能和模板
【发布时间】：2015-08-08 11:00:24
【问题描述】：

我正在使用 CUDA 7 并尝试将函数作为模板参数传递给设备函数，如下所示：

typedef float(*Op)(float, float);

template<typename Op>
__device__ bool is_maxima(float ax, float ay, cudaTextureObject_t current)
{
    // I try to use the passed function as:
    float cv = tex2D<float>(current, ax, ay);    
    float pv = tex2D<float>(current, ax - 1.f, ay);
    if (Op(cv, pv) != cv) return false;
    return true;
}

现在从我的global cuda 内核中，我这样称呼它：

__global__ void detect_keypoints(cudaTextureObject_t current
                                 float *result, int width, int height)
{
    const int x = __mul24(blockDim.x, blockIdx.x) + threadIdx.x;
    const int y = __mul24(blockDim.y, blockIdx.y) + threadIdx.y;

    float c = tex2D<float>(current, ax + 0.5f, ay + 0.5f);
    float ax = x + 1; float av = y + 1;

    if (is_maxima<fmaxf>(ax + 0.5f, ay + 0.5f, current))
        result[y * width + height] = 1.f;
}

但是，它给了我一个编译器错误：

error: no instance of function template "is_maxima" matches the argument 
list
        argument types are: (float, float, cudaTextureObject_t)

在 CUDA 设备函数中不允许将函数作为模板参数传递吗？我的印象是 CUDA 现在支持所有 C++ 功能。

为了完整起见，fmaxf 在 CUDA SDK 中定义为：

inline float fmaxf(float a, float b)
{
    return a > b ? a : b;
}

【问题讨论】：

我们又来了，缺少 MCVE。你有一些已经重现错误的东西。为什么不提供完整的代码？如果你认为这是因为你问的是一个“概念”问题，那你就错了，因为每个调试问题都可以映射成一个概念问题，以避免 MCVE 要求。
错误消息似乎在抱怨参数列表。 current 是有效的 cuda 纹理对象类型吗？此外，您的模板参数是带有大写 O 的 Op，并且您的代码使用带有小 o 的 op。看起来不对。
小“o”是我的一种。我纠正了它。纹理有效。无论如何，我在编译阶段遇到了错误。
好的，我知道了。签名应该是： templatedevice bool is_maxima(float ax, float ay, cudaTextureObject_t current) 和用法 if (op(cv, pv) != cv) return false;
根据我的测试应该是template<Op op> 而不是template<typename Op op>

标签： c++ templates cuda

【解决方案1】：

正如您已经指出的那样，这是不正确的：

template<typename Op>

我认为应该是：

template<Op op>

以下代码对我来说似乎可以正常工作（删除纹理，因为它与问题无关）：

$ cat t755.cu
#include <stdio.h>
#include <math.h>

typedef float(*Op)(float, float);


__device__ float mymax(float a, float b){

  return (a>b)?a:b;
}

template <Op op>
__device__ float test(float d1, float d2){
  return op(d1, d2);
}

__global__ void kernel(){

  float a1 = 1.0f;
  float a2 = 2.0f;
  printf("%f\n", test<mymax>(a1, a2));
}

int main(){

  kernel<<<1,1>>>();
  cudaDeviceSynchronize();
}
$ nvcc -arch=sm_35 -std=c++11 t755.cu -o t755
$ ./t755
2.000000
$

【讨论】：

是的，谢谢。经过多次反复试验，我也明白了这一点。

【解决方案2】：

您可以这样定义is_maxima，以避免typedef：

template<typename Op>
__device__ bool is_maxima(float ax, float ay, cudaTextureObject_t current, Op fun)
{
    float cv = tex2D<float>(current, ax, ay);    
    float pv = tex2D<float>(current, ax - 1.f, ay);
    if (fun(cv, pv) != cv) return false;
    return true;
}

然后使用is_maxima(ax + 0.5f, ay + 0.5f, current, fmaxf) 调用它。

更多详情请查看此答案：https://stackoverflow.com/a/1174193/678093

【讨论】：