0

CUDA はスカラー変数をレジスターに割り当てようとするだろうと思います。Fermi クラスの GPU では、各スレッドに 63 個のレジスターがあります。私のコードは次のようなものです：

__global__ void test20 (double a)
{
    double i1=1.0;
    double i2=2.0;
    double i3=3.0;
    double i4=4.0;
    double i5=5.0;
    double i6=6.0;
    double i7=7.0;
    double i8=8.0;
    double i9=9.0;
    double i10=10.0;
    ...

    a = i1+i2+i3 ... i20
 }

しかし、NVVP を使用してスレッドごとのレジスタ数を確認すると、スレッドごとに 2 つのレジスタしか割り当てられていないことがわかります。変数を 10 に減らしても、割り当てられるレジスタの量は変わりません。なぜこれが起こっているのですか? n 変数がある場合、CUDA は n レジスタを使用します (各変数を単一のレジスタに格納できることを考慮して)。

編集：

アドバイスに従って、次のようにコードを修正しました。

 __global__ void test (double *a)
{
    double reg1;
    double reg2;
    double reg3;
    double reg4;
    double reg5;
    double reg6;
    double reg7;
    double reg8;
    ....till 40
    reg1 = log10f(a[0]);
    reg2 = log10f(a[1]);
    reg3 = log10f(a[2]);
    reg4 = log10f(a[3]);
    reg5 = log10f(a[4]);
    reg6 = log10f(a[5]);
    reg7 = log10f(a[6]);
    reg8 = log10f(a[7]);
    reg9 = log10f(a[8]);
    ....till 40
    a[0] = reg1;
    a[1] = reg2;
    a[2] = reg3;
    a[3] = reg4;
    a[4] = reg5;
    a[5] = reg6;
    a[6] = reg7;
    a[7] = reg8;
   }

memcpyアレイをaホストに戻しています。スレッドごとに 63 個のレジスタすべてが使用されていることがわかりますptxas info : Used 62 registers, 40 bytes cmem[0]。レジスタに収まりきらないほど多くの変数を渡しましたが、ローカル メモリへのスピルは見られません。NVCC はレジスタのみを使用するようにコードを最適化していると思います。

4

1 に答える 1

1

実行時に評価できない式を使用するという @talonmies の提案に従った場合でも、宣言ごとにレジスタを取得できない場合があります (この場合は、double を保持するための 2 つのレジスタ)。期間中、変数を有効にしておく必要がある場合もあります。

__global__ void test20 (double a)
{
    double i1=1.0 * a;
    double i2=2.0 * i1;
    double i3=3.0 * i2;
    double i4=4.0 * i3;
    double i5=5.0 * i4;

    a = i1+i2+i3+i4+i5;

    printf("a = %f = %f + %f + %f + %f + %f\n", a, i1, i2, i3, i4, i5);
}

ブラウザで書いたサンプルコードです。目標は、レジスタに値を保持することです。コンパイラの目標は最小限のレジスタを使用することであるため、このサンプルには実用的なアプリケーションはありません。これの唯一の値は、スコープの期間中変数を有効に保つためのデバッグ用です。

レジスタの使用法を理解したい場合は、cuobjump -sass を使用してカーネルのアセンブリ コードをダンプする必要があります。

于 2012-09-22T00:11:43.787 に答える