【问题标题】:CUDA: passing class to device with a class member that is a pointer functionCUDA:使用作为指针函数的类成员将类传递给设备
【发布时间】:2017-12-08 00:54:28
【问题描述】:

我想编写一个 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();

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

谢谢

【问题讨论】:

  • f1 不是设备功能。无论函数指针设置是否正确,您都无法使用 from s 内核

标签: c++11 cuda


【解决方案1】:

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

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

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

  3. 您的 genericFunction 类已修改为能够同时保存 __host__ 和单独的 __device__ 入口点(函数指针)以用于所需的函数。此外,根据我们是在编译类的主机版本还是设备版本(__CUDA_ARCH__ 宏)修改类以选择正确的类,并修改类构造函数以接受和分配两个入口点。

  4. 最后,我们还需要在设备上初始化d_g2对象。在d_g1 对象的情况下,该对象没有类数据成员,因此我们可以“摆脱”创建一个由d_g1 指向的“空”对象,它可以正常工作,因为它的入口点对象的类成员函数在设备代码中是已知的。但是,在d_g2 的情况下,我们通过类数据成员间接访问函数,这些数据成员是指向函数的相应主机和设备版本(入口点)的指针。因此,在主机代码中初始化h_g2对象,并在设备代码中为d_g2对象建立存储之后,我们必须用h_g2的内容初始化h_g2,在cudaMallocManaged之后使用cudaMemcpy为@ 987654342@.

通过这些更改,您的代码按照我的测试编写:

$ 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
$

【讨论】:

  • 伙计,我当然希望你能得到报酬 :-)
  • 谢谢罗伯特。非常感谢您的帮助!
猜你喜欢
  • 1970-01-01
  • 1970-01-01
  • 2013-08-11
  • 2017-01-19
  • 1970-01-01
  • 2021-07-03
  • 2012-08-01
  • 1970-01-01
  • 1970-01-01
相关资源
最近更新 更多