How to pass non-static member function of a class into CUDA kernel function (__global__ function)

1
votes

Finally, i have been able to pass a host function as a function pointer in CUDA kernel function (__global__ function). Thanks to Robert Crovella and njuffa for the answer. I have been able to pass a class member function(cpu function) as a function pointer to a CUDA kernel. But, the main problem is, I can only pass the static class member function. I am not being able to pass the function not declared as static.

My question is: How to pass non static member function into CUDA kernel

For Example:

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

The above function work because this member function is declared as static member function. If i do not declare this member function as a static member as ,

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

then it does not work.

The complete code has four files.

fundef.h

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

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
);

};

solver.cu

#include "solver.h"
#include "fundef.h"
#include <stdio.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;

};

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");

}
1
c++cudagpufunction-pointersgpgpu

1 Answers

2
votes

The type of a member function is different:

typedef int (CalcVars::*MethodPtr)(int N);

__device__ MethodPtr pF1_d = &CalcVars::CellfunPtr;

You would then call it using:

__global__ void kernel(int*a, MethodPtr func, void* thisPtr_)
{
    int tid = threadIdx.x;
    CalcVars* c = ((CalcVars*)thisPtr_);
    a[tid] = (c->*func)(a[tid]);
};

BUT the this pointer you are passing to the kernel is a host pointer:

kernel<<<1,42>>>(a1_dev, pF1_h, this);

This will result in invalid memory access in the kernel.

You would have to pass a device pointer of a CalcVars instance to the kernel in order to make it work.

As requested, a complete compilable example, which is a condensed version of your's and still suffers from the this pointer issue I wrote above.

demo.cu

#include <stdio.h>

struct CalcVars
{
  void  CalcAdv();
  __host__ __device__
   int  CellfunPtr(int a);
};

typedef int (CalcVars::*MethodPtr)(int N);

__device__ MethodPtr pF1_d = &CalcVars::CellfunPtr;
MethodPtr pF1_h;

__global__ void kernel(int* a, MethodPtr func, void* thisPtr_)
{
    int tid = threadIdx.x;
    CalcVars* c = ((CalcVars*)thisPtr_);
    a[tid] = (c->*func)(a[tid]);
};

voidCalcVars::CalcAdv()
{
   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);
    cudaMemcpyFromSymbol(&pF1_h, pF1_d, sizeof(MethodPtr));

    // DON'T pass the host this pointer here in real code
    kernel<<<1,42>>>(a1_dev, pF1_h, this);
    cudaDeviceSynchronize();
    cudaMemcpy(a1, a1_dev, Num*sizeof(int), cudaMemcpyDeviceToHost);
};

int CalcVars::CellfunPtr(int a)
{
    printf("Printing from CPU function\n");
    return a-1;
};

int main()
{
   CalcVars calcvars;
   calcvars.CalcAdv();
}