puntatori di funzioni del dispositivo

Ho bisogno di una versione del dispositivo del seguente codice host:

double (**func)(double x); double func1(double x) { return x+1.; } double func2(double x) { return x+2.; } double func3(double x) { return x+3.; } void test(void) { double x; for(int i=0;i<3;++i){ x=func[i](2.0); printf("%g\n",x); } } int main(void) { func=(double (**)(double))malloc(10*sizeof(double (*)(double))); test(); return 0; } 

dove func1, func2, func3 devono essere __device__ funzioni e “test” deve essere un kernel __global__ (opportunamente modificato).

Ho una NVIDIA GeForce GTS 450 (capacità di calcolo 2.1) Grazie in anticipo Michele

================================================== ======

Una soluzione di lavoro

 #define REAL double typedef REAL (*func)(REAL x); __host__ __device__ REAL func1(REAL x) { return x+1.0f; } __host__ __device__ REAL func2(REAL x) { return x+2.0f; } __host__ __device__ REAL func3(REAL x) { return x+3.0f; } __device__ func func_list_d[3]; func func_list_h[3]; __global__ void assign_kernel(void) { func_list_d[0]=func1; func_list_d[1]=func2; func_list_d[2]=func3; } void assign(void) { func_list_h[0]=func1; func_list_h[1]=func2; func_list_h[2]=func3; } __global__ void test_kernel(void) { REAL x; for(int i=0;i<3;++i){ x=func_list_d[i](2.0); printf("%g\n",x); } } void test(void) { REAL x; printf("=============\n"); for(int i=0;i<3;++i){ x=func_list_h[i](2.0); printf("%g\n",x); } } int main(void) { assign_kernel<<>>(); test_kernel<<>>(); cudaThreadSynchronize(); assign(); test(); return 0; } 

puntatori di funzione sono consentiti su Fermi. Ecco come potresti farlo:

 typedef double (*func)(double x); __device__ double func1(double x) { return x+1.0f; } __device__ double func2(double x) { return x+2.0f; } __device__ double func3(double x) { return x+3.0f; } __device__ func pfunc1 = func1; __device__ func pfunc2 = func2; __device__ func pfunc3 = func3; __global__ void test_kernel(func* f, int n) { double x; for(int i=0;i>>(d_f,N); cudaFree(d_f); free(h_f); return 0; }