我有一个在 CUDA 上运行的简单矩阵乘法内核。
当使用
-lineinfo--ptxas-options -v-lineinfo使用的确切命令：
nvcc -g -G --ptxas-options -v -arch=sm_86 -o mmul_dbg mmul.cu
和
nvcc -lineinfo --ptxas-options -v -arch=sm_86 -o mmul_ncu mmul.cu
我也检查过
nvcc --ptxas-options -v -arch=sm_86 -o mmul_dbg mmul.cu
它产生 20 个寄存器。
__global__ void matrixMul(const int *a, const int *b, int *c, int N) {
// Compute each thread's global row and column index
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
// Iterate over row, and down column
c[row * N + col] = 0;
for (int k = 0; k < N; k++) {
// Accumulate results for a single element
c[row * N + col] += a[row * N + k] * b[k * N + col];
}
}
寄存器数量增加的原因可能是什么？
编辑：nvcc 是 12.3
编辑（2）：删除图像并添加文本输出
$ nvcc --ptxas-options -v -lineinfo -o wlineinfo -arch=sm_86 m mul.cu
ptxas info : 0 bytes gmem
ptxas info : Compiling entry function '_Z9matrixMulPKiS0_Pii' for 'sm_86'
ptxas info : Function properties for _Z9matrixMulPKiS0_Pii
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 28 registers, 380 bytes cmem[0]
$ nvcc --ptxas-options -v -g -G -o wlineinfo -arch=sm_86 mmul.cu
ptxas info : 0 bytes gmem
ptxas info : Compiling entry function '_Z9matrixMulPKiS0_Pii' for 'sm_86'
ptxas info : Function properties for _Z9matrixMulPKiS0_Pii
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 20 registers, 380 bytes cmem[0]
造成差异的原因是使用了
-G在调试模式下，许多/大多数优化被禁用。编译器可能使用但被禁用的一种优化是循环展开。
在非
-G因此，由于非
-Gcuobjdump-G# cuobjdump -sass wlineinfo
Fatbin elf code:
================
arch = sm_86
code version = [1,7]
host = linux
compile_size = 64bit
identifier = t128.cu
code for sm_86
Function : _Z9matrixMulPKiS0_Pii
.headerflags @"EF_CUDA_TEXMODE_UNIFIED EF_CUDA_64BIT_ADDRESS EF_CUDA_SM86 EF_CUDA_VIRTUAL_SM(EF_CUDA_SM86)"
/*0000*/ MOV R1, c[0x0][0x28] ; /* 0x00000a0000017a02 */
/* 0x000fc40000000f00 */
/*0010*/ S2R R9, SR_CTAID.Y ; /* 0x0000000000097919 */
/* 0x000e220000002600 */
/*0020*/ MOV R7, 0x4 ; /* 0x0000000400077802 */
/* 0x000fe20000000f00 */
/*0030*/ ULDC.64 UR4, c[0x0][0x118] ; /* 0x0000460000047ab9 */
/* 0x000fe40000000a00 */
/*0040*/ S2R R0, SR_TID.Y ; /* 0x0000000000007919 */
/* 0x000e280000002200 */
/*0050*/ S2R R8, SR_CTAID.X ; /* 0x0000000000087919 */
/* 0x000e680000002500 */
/*0060*/ S2R R3, SR_TID.X ; /* 0x0000000000037919 */
/* 0x000e620000002100 */
/*0070*/ IMAD R9, R9, c[0x0][0x4], R0 ; /* 0x0000010009097a24 */
/* 0x001fe200078e0200 */
/*0080*/ MOV R0, c[0x0][0x178] ; /* 0x00005e0000007a02 */
/* 0x000fc60000000f00 */
/*0090*/ IMAD R9, R9, c[0x0][0x178], RZ ; /* 0x00005e0009097a24 */
/* 0x000fe200078e02ff */
/*00a0*/ ISETP.GE.AND P0, PT, R0, 0x1, PT ; /* 0x000000010000780c */
/* 0x000fe20003f06270 */
/*00b0*/ IMAD R8, R8, c[0x0][0x0], R3 ; /* 0x0000000008087a24 */
/* 0x002fca00078e0203 */
/*00c0*/ IADD3 R2, R8, R9, RZ ; /* 0x0000000908027210 */
/* 0x000fca0007ffe0ff */
/*00d0*/ IMAD.WIDE R2, R2, R7, c[0x0][0x170] ; /* 0x00005c0002027625 */
/* 0x000fca00078e0207 */
/*00e0*/ STG.E [R2.64], RZ ; /* 0x000000ff02007986 */
/* 0x0001e2000c101904 */
/*00f0*/ @!P0 EXIT ; /* 0x000000000000894d */
/* 0x000fea0003800000 */
/*0100*/ IADD3 R4, R0, -0x1, RZ ; /* 0xffffffff00047810 */
/* 0x000fe40007ffe0ff */
/*0110*/ MOV R15, RZ ; /* 0x000000ff000f7202 */
/* 0x000fe40000000f00 */
/*0120*/ ISETP.GE.U32.AND P0, PT, R4, 0x3, PT ; /* 0x000000030400780c */
/* 0x000fe40003f06070 */
/*0130*/ LOP3.LUT R6, R0, 0x3, RZ, 0xc0, !PT ; /* 0x0000000300067812 */
/* 0x000fe400078ec0ff */
/*0140*/ MOV R11, RZ ; /* 0x000000ff000b7202 */
/* 0x000fd20000000f00 */
/*0150*/ @!P0 BRA 0xc80 ; /* 0x00000b2000008947 */
/* 0x000fea0003800000 */
/*0160*/ IADD3 R10, -R6, c[0x0][0x178], RZ ; /* 0x00005e00060a7a10 */
/* 0x000fe20007ffe1ff */
/*0170*/ IMAD.WIDE R4, R9, R7.reuse, c[0x0][0x160] ; /* 0x0000580009047625 */
/* 0x080fe200078e0207 */
/*0180*/ MOV R15, RZ ; /* 0x000000ff000f7202 */
/* 0x000fe40000000f00 */
/*0190*/ ISETP.GT.AND P0, PT, R10, RZ, PT ; /* 0x000000ff0a00720c */
/* 0x000fe20003f04270 */
/*01a0*/ IMAD.WIDE R12, R8, R7, c[0x0][0x168] ; /* 0x00005a00080c7625 */
/* 0x000fe200078e0207 */
/*01b0*/ IADD3 R4, P1, R4, 0x8, RZ ; /* 0x0000000804047810 */
/* 0x000fe40007f3e0ff */
/*01c0*/ MOV R11, RZ ; /* 0x000000ff000b7202 */
/* 0x000fe40000000f00 */
/*01d0*/ IADD3.X R5, RZ, R5, RZ, P1, !PT ; /* 0x00000005ff057210 */
/* 0x000fce0000ffe4ff */
/*01e0*/ @!P0 BRA 0xad0 ; /* 0x000008e000008947 */
/* 0x000fea0003800000 */
/*01f0*/ ISETP.GT.AND P1, PT, R10, 0xc, PT ; /* 0x0000000c0a00780c */
/* 0x000fe40003f24270 */
/*0200*/ PLOP3.LUT P0, PT, PT, PT, PT, 0x80, 0x0 ; /* 0x000000000000781c */
/* 0x000fd60003f0f070 */
/*0210*/ @!P1 BRA 0x7a0 ; /* 0x0000058000009947 */
/* 0x000fea0003800000 */
/*0220*/ PLOP3.LUT P0, PT, PT, PT, PT, 0x8, 0x0 ; /* 0x000000000000781c */
/* 0x000fc40003f0e170 */
/*0230*/ LDG.E R14, [R12.64] ; /* 0x000000040c0e7981 */
/* 0x000ea8000c1e1900 */
/*0240*/ LDG.E R16, [R4.64+-0x8] ; /* 0xfffff80404107981 */
/* 0x000ea4000c1e1900 */
/*0250*/ IMAD R19, R14, R16, R15 ; /* 0x000000100e137224 */
/* 0x004fe400078e020f */
/*0260*/ IMAD.WIDE R14, R0, 0x4, R12 ; /* 0x00000004000e7825 */
/* 0x008fc600078e020c */
/*0270*/ STG.E [R2.64], R19 ; /* 0x0000001302007986 */
/* 0x0003e8000c101904 */
/*0280*/ LDG.E R16, [R14.64] ; /* 0x000000040e107981 */
/* 0x000ea8000c1e1900 */
/*0290*/ LDG.E R17, [R4.64+-0x4] ; /* 0xfffffc0404117981 */
/* 0x000ea4000c1e1900 */
/*02a0*/ IMAD R21, R16, R17, R19 ; /* 0x0000001110157224 */
/* 0x004fc400078e0213 */
/*02b0*/ IMAD.WIDE R16, R0, 0x4, R14 ; /* 0x0000000400107825 */
/* 0x000fc600078e020e */
/*02c0*/ STG.E [R2.64], R21 ; /* 0x0000001502007986 */
/* 0x0005e8000c101904 */
/*02d0*/ LDG.E R18, [R16.64] ; /* 0x0000000410127981 */
/* 0x000ee8000c1e1900 */
/*02e0*/ LDG.E R12, [R4.64] ; /* 0x00000004040c7981 */
/* 0x000ee4000c1e1900 */
/*02f0*/ IMAD R23, R18, R12, R21 ; /* 0x0000000c12177224 */
/* 0x008fc400078e0215 */
/*0300*/ IMAD.WIDE R12, R0, 0x4, R16 ; /* 0x00000004000c7825 */
/* 0x000fc600078e0210 */
/*0310*/ STG.E [R2.64], R23 ; /* 0x0000001702007986 */
/* 0x0007e8000c101904 */
/*0320*/ LDG.E R18, [R12.64] ; /* 0x000000040c127981 */
/* 0x000e68000c1e1900 */
/*0330*/ LDG.E R14, [R4.64+0x4] ; /* 0x00000404040e7981 */
/* 0x000e64000c1e1900 */
/*0340*/ IMAD R19, R18, R14, R23 ; /* 0x0000000e12137224 */
/* 0x002fc400078e0217 */
/*0350*/ IMAD.WIDE R14, R0, 0x4, R12 ; /* 0x00000004000e7825 */
/* 0x000fc600078e020c */
/*0360*/ STG.E [R2.64], R19 ; /* 0x0000001302007986 */
/* 0x0003e8000c101904 */
/*0370*/ LDG.E R18, [R14.64] ; /* 0x000000040e127981 */
/* 0x000ea8000c1e1900 */
/*0380*/ LDG.E R16, [R4.64+0x8] ; /* 0x0000080404107981 */
/* 0x000ea4000c1e1900 */
/*0390*/ IMAD R21, R18, R16, R19 ; /* 0x0000001012157224 */
/* 0x004fc400078e0213 */
/*03a0*/ IMAD.WIDE R16, R0, 0x4, R14 ; /* 0x0000000400107825 */
/* 0x000fc600078e020e */
/*03b0*/ STG.E [R2.64], R21 ; /* 0x0000001502007986 */
/* 0x0005e8000c101904 */
/*03c0*/ LDG.E R18, [R16.64] ; /* 0x0000000410127981 */
/* 0x000ee8000c1e1900 */
/*03d0*/ LDG.E R12, [R4.64+0xc] ; /* 0x00000c04040c7981 */
/* 0x000ee4000c1e1900 */
/*03e0*/ IMAD R23, R18, R12, R21 ; /* 0x0000000c12177224 */
/* 0x008fc400078e0215 */
/*03f0*/ IMAD.WIDE R12, R0, 0x4, R16 ; /* 0x00000004000c7825 */
/* 0x000fc600078e0210 */
/*0400*/ STG.E [R2.64], R23 ; /* 0x0000001702007986 */
/* 0x0007e8000c101904 */
/*0410*/ LDG.E R18, [R12.64] ; /* 0x000000040c127981 */
/* 0x000e68000c1e1900 */
/*0420*/ LDG.E R14, [R4.64+0x10] ; /* 0x00001004040e7981 */
/* 0x000e64000c1e1900 */
/*0430*/ IMAD R19, R18, R14, R23 ; /* 0x0000000e12137224 */
/* 0x002fc400078e0217 */
/*0440*/ IMAD.WIDE R14, R0, 0x4, R12 ; /* 0x00000004000e7825 */
/* 0x000fc600078e020c */
/*0450*/ STG.E [R2.64], R19 ; /* 0x0000001302007986 */
/* 0x0003e8000c101904 */
/*0460*/ LDG.E R18, [R14.64] ; /* 0x000000040e127981 */
/* 0x000ea8000c1e1900 */
/*0470*/ LDG.E R16, [R4.64+0x14] ; /* 0x0000140404107981 */
/* 0x000ea4000c1e1900 */
/*0480*/ IMAD R21, R18, R16, R19 ; /* 0x0000001012157224 */
/* 0x004fc400078e0213 */
/*0490*/ IMAD.WIDE R16, R0, 0x4, R14 ; /* 0x0000000400107825 */
/* 0x000fc600078e020e */
/*04a0*/ STG.E [R2.64], R21 ; /* 0x0000001502007986 */
/* 0x0005e8000c101904 */
/*04b0*/ LDG.E R18, [R16.64] ; /* 0x0000000410127981 */
/* 0x000ee8000c1e1900 */
/*04c0*/ LDG.E R12, [R4.64+0x18] ; /* 0x00001804040c7981 */
/* 0x000ee4000c1e1900 */
/*04d0*/ IMAD R23, R18, R12, R21 ; /* 0x0000000c12177224 */
/* 0x008fc400078e0215 */
/*04e0*/ IMAD.WIDE R12, R0, 0x4, R16 ; /* 0x00000004000c7825 */
/* 0x000fc600078e0210 */
/*04f0*/ STG.E [R2.64], R23 ; /* 0x0000001702007986 */
/* 0x0007e8000c101904 */
/*0500*/ LDG.E R18, [R12.64] ; /* 0x000000040c127981 */
/* 0x000e68000c1e1900 */
/*0510*/ LDG.E R14, [R4.64+0x1c] ; /* 0x00001c04040e7981 */
/* 0x000e64000c1e1900 */
/*0520*/ IMAD R19, R18, R14, R23 ; /* 0x0000000e12137224 */
/* 0x002fc400078e0217 */
/*0530*/ IMAD.WIDE R14, R0, 0x4, R12 ; /* 0x00000004000e7825 */
/* 0x000fc600078e020c */
/*0540*/ STG.E [R2.64], R19 ; /* 0x0000001302007986 */
/* 0x000fe8000c101904 */
/*0550*/ LDG.E R18, [R14.64] ; /* 0x000000040e127981 */
/* 0x000ea8000c1e1900 */
/*0560*/ LDG.E R16, [R4.64+0x20] ; /* 0x0000200404107981 */
/* 0x000ea4000c1e1900 */
/*0570*/ IMAD R21, R18, R16, R19 ; /* 0x0000001012157224 */
/* 0x004fc400078e0213 */
/*0580*/ IMAD.WIDE R16, R0, 0x4, R14 ; /* 0x0000000400107825 */
/* 0x000fc600078e020e */
/*0590*/ STG.E [R2.64], R21 ; /* 0x0000001502007986 */
/* 0x0003e8000c101904 */
/*05a0*/ LDG.E R18, [R16.64] ; /* 0x0000000410127981 */
/* 0x000ee8000c1e1900 */
/*05b0*/ LDG.E R12, [R4.64+0x24] ; /* 0x00002404040c7981 */
/* 0x000ee4000c1e1900 */
/*05c0*/ IMAD R23, R18, R12, R21 ; /* 0x0000000c12177224 */
/* 0x008fc400078e0215 */
/*05d0*/ IMAD.WIDE R12, R0, 0x4, R16 ; /* 0x00000004000c7825 */
/* 0x000fc600078e0210 */
/*05e0*/ STG.E [R2.64], R23 ; /* 0x0000001702007986 */
/* 0x0005e8000c101904 */
/*05f0*/ LDG.E R18, [R12.64] ; /* 0x000000040c127981 */
/* 0x000ee8000c1e1900 */
/*0600*/ LDG.E R14, [R4.64+0x28] ; /* 0x00002804040e7981 */
/* 0x000ee4000c1e1900 */
/*0610*/ IMAD R25, R18, R14, R23 ; /* 0x0000000e12197224 */
/* 0x008fc400078e0217 */
/*0620*/ IMAD.WIDE R14, R0, 0x4, R12 ; /* 0x00000004000e7825 */
/* 0x000fc600078e020c */
/*0630*/ STG.E [R2.64], R25 ; /* 0x0000001902007986 */
/* 0x0007e8000c101904 */
/*0640*/ LDG.E R18, [R14.64] ; /* 0x000000040e127981 */
/* 0x000e68000c1e1900 */
/*0650*/ LDG.E R16, [R4.64+0x2c] ; /* 0x00002c0404107981 */
/* 0x000e64000c1e1900 */
/*0660*/ IMAD R21, R18, R16, R25 ; /* 0x0000001012157224 */
/* 0x002fc400078e0219 */
/*0670*/ IMAD.WIDE R16, R0, 0x4, R14 ; /* 0x0000000400107825 */
/* 0x000fc600078e020e */
/*0680*/ STG.E [R2.64], R21 ; /* 0x0000001502007986 */
/* 0x0007e8000c101904 */
/*0690*/ LDG.E R18, [R16.64] ; /* 0x0000000410127981 */
/* 0x000ea8000c1e1900 */
/*06a0*/ LDG.E R12, [R4.64+0x30] ; /* 0x00003004040c7981 */
/* 0x000ea2000c1e1900 */
/*06b0*/ IADD3 R10, R10, -0x10, RZ ; /* 0xfffffff00a0a7810 */
/* 0x000fe20007ffe0ff */
/*06c0*/ IMAD R23, R18, R12, R21 ; /* 0x0000000c12177224 */
/* 0x004fc400078e0215 */
/*06d0*/ IMAD.WIDE R18, R0, 0x4, R16 ; /* 0x0000000400127825 */
/* 0x000fc600078e0210 */
/*06e0*/ STG.E [R2.64], R23 ; /* 0x0000001702007986 */
/* 0x0007e8000c101904 */
/*06f0*/ LDG.E R12, [R18.64] ; /* 0x00000004120c7981 */
/* 0x000ea8000c1e1900 */
/*0700*/ LDG.E R15, [R4.64+0x34] ; /* 0x00003404040f7981 */
/* 0x0002a2000c1e1900 */
/*0710*/ ISETP.GT.AND P1, PT, R10, 0xc, PT ; /* 0x0000000c0a00780c */
/* 0x000fe40003f24270 */
/*0720*/ IADD3 R14, P2, R4, 0x40, RZ ; /* 0x00000040040e7810 */
/* 0x000fc40007f5e0ff */
/*0730*/ IADD3 R11, R11, 0x10, RZ ; /* 0x000000100b0b7810 */
/* 0x000fe40007ffe0ff */
/*0740*/ IADD3.X R5, RZ, R5, RZ, P2, !PT ; /* 0x00000005ff057210 */
/* 0x002fe400017fe4ff */
/*0750*/ MOV R4, R14 ; /* 0x0000000e00047202 */
/* 0x000fe20000000f00 */
/*0760*/ IMAD R15, R12, R15, R23 ; /* 0x0000000f0c0f7224 */
/* 0x004fe400078e0217 */
/*0770*/ IMAD.WIDE R12, R0, 0x4, R18 ; /* 0x00000004000c7825 */
/* 0x000fc600078e0212 */
/*0780*/ STG.E [R2.64], R15 ; /* 0x0000000f02007986 */
/* 0x0007e2000c101904 */
/*0790*/ @P1 BRA 0x230 ; /* 0xfffffa9000001947 */
在上面清单的末尾，您会发现一系列重复的指令，大致如下：
LDG
LDG
IMAD
IMAD
STG
该重复序列代表展开的循环体。如果您使用
cuobjdump-G