Hi everyone,
I am trying to call a member function using inheritance. I have been seen that calling a virtual function from a CUDA kernel it is not possible. However, I haven’t found a way of how this can be done. I have parent object called CKernel which has general properties, and I would like to have children objects that inherit those properties and also have specific functions, those functions I would like to call them from a CUDA kernel. I have tried to use nvfunctional with no success and also calling the device virtual function.
What would be the cleanest way to this conserving object orientation?
Thanks in advanced.
Regards.
#include <stdio.h>
#include <cuda.h>
#include <nvfunctional>
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
if (code != cudaSuccess)
{
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
// Kernel that executes on the CUDA device
class Test{
public:
__host__ __device__ float test(){return 1;};
__host__ __device__ virtual nvstd::function<float ()> getFunction()=0;
/*__host__ __device__ nvstd::function<float ()> getFunction(){return [](){
return 1.0f;
};};*/
};
class Child{
public:
__host__ __device__ nvstd::function<float ()> getFunction(){
return [](){
return 12.0f;
};
};
};
__global__ void square_array(float *a, int N, Hijole *test)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
printf("hola\n");
nvstd::function<float()> fn1 = test->getFunction();
float f = fn1();
printf("function returns: %f\n", f);
if (idx<N) a[idx] = a[idx] * a[idx];
}
// main routine that executes on the host
int main(void)
{
int num_gpus;
Child *test = new Child();
Child *d_test;
cudaMalloc((void **) &d_test, sizeof(Child*)); ;
cudaMemcpy(d_test, test, sizeof(Child*), cudaMemcpyHostToDevice);
cudaGetDeviceCount(&num_gpus);
cudaSetDevice(0);
printf("Num gpus: %d\n", num_gpus);
float *a_h, *a_d; // Pointer to host & device arrays
const int N = 10; // Number of elements in arrays
size_t size = N * sizeof(float);
a_h = (float *)malloc(size); // Allocate array on host
cudaMalloc((void **) &a_d, size); // Allocate array on device
// Initialize host array and copy it to CUDA device
for (int i=0; i<N; i++) a_h[i] = (float)i;
cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);
// Do calculation on device:
int block_size = 4;
int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);
square_array <<< n_blocks, block_size >>> (a_d, N, test);
gpuErrchk(cudaDeviceSynchronize());
// Retrieve result from device and store it in host array
cudaMemcpy(a_h, a_d, size, cudaMemcpyDeviceToHost);
// Print results
for (int i=0; i<N; i++) printf("%d %f\n", i, a_h[i]);
// Cleanup
free(a_h); cudaFree(a_d);
}