具有 CUDA 设备函数的函数指针

Question

我想在我的 Cuda C++ 代码中使用函数指针，如下所示，

typedef __device__ void customFunc(const char*, uint64_t, char*, const uint64_t);

这就是我所追求的。没有“__device__”的等效项确实可以很好地工作。

Cuda 支持函数指针吗？

编辑：

我对如何使用 __device__ 函数作为指向 __device__ 函数的函数指针特别感兴趣

Answer 1

在设备代码中使用设备函数指针并没有什么神奇之处。它在功能和语法上与标准 C++ 相同。

例如：

#include <cstdio>

typedef int (*ufunc)(int args);

__device__ int f1(int x)
{
    int res = 2*x;
    printf("f1 arg = %d, res = %d\n", x, res);
    return res;
}

__device__ int f2(int x, int y, ufunc op)
{
    int res = x + op(y);
    printf("f2 arg = %d, %d, res = %d\n", x, y, res);
    return res;
}


__global__ void kernel(int *z) 
{

    int x = threadIdx.x;
    int y = blockIdx.x;
    int tid = threadIdx.x + blockDim.x * blockIdx.x;

    z[tid] = f2(x, y, &f1);
}

int main()
{
    const int nt = 4, nb = 4;
    int* a_d;
    cudaMalloc(&a_d, sizeof(float) * nt *nb);

    kernel<<<nb, nt>>>(a_d);
    cudaDeviceSynchronize();
    cudaDeviceReset();

    return 0;
}
#include <cstdio>

typedef int (*bfunc)(int args);

__device__ int f1(int x)
{
    int res = 2*x;
    printf("f1 arg = %d, res = %d\n", x, res);
    return res;
}

__device__ int f2(int x, int y, bfunc op)
{
    int res = x + f1(y);
    printf("f2 arg = %d, %d, res = %d\n", x, y, res);
    return res;
}


__global__ void kernel(int *z) 
{

    int x = threadIdx.x;
    int y = blockIdx.x;
    int tid = threadIdx.x + blockDim.x * blockIdx.x;

    z[tid] = f2(x, y, &f1);
}

int main()
{
    const int nt = 4, nb = 4;
    int* a_d;
    cudaMalloc(&a_d, sizeof(float) * nt *nb);

    kernel<<<nb, nt>>>(a_d);
    cudaDeviceSynchronize();
    cudaDeviceReset();

    return 0;
}

在这里，我们将一个指向一元仿函数的简单指针定义为一个类型，然后定义一个设备函数，该函数将该类型作为参数。内核调用中函数指针的静态分配在编译时处理，一切正常。如果您想在运行时选择函数指针，那么您需要按照您已经提供的link 中给出的说明进行操作。

这里要记住的重要一点是，在 CUDA 中，在类型定义中包含 CUDA 说明符（__device__、__constant__、__global__ 等）是不合法的。每个变量实例都有一个说明符作为其定义的一部分。