11 
__global__ void add( int *c, const int* a, const int* b )
{
    int x = blockIdx.x;
    int y = blockIdx.y;
    int offset = x + y * gridDim.x;
    c[offset] = a[offset] + b[offset];
}

上記の例では、、、はレジスタに保存されていると思いますがx、yoffset

  • nvcc-Xptxas-vは4 registers, 24+16 bytes smem

  • プロファイラーは4つのレジスターを表示します

  • およびptxファイルのヘッド： 

    .reg .u16 %rh<4>;
.reg .u32 %r<9>;    
.reg .u64 %rd<10>;  
.loc    15  21  0   

$LDWbegin__Z3addPiPKiS1_:   
.loc    15  26  0

レジスターの使い方を明確にできる人はいますか？Fermiでは、レジスタの最大数はスレッドごとに63です。私のプログラムでは、カーネルがあまりにも多くのレジスタを消費する場合をテストしたいと思います（したがって、変数をローカルメモリに自動的に格納する必要があり、パフォーマンスが低下する可能性があります）。次に、この時点で、各スレッドに十分なレジスタがあるように、1つのカーネルを2つに分割できます。SMリソースが並行カーネルに十分であると想定します。

私が正しいかどうかはわかりません。

4

1 に答える 1

16

PTX でのレジスタ割り当ては、カーネルの最終的なレジスタ消費とはまったく関係ありません。PTX は、最終的なマシン コードの中間表現にすぎず、静的単一代入形式を使用します。つまり、PTX の各レジスタは 1 回だけ使用されます。数百のレジスターを持つ PTX の一部は、ほんの数個のレジスターを持つカーネルにコンパイルできます。

レジスターの割り当てはptxas、完全にスタンドアロンのコンパイル パスとして (ドライバーによる静的またはジャストインタイム、またはその両方) によって行われ、入力 PTX で多くのコードの並べ替えと最適化を実行して、スループットを向上させ、レジスターを節約できます。つまり、元の C の変数または PTX のレジスターと、アセンブルされたカーネルの最終的なレジスター数との間には、ほとんどまたはまったく関係がないことがわかります。

nvccアセンブラのレジスタ割り当て動作に影響を与えるいくつかの方法を提供します。__launch_bounds__レジスタの割り当てに影響を与える可能性のあるヒューリスティックなヒントをコンパイラに提供する必要があり、コンパイラ/アセンブラは引数-maxrregcountを取得します (パフォーマンスが低下する可能性があるローカルメモリへのレジスタのスピルという潜在的な犠牲を払って)。volatile キーワードは、nvopen64 ベースのコンパイラの古いバージョンに違いをもたらすために使用され、ローカル メモリ スピル動作に影響を与える可能性があります。ただし、元の C コードまたは PTX アセンブリ言語コードでは、レジスタの割り当てを任意に制御または操作することはできません。

于 2012-07-14T13:15:23.487 に答える