punteros de función de dispositivo

Question

I necesitan una versión del dispositivo del siguiente código 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; 
} 

donde func1, func2, FUNC3 tienen que ser funciones __device__ y "test" tiene que ser un (adecuadamente modificado) __global__ kernel.

Tengo una GeForce GTS 450 (capacidad de cómputo 2.1) gracias de antemano Michele

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

Una solución de trabajo

#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<<<1,1>>>(); 
    test_kernel<<<1,1>>>(); 
    cudaThreadSynchronize(); 

    assign(); 
    test(); 

    return 0; 
} 

Answer 1

punteros de función son permitido en Fermi. Esta es la forma en que podría hacerlo:

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<n;++i){ 
    x=f[i](2.0); 
    printf("%g\n",x); 
    } 
} 

int main(void) 
{ 
    int N = 5; 
    func* h_f; 
    func* d_f; 
    h_f = (func*)malloc(N*sizeof(func)); 
    cudaMalloc((void**)&d_f,N*sizeof(func)); 

    cudaMemcpyFromSymbol(&h_f[0], pfunc1, sizeof(func)); 
    cudaMemcpyFromSymbol(&h_f[1], pfunc1, sizeof(func)); 
    cudaMemcpyFromSymbol(&h_f[2], pfunc2, sizeof(func)); 
    cudaMemcpyFromSymbol(&h_f[3], pfunc3, sizeof(func)); 
    cudaMemcpyFromSymbol(&h_f[4], pfunc3, sizeof(func)); 

    cudaMemcpy(d_f,h_f,N*sizeof(func),cudaMemcpyHostToDevice); 

    test_kernel<<<1,1>>>(d_f,N); 

    cudaFree(d_f); 
    free(h_f); 

    return 0; 
}