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++ スタイルの参照渡しをサポートしているため、投稿した 2 番目のコードで記述したデバイス関数のいずれかは、次のように簡単に記述できます。

__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 に答える