Cleanest way to call a device member function (of a children object) from a kernel

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.


#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{
  __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{
        __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;
  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);
    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);
    // 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);