CUDA __constant__ 内存中的函数指针

Question

我在 CUDA 中试验函数指针时发现了一些奇怪的运行时行为。

目标
我的目标是让我的函数指针根据后者的内部属性选择将哪个函数应用于两个对象。 简而言之，我想用 CUDA 内核模拟 C++ 模板——实际上不使用模板参数或switch 子句，而是使用函数指针和class/struct 成员。

接近

• 使用一个属性 (int type) 定义我的自定义对象 struct customObj，该属性将模拟模板的参数。
• 定义一堆虚拟函数（Sum()、Subtract() 等）以供选择。
• 在 __constant__ 内存中保留要应用的函数列表 (functionsList) 和要查找的各个 type 成员 (first_types, second_types)，以便函数 functionsList[i](obj1,obj2) 应用于具有obj1.type == first_types[i] 和 obj2.type == second_types[i]。

工作代码
以下代码已针对具有 CUDA 5.0 的 Linux x86_64 在具有计算能力 3.0 (GeForce GTX 670) 的 GPU 上编译，并且可以运行。

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

问题
我发现的问题是，一旦我替换成员int type 和相关变量和函数（__constant__ int first_types[...] 等）的数据类型......代码编译但停止工作！

• 如果我将数据类型从int 更改为char 或int8_t，内存检查器会在我调用cudaDeviceSynchronize() 时抛出error 4。
• 如果我将数据类型更改为unsigned short int，则会出现硬件堆栈溢出。

那么，在使用__constant__ 内存时，是否有人遇到过类似的问题？我真的不知道发生了什么。据我所知，char和int8_t是1字节长度的内置类型，而int的大小是4字节，所以可能是关于数据对齐的问题，但我只是在这里猜测。此外，CUDA 应该支持 GPU 上的函数指针，因为计算能力 2.0。我缺少的__constant__ 内存中的函数指针是否有任何特殊限制？

Answer 1

我能够在 64 位 RHEL 5.5 上的 CUDA 5.0 上重现该问题（错误 4，未指定的启动失败），但在 CUDA 6.0 上却没有。

请更新/升级到 CUDA 6。