4

CUDAマニュアルでは、マルチプロセッサあたりの32ビットレジスタの数を指定しています。それは次のことを意味しますか？

  1. ダブル変数は2つのレジスタを取りますか？

  2. ポインタ変数は2つのレジスタを取りますか？-6 GBのメモリを搭載したFermiの複数のレジスタである必要がありますよね？

  3. 質問2の答えが「はい」の場合は、使用するポインター変数を減らし、intインデックスを増やす方がよいでしょう。

    例：このカーネルコード：

    float* p1;               // two regs
float* p2 = p1 + 1000;   // two regs
int i;                   // one reg
for ( i = 0; i < n; i++ )
{
    CODE THAT USES p1[i] and p2[i]
}

    理論的には、このカーネルコードよりも多くのレジスタが必要です。

    float* p1;               // two regs
int i;                   // one reg
int j;                   // one reg
for ( i = 0, j = 1000; i < n; i++, j++ )
{
    CODE THAT USES p1[i] and p1[j]
}
4

1 に答える 1

9

あなたの3つの質問に対する簡単な答えは次のとおりです。

  1. はい。
  2. はい (コードが 64 ビット ホスト オペレーティング システム用にコンパイルされている場合)。デバイス ポインターのサイズは、常に CUDA のホスト アプリケーションのポインター サイズと一致します。
  3. いいえ。

ポイント 3 を拡張するには、次の 2 つの単純なメモリ コピー カーネルを検討してください。

__global__
void debunk(float *in, float *out, int n)
{
    int i = n * (threadIdx.x + blockIdx.x*blockDim.x);

    for(int j=0; j<n; j++) {
        out[i+j] = in[i+j];
    }
}

__global__
void debunk2(float *in, float *out, int n)
{
    int i = n * (threadIdx.x + blockIdx.x*blockDim.x);
    float *x = in + i;
    float *y = out + i;

    for(int j=0; j<n; j++, x++, y++) {
        *x = *y;
    }
}

あなたの推測では、ローカル整数変数が 2 つしかないのに対し、追加のポインターが 2 つあるため、使用するレジスターを少なくするdebunk 必要があります。debunk2それでも、CUDA 5 リリース ツールチェーンを使用してそれらをコンパイルすると、次のようになります。

$ nvcc -m64 -arch=sm_20 -c -Xptxas="-v"  pointer_size.cu 
ptxas info    : 0 bytes gmem
ptxas info    : Compiling entry function '_Z6debunkPfS_i' for 'sm_20'
ptxas info    : Function properties for _Z6debunkPfS_i
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 8 registers, 52 bytes cmem[0]
ptxas info    : Compiling entry function '_Z7debunk2PfS_i' for 'sm_20'
ptxas info    : Function properties for _Z7debunk2PfS_i
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 8 registers, 52 bytes cmem[0]

それらはまったく同じレジスター数にコンパイルされます。ツールチェーンの出力を逆アセンブルすると、セットアップ コードを除けば、最終的な命令ストリームはほぼ同じであることがわかります。これにはいくつかの理由がありますが、基本的には次の 2 つの単純なルールに帰着します。

  1. C コード (または PTX アセンブラー) からレジスター数を決定しようとしても、ほとんど無駄です。
  2. 非常に洗練されたコンパイラとアセンブラを推測しようとしても、ほとんど無駄です。
于 2013-02-13T18:17:56.197 に答える