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