Device function pointers as structure elements
I have this (working) processor code:
#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;
}
Now I will try a GPU version like this
#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;
}
but the assignment of device function pointers does not work. Where is the mistake? And how can this be fixed? Thank you very much in advance. Michelle
source to share
Hope this helps someone
#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;
}
source to share
According to the CUDA C Programming Guide ,
D.2.4.3 Function pointers
Function pointers are
__global__
supported in host code but not device code.Function pointers to functions are
__device__
supported only in device code compiled for devices with 2.x computability.You cannot accept a function address
__device__
in the main code.
I am assuming you are compiling for computational power that is below 2.0.
source to share