ボルタのスレッドごとのプログラムカウンタのレジスタ消費

Question

Volta whitepaperの18ページの表2の末尾にある脚注の意味を知りたいのは興味深かったです。表は、ボルタは前任者と同様SM当たり256キロバイトのレジスタを有することを示すが、フットプリントは、改善されたSIMTモデルの一部を形成する

スレッドごとのプログラム・カウンタ（PC）は、典型的に のうちの2つを必要とすることを言及しますスレッドごとにスロットを登録します。

Voltaの実行中のスレッドごとに、PCを追跡する2つの予約済み32ビットレジスタがありますか？ Yesの場合、この予約は、SM上に存在するスレッドの数に関係なく、2048（SM上で許可されるスレッドの最大数）* 2 = 4096個のレジスタが使用されるという意味で、この予約が静的であることを意味しますか？また、7.0より低いCCのためにコンパイルすることによって、この予約を取り除くことができますか？

Answer 1

実行中のスレッドごとに、Compute Capability 7.0用にコンパイルするときに、SMのレジスタファイルから2つの追加レジスタが割り当てられているようです。 CUDA 9.1を使用して

は、私は最大のコンパイラ最適化フラグと次のような単純なSAXPYカーネルCC 6.1および7.0のための

__global__ void saxpy(float* out, float a, float* x, float* y) { 
    out[ threadIdx.x ] = a * x[ threadIdx.x ] + y[ threadIdx.x ]; 
} 

（-03）が適用されてコンパイル。 CC 6.1のバイナリにcuobjdump -reg-usageを使用しているのに対し、カーネルのすべてのスレッドで8つのレジスタが使用されていることを示しています。CC 7.0のバイナリでは、スレッドごとのレジスタ使用量が10であると報告しています。cuobjdump -sassを使用して、以下は、CC 6.1のバイナリの内容です。インデックス0から7までのすべてが使用されているアーキテクチャ化されたレジスタを見ることができます。

code for sm_61 
      Function : _Z5saxpyPffS_S_ 
    .headerflags @"EF_CUDA_SM61 EF_CUDA_PTX_SM(EF_CUDA_SM61)" 
                       /* 0x083fc400e3e007f6 */ 
    /*0008*/     MOV R1, c[0x0][0x20];       /* 0x4c98078000870001 */ 
    /*0010*/     S2R R0, SR_TID.X;        /* 0xf0c8000002170000 */ 
    /*0018*/     SHL R6, R0.reuse, 0x2;       /* 0x3848000000270006 */ 
                       /* 0x081fc840fec007f5 */ 
    /*0028*/     SHR.U32 R0, R0, 0x1e;       /* 0x3828000001e70000 */ 
    /*0030*/     IADD R2.CC, R6.reuse, c[0x0][0x150];   /* 0x4c10800005470602 */ 
    /*0038*/     IADD.X R3, R0.reuse, c[0x0][0x154];   /* 0x4c10080005570003 */ 
                       /* 0x001f8800eec007f0 */ 
    /*0048*/   {   IADD R4.CC, R6, c[0x0][0x158];     /* 0x4c10800005670604 */ 
    /*0050*/     LDG.E R2, [R2];  }      /* 0xeed4200000070202 */ 
    /*0058*/     IADD.X R5, R0, c[0x0][0x15c];     /* 0x4c10080005770005 */ 
                       /* 0x001fdc00fec00771 */ 
    /*0068*/     LDG.E R4, [R4];        /* 0xeed4200000070404 */ 
    /*0070*/     IADD R6.CC, R6, c[0x0][0x140];     /* 0x4c10800005070606 */ 
    /*0078*/     IADD.X R7, R0, c[0x0][0x144];     /* 0x4c10080005170007 */ 
                       /* 0x001ffc001e2047f2 */ 
    /*0088*/     FFMA R0, R2, c[0x0][0x148], R4;    /* 0x4980020005270200 */ 
    /*0090*/     STG.E [R6], R0;        /* 0xeedc200000070600 */ 
    /*0098*/     EXIT;           /* 0xe30000000007000f */ 
                       /* 0x001f8000fc0007ff */ 
    /*00a8*/     BRA 0xa0;          /* 0xe2400fffff07000f */ 
    /*00b0*/     NOP;           /* 0x50b0000000070f00 */ 
    /*00b8*/     NOP;           /* 0x50b0000000070f00 */ 
    .......................... 

ここでCC 7.0。あなたがもう一度だけ（R3を除く）7にレジスタ0を構築さていることがわかり

code for sm_70 
      Function : _Z5saxpyPffS_S_ 
    .headerflags @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)" 
    /*0000*/    @!PT SHFL.IDX PT, RZ, RZ, RZ, RZ;       /* 0x000000fffffff389 */ 
                        /* 0x000fe200000e00ff */ 
    /*0010*/     MOV R1, c[0x0][0x28];         /* 0x00000a0000017a02 */ 
                        /* 0x000fd00000000f00 */ 
    /*0020*/     S2R R6, SR_TID.X;          /* 0x0000000000067919 */ 
                        /* 0x000e220000002100 */ 
    /*0030*/     MOV R7, 0x4;           /* 0x0000000400077802 */ 
                        /* 0x000fca0000000f00 */ 
    /*0040*/     IMAD.WIDE.U32 R2, R6.reuse, R7.reuse, c[0x0][0x170]; /* 0x00005c0006027625 */ 
                        /* 0x0c1fe400078e0007 */ 
    /*0050*/     IMAD.WIDE.U32 R4, R6, R7, c[0x0][0x178];    /* 0x00005e0006047625 */ 
                        /* 0x000fd000078e0007 */ 
    /*0060*/     LDG.E.SYS R2, [R2];         /* 0x0000000002027381 */ 
                        /* 0x000e2800001ee900 */ 
    /*0070*/     LDG.E.SYS R4, [R4];         /* 0x0000000004047381 */ 
                        /* 0x000e2200001ee900 */ 
    /*0080*/     IMAD.WIDE.U32 R6, R6, R7, c[0x0][0x160];    /* 0x0000580006067625 */ 
                        /* 0x000fe400078e0007 */ 
    /*0090*/     FFMA R0, R2, c[0x0][0x168], R4;      /* 0x00005a0002007a23 */ 
                        /* 0x001fd00000000004 */ 
    /*00a0*/     STG.E.SYS [R6], R0;         /* 0x0000000006007386 */ 
                        /* 0x0001e2000010e900 */ 
    /*00b0*/     EXIT;             /* 0x000000000000794d */ 
                        /* 0x000fea0003800000 */ 
    /*00c0*/     BRA 0xc0;            /* 0xfffffff000007947 */ 
                        /* 0x000fc0000383ffff */ 
    /*00d0*/     NOP;             /* 0x0000000000007918 */ 
                        /* 0x000fc00000000000 */ 
    /*00e0*/     NOP;             /* 0x0000000000007918 */ 
                        /* 0x000fc00000000000 */ 
    /*00f0*/     NOP;             /* 0x0000000000007918 */ 
                        /* 0x000fc00000000000 */ 

は、コードブロック内で使用されています。また、カーネルの最初にRZを使用しています（シャッフル命令がどのように/なぜ使われているのかわかりません）。今では2つのレジスタがどこにあるのかわかりません。スレッドのPCを追跡するために2つのレジスタが予約されていると思う傾向があります。

とにかく、私は明らかに不十分な観測で投稿の冒頭に述べた結論に達しました。この回答を改善するための寄稿は感謝しています。