0

CUDAを使用して、ベクトル（私の場合はint2）で構成される配列を使用してコンパイル時の最適化を試みていますが、これをクリーンな方法で実現することはできません。具体的には、2つの定数配列cとwを使用する問題に取り組んでいます。配列wはfloatで構成され、配列cはint2で構成されます。これらの配列は一定であるため、コンパイラにコンパイル時の最適化を実行させ、それによって配列アクセスを効果的に最適化してもらいたいと思います。たとえば、次の2つのデバイス関数の場合、コンパイラはループを展開し、配列アクセスをcとwの値に直接置き換えることで配列アクセスを最適化します。

__forceinline__ __device__ float someFunction1() {
  const int2  c[9] = {make_int2(0, 0), make_int2(1, 0), make_int2(0, 1), make_int2(-1, 0), make_int2(0, -1),
                      make_int2(1, 1), make_int2(-1, 1), make_int2(-1, -1), make_int2(1, -1)};
  const float w[9] = {4.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/36.0f, 1.0f/36.0f, 1.0f/36.0f, 1.0f/36.0f};
  #pragma unroll
  for (int i = 0; i < 9; ++i) {
    //Do something here, accessing c[i] and w[i]
  }
}

__forceinline__ __device__ float someFunction2() {
  const int2  c[9] = {make_int2(0, 0), make_int2(1, 0), make_int2(0, 1), make_int2(-1, 0), make_int2(0, -1),
                           make_int2(1, 1), make_int2(-1, 1), make_int2(-1, -1), make_int2(1, -1)};
  const float w[9] = {4.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/36.0f, 1.0f/36.0f, 1.0f/36.0f, 1.0f/36.0f};
  #pragma unroll
  for (int i = 0; i < 9; ++i) {
    //Do something here, accessing c[i] and w[i]
  }
}

さて、問題は、cとwを使用する各デバイス関数でcとwを継続的に宣言したくないということです。wをグローバルに宣言することはできますが、cをグローバルに宣言することはできません。CUDAではグローバル変数でmake_int2コンストラクターを呼び出すことができないためです。つまり、以下のプログラムは、「デバイス上で空でないコンストラクタまたはデストラクタのコードを生成できません」というエラーを出します。

//Declaring array c like this is not allowed
__device__ const int2  c[9] = {make_int2(0, 0), make_int2(1, 0), make_int2(0, 1), make_int2(-1, 0), make_int2(0, -1),
                               make_int2(1, 1), make_int2(-1, 1), make_int2(-1, -1), make_int2(1, -1)};
__device__ const float w[9] = {4.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/9.0f, 1.0f/36.0f, 1.0f/36.0f, 1.0f/36.0f, 1.0f/36.0f};

__forceinline__ __device__ float someFunction() {
  #pragma unroll
  for (int i = 0; i < 9; ++i) {
    //Do something here, accessing c[i] and w[i]
  }
}

私の質問は、これらの変数にアクセスする各関数でcとwの宣言を防ぎ、必要なコンパイル時の最適化を維持するにはどうすればよいですか？または、別の言い方をすれば、ベクトルの配列をグローバルに宣言するための回避策はありますか？

注意：cとwをグローバルメモリまたは__constant__メモリに格納できることは承知していますが、コンパイル時の最適化はできません。__constant__メモリも、不定期にアクセスすると問題になる可能性があります。

4

1 に答える 1

1

これがコンパイラの最適化に関してあなたが探しているものを本当に可能にするかどうかはわかりませんが、intからint2へのポインタのキャストは私にとってはうまくいくようです：

#include <stdio.h>

  __device__ const int  ci[16] = {0, 0, 1, 0, 0, 1, -1, 0, 0, -1, 1, 1, -1, 1, -1, -1};

  __device__ const int2 *c = (const int2 *)ci;

  __global__ void mykernel(){

  int2 temp = c[1];
  int2 temp1 = c[4];
  printf("c[1].x = %d\n", temp.x);
  printf("c[4].y = %d\n", temp1.y);
  }

int main(){

  mykernel<<<1,1>>>();
  cudaDeviceSynchronize();
  printf("Done\n");
  return 0;
}

constと__定数__は同じものではないことに注意してください。constとの変数定義から宣言を削除することはできますが、宣言をそこに含めるcとci、コンパイラーが目的を達成するのに役立つと思います。

于 2013-03-08T05:54:03.390 に答える