CUDA: copy dynamically created array of function pointers to processor to GPU memory

I would like to dynamically create a list of function pointers on the cpu (using some method push_back()

called from main()

) and copy it to a GPU array __constant__

or __device__

, without having to refer to static function pointers __device__

. I believe this question is related to my problem; however, my goal is to create an array of function pointers __host__

iteratively and then copy it into the array of function pointers __constant__

instead of initializing the latter in the declaration.

An example of working code with static function pointers (see here or here ):

common.h:

#ifndef COMMON_H
#define COMMON_H

#include <stdio.h>
#include <iostream>

#define num_functions 3

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

// fptr_t: Pointer to void function that takes two integer lvalues
typedef void (*fptr_t)(int&, int&);

// some examples of void(int&, int&) functions...
__device__ void Add(int &a, int &b) {printf("Add... %i + %i = %i\n", a, b, a+b);}
__device__ void Subtract(int &a, int &b) {printf("Subtract... %i - %i = %i\n", a, b, a-b);}
__device__ void Multiply(int &a, int &b) {printf("Multiply... %i * %i = %i\n", a, b, a*b);}

// List of function pointers in device memory
__constant__ fptr_t constant_fList[num_functions];

// Kernel called from main(): choose the function to apply whose index is equal to thread ID
__global__ void kernel(int a, int b) {
  fptr_t f;
  if (threadIdx.x < num_functions) {
    f = constant_fList[threadIdx.x];
    f(a,b);
  }
}

#endif

      

        
        
        
      

    

main.cu:

#include "common.h"

// Static device function pointers
__device__ fptr_t p_Add = Add;
__device__ fptr_t p_Sub = Subtract;
__device__ fptr_t p_Mul = Multiply;

// Load function list to constant memory
void loadList_staticpointers() {
  fptr_t h_fList[num_functions];
  gpuErrchk( cudaMemcpyFromSymbol(&h_fList[0], p_Add, sizeof(fptr_t)) );
  gpuErrchk( cudaMemcpyFromSymbol(&h_fList[1], p_Sub, sizeof(fptr_t)) );
  gpuErrchk( cudaMemcpyFromSymbol(&h_fList[2], p_Mul, sizeof(fptr_t)) );
  gpuErrchk( cudaMemcpyToSymbol(constant_fList, h_fList, num_functions * sizeof(fptr_t)) );
}

int main() {

  loadList_staticpointers();
  int a = 12, b = 15;
  kernel<<<1,3>>>(a, b);
  gpuErrchk(cudaGetLastError());
  gpuErrchk(cudaDeviceSynchronize());

  return 0;
}

      

        
        
        
      

    

Specs: GeForce GTX 670 compiled for -arch=sm_30

, CUDA 6.5, Ubuntu 14.04

I want to avoid using static device function pointers, since adding each function would require user-side code maintenance - declaring a new static pointer like p_Add

or p_Mul

, manipulating with void loadList_functionpointers()

, etc. To clarify the situation, I am trying to do something like the following (emergency) code:

main_wrong.cu:

#include "common.h"
#include <vector>

// Global variable: list of function pointers in host memory
std::vector<fptr_t> vec_fList;

// Add function to functions list
void addFunc(fptr_t f) {vec_fList.push_back(f);}

// Upload the functions in the std::vector<fptr_t> to GPU memory
// Copies CPU-side pointers to constant_fList, therefore crashes on kernel call 
void UploadVector() {
  fptr_t* h_vpointer = vec_fList.data();
  gpuErrchk( cudaMemcpyToSymbol(constant_fList, h_vpointer, vec_fList.size() * sizeof(fptr_t)) );
}

int main() {

  addFunc(Add);
  addFunc(Subtract);
  addFunc(Multiply);
  int a = 12, b = 15;

  UploadVector();

  kernel<<<1,3>>>(a, b); // Wrong to call a host-side function pointer from a kernel
  gpuErrchk(cudaGetLastError());
  gpuErrchk(cudaDeviceSynchronize());

  return 0;
}

      

        
        
        
      

    

I understand that function pointers pointing to host addresses are copied to the GPU and not valid for the kernel, which requires pointers pointing to GPU addresses when the function is called f(a,b)

. Populating a node-side array with device-side pointers will work for me with raw data (see this question ), but not with function pointers. Trivial attempts with Unified Memory have also failed ... so far I've only found static pointers on the device side. Is there no other way to copy a dynamically created array of arrays of function pointers to the GPU?