CUDA函数指针

Question

CUDA函数指针

7

我正在尝试在CUDA中编写类似以下内容的东西（实际上我需要编写一些集成函数）

#include <iostream>
using namespace std;

float f1(float x) {
    return x * x;
}

float f2(float x) {
    return x;
}

void tabulate(float p_f(float)) {
    for (int i = 0; i != 10; ++i) {
        std::cout << p_f(i) << ' ';
    }
    std::cout << std::endl;
}

int main() {
    tabulate(f1);
    tabulate(f2);
    return 0;
}

输出：

0 1 4 9 16 25 36 49 64 81
0 1 2 3 4 5 6 7 8 9

我尝试了以下操作，但只收到了错误信息：

错误：不支持函数指针和函数模板参数在sm_1x中使用。

float f1(float x) {
    return x;
}

__global__ void tabulate(float lower, float upper, float p_function(float), float* result) {
    for (lower; lower < upper; lower++) {
        *result = *result + p_function(lower);
    }
}

int main() {
    float res;
    float* dev_res;

    cudaMalloc( (void**)&dev_res, sizeof(float) ) ;

    tabulate<<<1,1>>>(0.0, 5.0, f1, dev_res);
    cudaMemcpy(&res, dev_res, sizeof(float), cudaMemcpyDeviceToHost);

    printf("%f\n", res);
    /************************************************************************/
    scanf("%s");

    return 0;
}

- DanilGholtsman

1

你使用的是哪张显卡？你似乎正在将代码编译为计算能力1.x，而我认为函数指针是计算能力2.x的特性。如果你的显卡支持，你可以更改nvcc调用以具有-gencode arch=compute_20,code=sm_20。 - alrikai

@alrikai GeForce 560Ti - DanilGholtsman

然后，您应该将编译器从1.x更改为2.x，这将消除您的编译错误。但是，您可能仍然会遇到一些运行时问题... - alrikai

@alrikai 哦，好的，但是在1.x中有可能实现这样的功能吗？ - DanilGholtsman

我认为不行，看起来你需要一个指向设备函数的函数指针。根据CUDA编程指南所述：“只有在为计算能力为2.x及更高版本的设备编译的设备代码中才支持指向设备函数的函数指针。” 然而，你的560Ti的计算能力是2.1，因此如果你切换到-gencode arch=compute_20,code=sm_20进行编译，这就变得可行了。 - alrikai

@DanilGholtsman：不使用函数指针。但是您可以使用模板参数和switch-dispatch表：enum F{f1, f2, …}; template<F f> __global__ g(){ switch(f){ case f1: func1(); break; case f2: func2(); break; … }} 然后 void h(F f){ switch(f){ case f1: g<f1><<<…>>>(); break; case f2: g<f2><<<…>>>(); break; … } }。但要注意组合爆炸；如果堆栈足够深，您可能会得到数百兆字节的PTX代码。BT;DT. - datenwolf

3个回答

1

即使您可以编译此代码（请参见@Robert Crovella的答案），但此代码将无法工作。您不能将函数指针从主机代码传递，因为主机编译器无法找出函数地址。

- Eugene

1

这是一个简单的函数指针类，可从我基于this问题编写的内核中调用：

template <typename T>
struct cudaCallableFunctionPointer
{
public:
  cudaCallableFunctionPointer(T* f_)
  {
    T* host_ptr = (T*)malloc(sizeof(T));
    cudaMalloc((void**)&ptr, sizeof(T));

    cudaMemcpyFromSymbol(host_ptr, *f_, sizeof(T));
    cudaMemcpy(ptr, host_ptr, sizeof(T), cudaMemcpyHostToDevice);
    
    cudaFree(host_ptr)
  }

  ~cudaCallableFunctionPointer()
  {
    cudaFree(ptr);
  }

  T* ptr;
};

你可以像这样使用它：

你的文本内容

__device__ double func1(double x)
{
    return x + 1.0f;
}

typedef double (*func)(double x);
__device__ func f_ = func1;



__global__ void test_kernel(func* f)
{
    double x = (*f)(2.0);
    printf("%g\n", x);
}



int main()
{
    cudaCallableFunctionPointer<func> f(&f_);

    test_kernel << < 1, 1 >> > (f.ptr);
}

输出：

- jakob

1

host_ptr 上不是有一个内存泄漏吗？你从未调用 free。为什么不直接将对象放在堆栈上，而是使用 malloc 呢？ - Russell Trahan

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- alrikai · Accepted Answer

为了解决编译错误，编译代码时必须使用-gencode arch=compute_20,code=sm_20作为编译器参数。但是你可能会遇到一些运行时问题：

引用自CUDA编程指南http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#functions “__global__”函数的函数指针在主机代码中受支持，但不受设备代码支持。只有在计算能力为2.x及更高的设备上编译的设备代码才支持对“__device__”函数的函数指针。在主机代码中不允许取一个“__device__”函数的地址。

因此，您可以像这样实现（改编自“FunctionPointers”示例）：

//your function pointer type - returns unsigned char, takes parameters of type unsigned char and float
typedef unsigned char(*pointFunction_t)(unsigned char, float);

//some device function to be pointed to
__device__ unsigned char
Threshold(unsigned char in, float thresh)
{
   ...
}

//pComputeThreshold is a device-side function pointer to your __device__ function
__device__ pointFunction_t pComputeThreshold = Threshold;
//the host-side function pointer to your __device__ function
pointFunction_t h_pointFunction;

//in host code: copy the function pointers to their host equivalent
cudaMemcpyFromSymbol(&h_pointFunction, pComputeThreshold, sizeof(pointFunction_t))

您可以将 h_pointFunction 作为参数传递给您的内核，内核可以使用它来调用您的 __device__ 函数。

然后您可以将 h_pointFunction 作为参数传递给您的内核，内核可以使用它来调用您的 __device__ 函数。

//your kernel taking your __device__ function pointer as a parameter
__global__ void kernel(pointFunction_t pPointOperation)
{
    unsigned char tmp;
    ...
    tmp = (*pPointOperation)(tmp, 150.0)
    ...
}

//invoke the kernel in host code, passing in your host-side __device__ function pointer
kernel<<<...>>>(h_pointFunction);

希望这有些说服力。总之，看起来你需要将你的f1函数改为__device__函数，并遵循类似的过程（typedef不是必需的，但它们确实使代码更好）以便在主机端将其作为有效函数指针传递给内核。我还建议您查看FunctionPointers CUDA示例。