CUDA：将类传递给具有指针函数的类成员的设备

Question

我想编写一个 c++ CUDA 程序，我将一个类传递给内核。该类仅通过调用 operator() 评估内核上的函数。如果我在类中硬连接函数，一切都会按照我的意愿工作。但是，我希望该类具有一定的灵活性，因此我希望该类能够使用不同的函数进行实例化。说是传入一个指针函数。我无法让指针函数实现工作。下面我定义了两个类，一个定义了函数（fixedFunction），另一个带有指向函数的指针（genericFunction）

//Functions.hh
#include <iostream>
#include <stdio.h>

class fixedFunction{
public:
 __host__ fixedFunction() {}
 __host__ __device__ double operator()(double x) {
    return x*x;
 }
};

double f1(double x){
  return x*x;
}

typedef double (*pf) (double var);

class genericFunction{
public:
  __host__ genericFunction(double (*infunc)(double)) : func(infunc){}
  __host__ __device__ double operator()(double x) {
    return func(x);
  }
private:
  pf func;
};

__global__ void kernel1(fixedFunction* g1){
  unsigned int tid = blockIdx.x *blockDim.x + threadIdx.x;
  printf("Func val is: %f\n", (*g1)(tid));
}

__global__ void kernel2(genericFunction* g1){
  unsigned int tid = blockIdx.x *blockDim.x + threadIdx.x;
  printf("Func val is: %f\n", (*g1)(tid));
}

实例化两个类并在主机上运行它们。传递给相关内核，我看到该类调用指针函数的 kernel2 失败

#include "Functions.hh"

