CUDA中的函数指针__constant__内存

Question

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

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

方法

• 定义我的自定义对象struct customObj有一个属性（int type），将效仿的模板的参数。
• 定义一组虚拟函数（Sum()，Subtract()等）以供选择。
• 保持的功能的列表应用（functionsList）和相应的type成员查找（first_types，second_types）在__constant__存储器，使得功能functionsList[i](obj1,obj2)被施加到对象与obj1.type == first_types[i]和obj2.type == second_types[i]。

工作代码
下面的代码已经被编译的Linux x86_64的与CUDA 5.0，在计算能力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字节，所以也许是关于数据对齐，但我只是在这里猜测。此外，自计算能力2.0以来，CUDA应该支持GPU上的函数指针。我错过了__constant__内存中的函数指针是否有特殊限制？

Answer 1

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

请更新/升级到CUDA 6。