Function pointers in CUDA constant memory

Question

I have found some strange runtime behaviour while experimenting with function pointers in CUDA.

Goal
My goal is to make my function pointers choose which function to apply to two objects according to an internal property of the latter. In short, I want to emulate C++ templates with a CUDA kernel - without actually using template arguments or switch clauses, but function pointers and class/struct members instead.

Approach

Define my custom objects struct customObj with one property (int type) that will emulate the arguments of a template.
Define a bunch of dummy functions (Sum(), Subtract(), etc) to choose from.
Keep the list of functions to apply (functionsList) and respective type members to look up (first_types, second_types) in __constant__ memory, such that function functionsList[i](obj1,obj2) is applied to objects with obj1.type == first_types[i] and obj2.type == second_types[i].

Working code
The following code has been compiled for Linux x86_64 with CUDA 5.0, on a GPU with compute capability 3.0 (GeForce GTX 670), and works.

#include <stdio.h>
#include <iostream>
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, 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);
   }
}

struct customObj
{
  int type;
  double d;
  // Constructors
  __device__ __host__ customObj() {}
  __device__ __host__ customObj(const int& _type, const double& _d) : type(_type), d(_d) {}
};

typedef void (*function_t)(customObj&, customObj&);
// Define a bunch of functions
__host__ __device__ void Sum(customObj& obj1, customObj& obj2) {printf("Sum chosen! d1 + d2 = %f\n", obj1.d + obj2.d);}
__host__ __device__ void Subtract(customObj& obj1, customObj& obj2) {printf("Subtract chosen! d1 - d2 = %f\n", obj1.d - obj2.d);}
__host__ __device__ void Multiply(customObj& obj1, customObj& obj2) {printf("Multiply chosen! d1 * d2 = %f\n", obj1.d * obj2.d);}

#define ARRAYLENGTH 3
__constant__ int first_type[ARRAYLENGTH] = {1, 2, 3};
__constant__ int second_type[ARRAYLENGTH] = {1, 1, 2};
__constant__ function_t functionsList[ARRAYLENGTH] = {Sum, Sum, Subtract};

// Kernel to loop through functions list
__global__ void choosefunction(customObj obj1, customObj obj2) {
   int i = 0;
   function_t f = NULL;
   do {
     if ((obj1.type == first_type[i]) && (obj2.type == second_type[i])) {
       f = functionsList[i];
       break;
    }
    i++;
  } while (i < ARRAYLENGTH);
  if (f == NULL) printf("No possible interaction!\n");
  else f(obj1,obj2);
}

int main() {
  customObj obj1(1, 5.2), obj2(1, 2.6);
  choosefunction<<<1,1>>>(obj1, obj2);
  gpuErrchk(cudaPeekAtLastError());
  gpuErrchk(cudaDeviceSynchronize()); 

  return 0;
}

The issue
The problem that I have found is that, as soon as I replace the datatype of member int type and related variables and functions (__constant__ int first_types[...] and so on)... the code compiles but stops working!

If I change the datatype from int to char or int8_t, the memory checker throws error 4 on my call to cudaDeviceSynchronize().
If I change the datatype to unsigned short int, I get a hardware stack overflow.

So, is anybody having similar issues when working with __constant__ memory? I really have no clue about what is going on. As far as I know, char and int8_t are built-in types of 1 byte length, while the size of int is 4 bytes, so maybe it is about data alignment, but I'm just guessing here. Besides, CUDA is supposed to support function pointers on the GPU since compute capability 2.0. Are there any special constraints for function pointers in __constant__ memory that I'm missing?

I'm not able to repro the problem on linux with CUDA 6. Here is my example. Perhaps you should modify your code with a unique typedef that shows exactly which elements you are switching from int to some other type. — Robert Crovella

Robert Crovella Robert Crovella · Accepted Answer · 2014-06-25T16:27:37

I was able to reproduce the problem (error 4, unspecified launch failure) on CUDA 5.0 on 64bit RHEL 5.5, but not on CUDA 6.0.

Please update/upgrade to CUDA 6.

Function pointers in CUDA __constant__ memory

1 Answers

Function pointers in CUDA constant memory