int main(){

  fixedFunction h_g1;
  fixedFunction* d_g1;
  cudaMallocManaged(&d_g1, sizeof(h_g1));

  //Host call
  std::cout << h_g1(2.0) << "\n";

  //device call
  kernel1<<<1,32>>>(d_g1);
  cudaDeviceSynchronize();

  genericFunction h_g2(f1);
  genericFunction* d_g2;
  cudaMallocManaged(&d_g2, sizeof(h_g2));

  //Host call
  std::cout << h_g2(3.0) << "\n";

  //device call
  kernel2<<<1,32>>>(d_g2);
  cudaDeviceSynchronize();

我可以看到指针函数中的一个问题可以是任何大小，并且在设备上没有考虑到。那么有没有办法将指针函数传递给一个类并在设备上运行它呢？

谢谢

Answer 1

这大约是我可以对您的代码进行的“最小”更改数量，以使其大致按照您的意图运行。另请注意，关于 CUDA 中的函数指针还有许多其他问题，此答案链接到几个。

1. 装饰f1用__host__ __device__。这是让编译器为其生成设备可调用例程所必需的。否则，仅生成主机代码。

2. 我们需要为f1上面 1中创建的设备可调用版本捕获设备入口地址。有多种方法可以做到这一点。我将使用另一个__device__变量 ( f1_d) “静态地”捕获它，然后将cudaMemcpyFromSymbol其拉入主机代码。

3. 您的genericFunction类被修改为能够保存所需函数的入口点（函数指针）__host__和单独的__device__入口点（函数指针）。此外，根据我们是编译类（__CUDA_ARCH__宏）的主机版本还是设备版本，修改类以选择合适的类，并且修改类构造函数以接受和分配两个入口点。

4. 最后，我们还需要初始化d_g2设备上的对象。在d_g1对象的情况下，该对象没有类数据成员，因此我们可以“摆脱”创建指向的“空”对象，d_g1并且它可以正常工作，因为该对象的类成员函数的入口点已经在设备代码中已知。然而，在 的情况下d_g2，我们通过类数据成员间接访问函数，这些成员是指向函数各自主机和设备版本（入口点）的指针。因此，h_g2在主机代码中初始化对象，并d_g2在设备代码中为对象建立存储后，我们必须在for之后使用 usingd_g2的内容进行初始化。h_g2cudaMemcpycudaMallocManagedd_g2

通过这些更改，您的代码可以按照我的测试编写的那样工作：

$ cat t353.cu
#include <iostream>
#include <stdio.h>

class fixedFunction{
public:
 __host__ fixedFunction() {}
 __host__ __device__ double operator()(double x) {
    return x*x;
 }
};

__host__ __device__ double f1(double x){
  return x*x;
}

typedef double (*pf) (double var);

__device__ pf f1_d = f1;

class genericFunction{
public:
  __host__ genericFunction(double (*h_infunc)(double), double (*d_infunc)(double)) : h_func(h_infunc),d_func(d_infunc){}
  __host__ __device__ double operator()(double x) {
#ifdef __CUDA_ARCH__
    return d_func(x);
#else
    return h_func(x);
#endif
  }
private:
  pf h_func;
  pf d_func;
};

__global__ void kernel1(fixedFunction* g1){
  unsigned int tid = blockIdx.x *blockDim.x + threadIdx.x;
  printf("Func val is: %f\n", (*g1)(tid));
}

__global__ void kernel2(genericFunction* g1){
  unsigned int tid = blockIdx.x *blockDim.x + threadIdx.x;
  printf("Func val is: %f\n", (*g1)(tid));
}

int main(){

  fixedFunction h_g1;
  fixedFunction* d_g1;
  cudaMallocManaged(&d_g1, sizeof(h_g1));

  //Host call
  std::cout << h_g1(2.0) << "\n";

  //device call
  kernel1<<<1,32>>>(d_g1);
  cudaDeviceSynchronize();
  pf d_f1;
  cudaMemcpyFromSymbol(&d_f1, f1_d, sizeof(void*));
  genericFunction h_g2(f1, d_f1);
  genericFunction* d_g2;
  cudaMallocManaged(&d_g2, sizeof(h_g2));
  cudaMemcpy(d_g2, &h_g2, sizeof(h_g2), cudaMemcpyDefault);
  //Host call
  std::cout << h_g2(3.0) << "\n";

  //device call
  kernel2<<<1,32>>>(d_g2);
  cudaDeviceSynchronize();
}
$ nvcc -arch=sm_61 -o t353 t353.cu
$ cuda-memcheck ./t353
========= CUDA-MEMCHECK
4
Func val is: 0.000000
Func val is: 1.000000
Func val is: 4.000000
Func val is: 9.000000
Func val is: 16.000000
Func val is: 25.000000
Func val is: 36.000000
Func val is: 49.000000
Func val is: 64.000000
Func val is: 81.000000
Func val is: 100.000000
Func val is: 121.000000
Func val is: 144.000000
Func val is: 169.000000
Func val is: 196.000000
Func val is: 225.000000
Func val is: 256.000000
Func val is: 289.000000
Func val is: 324.000000
Func val is: 361.000000
Func val is: 400.000000
Func val is: 441.000000
Func val is: 484.000000
Func val is: 529.000000
Func val is: 576.000000
Func val is: 625.000000
Func val is: 676.000000
Func val is: 729.000000
Func val is: 784.000000
Func val is: 841.000000
Func val is: 900.000000
Func val is: 961.000000
9
Func val is: 0.000000
Func val is: 1.000000
Func val is: 4.000000
Func val is: 9.000000
Func val is: 16.000000
Func val is: 25.000000
Func val is: 36.000000
Func val is: 49.000000
Func val is: 64.000000
Func val is: 81.000000
Func val is: 100.000000
Func val is: 121.000000
Func val is: 144.000000
Func val is: 169.000000
Func val is: 196.000000
Func val is: 225.000000
Func val is: 256.000000
Func val is: 289.000000
Func val is: 324.000000
Func val is: 361.000000
Func val is: 400.000000
Func val is: 441.000000
Func val is: 484.000000
Func val is: 529.000000
Func val is: 576.000000
Func val is: 625.000000
Func val is: 676.000000
Func val is: 729.000000
Func val is: 784.000000
Func val is: 841.000000
Func val is: 900.000000
Func val is: 961.000000
========= ERROR SUMMARY: 0 errors
$