cuda - convolución - convolucion en python

Pasar la función de host como un puntero de función en__global__ OR__device__ function en CUDA

Question

Actualmente estoy desarrollando una versión de GPU de una función de CPU (por ejemplo, la función Calc (int a, int b, double * c, souble * d, CalcInvFunction GetInv)), en la que una función de host se 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 la GPU, tengo que pasar la función GetInv como un argumento de puntero de función en la función del dispositivo / función del núcleo, y ¿es eso posible?

Answer 1

Finalmente, pude pasar una función de host como puntero de función en la función de núcleo cuda (función __global__). Gracias a Robert Crovella y njuffa por la respuesta. He podido pasar una función de miembro de clase (función de CPU) como puntero de función a un núcleo de Cuda. Pero, el problema principal es que solo puedo pasar la función 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 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*/

2. 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*/

3. 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 <Num; i++){
            a1[i] = i;
        }

        cudaMemcpy(a1_dev, a1, Num*sizeof(int), cudaMemcpyHostToDevice);

        //copy addresses of device functions to host 
        cudaMemcpyFromSymbol(&pF1_h, pF1_d, sizeof(pFunc_t));


        kernel<<<1,42>>>(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*/

4. 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*/

Answer 2

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

Es posible, aquí hay algunos ejemplos trabajados:

Ex. 1

Ex. 2

Ex. 3

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