设备函数指针作为结构成员答案

【问题标题】：device function pointers as struct members设备函数指针作为结构成员
【发布时间】：2012-06-21 18:12:11
【问题描述】：

我有这个（工作的）CPU 代码：

#define NF 3
int ND;

typedef double (*POT)(double x, double y);

typedef struct {
    POT pot[NF];
} DATAMPOT;

DATAMPOT *datampot;

double func0(double x, double y);
double func1(double x, double y);
double func2(double x, double y);


int main(void)
{
    int i;

    ND=5;
    datampot=(DATAMPOT *)malloc(ND*sizeof(DATAMPOT));

    for(i=0;i<ND;i++){
        datampot[i].pot[0]=func0;
        datampot[i].pot[1]=func1;
        datampot[i].pot[2]=func2;
    }

    return 0;
}

现在我尝试这样的 GPU 版本

#define NF 3
int ND;

typedef double (*POT)(double x, double y);

typedef struct {
    POT pot[NF];
} DATAMPOT;

DATAMPOT *dev_datampot;

__device__ double z_func0(double x, double y);
__device__ double z_func1(double x, double y);
__device__ double z_func2(double x, double y);

__global__ void assign(DATAMPOT *dmp, int n)
{
    int i;

    for(i=0;i<n;i++){
        (dmp+i)->pot[0]=z_func0;
        (dmp+i)->pot[1]=z_func1;
        (dmp+i)->pot[2]=z_func2;
    }

}

int main(void)
{
    int i;

    ND=5;
    cudaMalloc((void**)&dev_datampot,ND*sizeof(DATAMPOT));

    assign<<<1,1>>>(dev_datampot,ND);

    return 0;
}

但是设备函数指针的分配不起作用。错误在哪里？以及如何纠正？非常感谢你提前。米歇尔

【问题讨论】：

更具体地说，它怎么不起作用？编译器会报错吗？

标签： cuda gpu

【解决方案1】：

根据CUDA C Programming Guide，

D.2.4.3 函数指针

主机代码支持指向__global__函数的函数指针，但设备代码不支持。

仅在为计算能力 2.x 的设备编译的设备代码中支持指向 __device__ 函数的函数指针。

不允许在宿主代码中取__device__函数的地址。

我的猜测是您正在为低于 2.0 的计算能力进行编译。

【讨论】：

我使用 GeForce GTS 450，计算能力 2.1。在我做 cudaMalloc((void**)&dev_datampot,ND*sizeof(DATAMPOT));是否可以将成员数组pot的三个函数指针链接到设备函数z_func1，z_func2，z_func3？
@micheletuttafesta：您必须在设备功能中执行此操作，这就是您在示例中所做的。您是否正在为计算能力 2.0 进行编译，例如-arch=sm_20?
抱歉，我的回答很晚 Pedro... 是的，我使用 -arch=sm_20 选项进行编译。但是，我可能已经为我的问题找到了解决方案。我会尽快写出来的

【解决方案2】：

希望这会对某人有所帮助

#define NF 3
int ND;

typedef double (*POT)(double x, double y);

typedef struct {
    POT pot[NF];
} DATAMPOT;

DATAMPOT *dev_datampot;

__device__ double z_func0(double x, double y);
__device__ double z_func1(double x, double y);
__device__ double z_func2(double x, double y);

//Static pointers to the above device functions    
__device__ POT z_func0_pointer=z_func0;  
__device__ POT z_func1_pointer=z_func1;
__device__ POT z_func2_pointer=z_func2;



int main(void)
{
    int i;
    POT pot_pointer;

    ND=5;
    cudaMalloc((void**)&dev_datampot,ND*sizeof(DATAMPOT));

    for(i=0;i<ND;++i){  
     cudaMemcpyFromSymbol( &pot_pointer,z_func0_pointer, sizeof( POT ) );
  cudaMemcpy(&dev_datampot[i].pot[0]),&pot_pointer,sizeof(POT),cudaMemcpyHostToDevice);

     cudaMemcpyFromSymbol( &pot_pointer,z_func1_pointer, sizeof( POT ) );
  cudaMemcpy(&dev_datampot[i].pot[1]),&pot_pointer,sizeof(POT),cudaMemcpyHostToDevice);

     cudaMemcpyFromSymbol( &pot_pointer,z_func2_pointer, sizeof( POT ) );
  cudaMemcpy(&dev_datampot[i].pot[2]),&pot_pointer,sizeof(POT),cudaMemcpyHostToDevice);
    }

    return 0;
}

【讨论】：

【解决方案3】：

您的编译器选项是什么？在计算能力为 1.3 或更低的设备上，设备函数必须内联，因此您不能使用设备函数指针。

【讨论】：