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

您的用例是一个非常糟糕的想法，我不会向我最大的敌人推荐这种设计模式。暂时不考虑代码的优点，正如我在 cmets 中所暗示的那样，您可以通过将它们所依赖的 __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 识别。

我猜所有应该只有单个线程范围的变量必须只在函数内部声明。