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。