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

Question

我有这个（工作）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; 
} 

但设备函数指针的分配不工作。 错误在哪里？以及如何纠正？ 非常感谢您提前。 Michele

Answer 1

希望这会帮助别人

#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; 
} 

Answer 2

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

Answer 3

按照CUDA C Programming Guide，

D.2.4.3函数指针

函数指针到__global__功能在主代码的支持，而不是在设备的代码。

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

在主机代码中不允许使用__device__函数的地址。

我的猜测是你编译的计算能力低于2.0。