如何将类的非静态成员函数传递给 CUDA 内核函数（__global__ 函数）

Question

最后，我已经能够将主机函数作为函数指针传递给 CUDA 内核函数（__global__ 函数）。感谢 Robert Crovella 和 njuffa 的回答。我已经能够将类成员函数（cpu 函数）作为函数指针传递给 CUDA 内核。但是，主要问题是，我只能传递静态类成员函数。我无法传递未声明为静态的函数。

我的问题是：如何将非静态成员函数传递给 CUDA 内核

例如：

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

上述函数之所以起作用，是因为该成员函数被声明为静态成员函数。如果我不将该成员函数声明为静态成员，

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

那就不行了。

完整的代码有四个文件。

---

fundef.h

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

---

solver.h 文件

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;

};

ma​​in.cpp 文件

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

}

Answer 1

成员函数的类型不同：

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

__device__ MethodPtr pF1_d = &CalcVars::CellfunPtr;

然后你可以使用：

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

但是你传递给内核的this指针是一个host指针：

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

这将导致内核中的内存访问无效。

您必须将 CalcVars 实例的 device 指针传递给内核才能使其工作。

---

根据要求，一个完整的可编译示例，它是您的压缩版本，仍然存在我上面写的this 指针问题。

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