It seems that for every running thread, 2 additional registers are allocated from SM's register file when compiling for Compute Capability 7.0.
Using CUDA 9.1, I compiled the following simple saxpy kernel
__global__ void saxpy(float* out, float a, float* x, float* y) {
out[ threadIdx.x ] = a * x[ threadIdx.x ] + y[ threadIdx.x ];
}
for CC 6.1 and 7.0 with maximum compiler optimization flag (-03) applied. While using cuobjdump -res-usage on the binary for CC 6.1 shows that 8 registers are used for every thread in the kernel, the same command on the binary for CC 7.0 reports that register usage per thread is 10. I also printed the sass using cuobjdump -sass. Below is the content for the binary for CC 6.1. You can see architected registers with indices 0 to 7 are all used.
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 */
..........................
Now for CC 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 */
You see that again only architected register 0 to 7 (except for R3 and R5) are used inside the code block. There is also the use of RZ at the beginning of the kernel. Now I do not see where three other registers are, which makes me inclined to believe that two registers are reserved for tracking thread's PC.
Anyway, I came to the conclusion I stated at the beginning of the post with clearly insufficient observations. Any contribution to improve this answer is appreciated.
