Pasar la función de host como un puntero de función en la función __global__ O __device__ en CUDA

Actualmente estoy desarrollando una versión de GPU de una función de CPU (por ejemplo, función Calc (int a, int b, doble * c, souble * d, CalcInvFunction GetInv)), en la que una función de host pasa como un puntero de función (por ejemplo, arriba) ejemplo GetInv es la función de host del tipo CalcInvFunction). Mi pregunta es, si tengo que poner la función Calc () completamente en GPU, tengo que pasar la función GetInv como un argumento de puntero a función en la función del dispositivo / función kernel, ¿y eso es posible?

Sí, para una implementación de GPU de Calc , debe pasar el GetInv como un puntero a la función __device__ .

Es posible, aquí hay algunos ejemplos trabajados:

Ex. 1

Ex. 2

Ex. 3

La mayoría de los ejemplos anteriores demuestran llevar el puntero de la función del dispositivo hasta el código de host . Esto puede no ser necesario para su caso particular. Pero debería ser bastante obvio desde arriba cómo agarrar un puntero a la función __device__ (en el código del dispositivo) y usarlo en un kernel.

Finalmente, he podido pasar una función de host como un puntero de función en la función de kernel cuda (función __global__). Gracias a Robert Crovella y njuffa por la respuesta. Pude pasar una función de miembro de clase (función cpu) como un puntero de función a un kernel de cuda. Pero, el problema principal es que solo puedo pasar la función de miembro de clase estática. No puedo pasar la función no declarada como estática. Por ejemplo:

/**/ __host__ __device__ static int CellfunPtr( void*ptr, int a ); /**/

La función anterior funciona porque esta función miembro se declara como función miembro estática. Si no declaro esta función miembro como un miembro estático como, /**/ __host__ __device__ int CellfunPtr( void*ptr, int a ); /**/ /**/ __host__ __device__ int CellfunPtr( void*ptr, int a ); /**/

entonces no funciona.

El código completo tiene cuatro archivos.


  1. Primer archivo

/*start of fundef.h file*/

typedef int (*pFunc_t)(void* ptr, int N);

/*end of fundef.h file*/


  1. Segundo archivo

/*start of solver.h file*/

  class CalcVars { int eqnCount; int numCell; int numTri; int numTet; public: double* cellVel; double* cellPre; /** Constructor */ CalcVars( const int eqnCount_, const int numCell_, const int numTri_, const int numTet_ ); /** Destructor */ ~CalcVars(void); public: void CalcAdv(); __host__ __device__ static int CellfunPtr( void*ptr, int a ); }; 

/*end of solver.h file*/


  1. Tercer archivo

/*start of solver.cu file*/

  #include "solver.h" __device__ pFunc_t pF1_d = CalcVars::CellfunPtr; pFunc_t pF1_h ; __global__ void kernel(int*a, pFunc_t func, void* thisPtr_){ int tid = threadIdx.x; a[tid] = (*func)(thisPtr_, a[tid]); }; /* Constructor */ CalcVars::CalcVars( const int eqnCount_, const int numCell_, const int numTri_, const int numTet_ ) { this->eqnCount = eqnCount_; this->numCell = numCell_; this->numTri = numTri_; this->cellVel = (double*) calloc((size_t) eqnCount, sizeof(double)); this->cellPre = (double*) calloc((size_t) eqnCount, sizeof(double)); } /* Destructor */ CalcVars::~CalcVars(void) { free(this->cellVel); free(this->cellPre); } void CalcVars::CalcAdv( ){ /*int b1 = 0; b1 = CellfunPtr(this, 1);*/ int Num = 50; int *a1, *a1_dev; a1 = (int *)malloc(Num*sizeof(int)); cudaMalloc((void**)&a1_dev, Num*sizeof(int)); for(int i = 0; i >>(a1_dev, pF1_h, this); cudaDeviceSynchronize(); cudaMemcpy(a1, a1_dev, Num*sizeof(int), cudaMemcpyDeviceToHost); }; int CalcVars::CellfunPtr( void* ptr, int a ){ //CalcVars* ClsPtr = (CalcVars*)ptr; printf("Printing from CPU function\n"); //int eqn_size = ClsPtr->eqnCount; //printf("The number is %d",eqn_size); return a-1; }; 

/*end of solver.cu file*/


  1. Cuarto archivo

/*start of main.cpp file*/

  #include "solver.h" int main(){ int n_Eqn, n_cell, n_tri, n_tetra; n_Eqn = 100; n_cell = 200; n_tri = 300; n_tetra = 400; CalcVars* calcvars; calcvars = new CalcVars(n_Eqn, n_cell, n_tri, n_tetra ); calcvars->CalcAdv(); system("pause"); } 

/*end of main.cpp file*/