볼타의 스레드 별 프로그램 카운터의 레지스터 사용량

Question

18 페이지의 표 2 하단의 각주 뒤에 나오는 의미를 알고 싶습니다. Volta whitepaper. 볼타 전임자 유사한 SM 당 2백56킬로바이트 레지스터를 보유하는 테이블이 나타내고 있지만, 공간이 개선 SIMT 모델의 일부를 형성

당 사용 스레드 프로그램 카운터 (PC)는 일반적 두 필요하다고 언급 스레드 당 슬롯을 등록하십시오.

Volta에서 실행중인 모든 스레드에 대해 PC를 추적하는 2 개의 예약 된 32 비트 레지스터가 있다는 것을 의미합니까? 그렇다면이 예약은 SM에 상주하는 스레드 수에 관계없이 2048 (SM에서 허용되는 최대 스레드 수) * 2 = 4096 개의 레지스터가 사용된다는 의미에서 정적이라는 것을 의미합니까? 또한 7.0보다 낮은 CC에 대해 컴파일하여이 예약을 제거 할 수 있습니까?

Answer 1

실행중인 모든 스레드에 대해 Compute Capability 7.0 용으로 컴파일 할 때 SM의 레지스터 파일에서 2 개의 추가 레지스터가 할당 된 것처럼 보입니다.

CUDA 9.1을 사용하여, 나는 최대 컴파일러 최적화 플래그 (-03)와 함께 다음과 같은 간단한 saxpy 커널에게

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

CC 6.1

및 7.0을 적용 컴파일. 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. 그 아키텍처 만 다시 레지스터 0 (R3 제외) 7 참조

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을 사용하고 있습니다. (어떻게 셔플 명령이 사용되는 지/왜 모르겠습니다). 이제 두 개의 다른 레지스터가있는 곳을 볼 수 없으므로 스레드의 PC를 추적하기 위해 두 개의 레지스터가 예약되어 있다고 생각하게됩니다.

어쨌든 나는 분명히 관찰이 불충분 한 게시물의 시작 부분에서 언급 한 결론에 도달했습니다. 이 답변을 개선하기위한 모든 노력에 감사드립니다.