mardi 4 juillet 2017

CUDA: passing class to device with a class member that is a pointer function

I want to write a c++ CUDA program where I pass a class to the kernel. The class just evaluates a function on the kernel via the call operator(). If I hardwire the function in the class everything works as I'd like. However I want some flexibility with the class so I'd like the class to be able to be instantiated with different functions. Say by passing in a pointer function. I can't get the pointer function implementation to work. Below I define two classes, one that has the function defined (fixedFunction) and another that takes a pointer to function (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));
}

Instantiating both classes and running them on the host works. Passing to the relevant kernels I see that kernel2 where that class calls a pointer function fails

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

I can see an issue in the pointer function can be any size and that is not accounted for on the device. So is there a way to pass a pointer function to a class and run it on the device?

Thanks

Aucun commentaire:

Enregistrer un commentaire