0

是否可以让 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);
}
4

2 回答 2

3

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

这更清洁，更容易阅读。

于 2013-03-11T10:04:11.450 回答
-1

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

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

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

于 2013-03-10T22:31:42.827 回答