Cuda单线程作用域变量

Question

是否有可能使cuda使用在函数外声明的单线程作用域变量（寄存器或本地内存）？

我的大部分设备功能需要使用相同的变量。

我不想将与变量相同的变量作为参数传递给我的所有设备函数，我想在函数之外声明变量。

这可能吗？

我的计算能力是1.2。

编辑：一个例子：

__device__ __local__ int id; 
__device__ __local__ int variable1 = 3; 
__device__ __local__ int variable2 = 5; 
__device__ __local__ int variable3 = 8; 
__device__ __local__ int variable4 = 8; 

// 
__device__ int deviceFunction3() { 
    variable1 += 8; 
    variable4 += 7; 
    variable2 += 1; 
    variable3 += id; 

    return variable1 + variable2 + variable3; 
} 

__device__ int deviceFunction2() { 
    variable3 += 8; 
    variable1 += deviceFunction3(); 
    variable4 += deviceFunction3(); 

    return variable3 + variable4; 
} 

__device__ int deviceFunction1() { 
    variable1 += id; 
    variable4 += 2; 
    variable2 += deviceFunction2(); 
    variable3 += variable2 + variable4; 
    return variable1 + variable2 + variable3 + variable4; 
} 

// Kernel 
__global__ void kernel(int *dev_a, int *dev_b, int *dev_c) { 
    id = get_id(); 

    dev_c[id] = deviceFunction1(); 
} 

3个设备功能需要操纵相同的变量。每个变量都依赖于每个线程计算。据我的理解，我不能使用上面的代码，因为我不能声明变量，以便它们对于每个线程都是本地的。

我必须做的，而不是正在申报核函数中的所有变量，然后将指针传递给变量的所有其他功能：

__device__ int deviceFunction3(int* id,int* variable1,int* variable2,int* variable3,int* variable4) { 
    *variable1 += 8; 
    *variable4 += 7; 
    *variable2 += 1; 
    *variable3 += 2; 

    return *variable1 + *variable2 + *variable3; 
} 

__device__ int deviceFunction2(int* id,int* variable1,int* variable2,int* variable3,int* variable4) { 
    *variable3 += 8; 
    *variable1 += deviceFunction3(id,variable1,variable2,variable3,variable4); 
    *variable4 += deviceFunction3(id,variable1,variable2,variable3,variable4); 

    return *variable3 + *variable4; 
} 

__device__ int deviceFunction1(int* id,int* variable1,int* variable2,int* variable3,int* variable4) { 
    *variable1 += *id; 
    *variable4 += 2; 
    *variable2 += deviceFunction2(id,variable1,variable2,variable3,variable4); 
    *variable3 += *variable2 + *variable4; 
    return *variable1 + *variable2 + *variable3 + *variable4; 
} 

// Kernel 
__global__ void kernel(int *dev_a, int *dev_b, int *dev_c) { 
    int id = get_id(); 
    int variable1 = 3; 
    int variable2 = 5; 
    int variable3 = 8; 
    int variable4 = 8; 

    dev_c[id] = deviceFunction1(&id,&variable1,&variable2,&variable3,&variable4); 
} 

Answer 1

您的使用案例是一个真正糟糕的想法，我不会向我最大的敌人推荐这种设计模式。撇开代码的优点了一会儿，因为我在评论中暗示，可以实现线程局部变量的作用域你通过封装__device__函数和变量的愿望，他们依靠的结构，像这样：

struct folly 
{ 
    int id; 
    int variable1; 
    int variable2; 
    int variable3; 
    int variable4; 

    __device__ folly(int _id) { 
     id = _id; 
     variable1 = 3; 
     variable2 = 5; 
     variable3 = 8; 
     variable4 = 8; 
    } 

    __device__ int deviceFunction3() { 
     variable1 += 8; 
     variable4 += 7; 
     variable2 += 1; 
     variable3 += id; 

     return variable1 + variable2 + variable3; 
    } 

    __device__ int deviceFunction2() { 
     variable3 += 8; 
     variable1 += deviceFunction3(); 
     variable4 += deviceFunction3(); 

     return variable3 + variable4; 
    } 

    __device__ int deviceFunction1() { 
     variable1 += id; 
     variable4 += 2; 
     variable2 += deviceFunction2(); 
     variable3 += variable2 + variable4; 
     return variable1 + variable2 + variable3 + variable4; 
    } 
}; 

__global__ void kernel(int *dev_a, int *dev_b, int *dev_c) { 
    int id = threadIdx.x + blockIdx.x * blockDim.x; 
    folly do_calc(id); 
    dev_c[id] = do_calc.deviceFunction1(); 
} 

还要注意，CUDA支持通过引用C++式通，所以已写入第二片的您发布代码的设备功能中的任何一个可以很容易地这样写的：

__device__ int deviceFunction3(int & variable1, int & variable2, 
           int & variable3, int & variable4) 
{ 
    variable1 += 8; 
    variable4 += 7; 
    variable2 += 1; 
    variable3 += 2; 

    return variable1 + variable2 + variable3; 
} 

远远清洁和更容易读书。

Answer 2

我只是想补充一点，我的结论是这不可能。我发现它是CUDA C的一个主要设计问题。

我在某些幻灯片中看到了一个名为__local__的关键字，但我找不到任何文档，并且nvcc也无法识别它。

我想所有应该只有一个线程范围的变量只能在函数内部声明。