问题描述
我正在编写一个简单的 cuda 内核，我正在测量 DRAM 访问的时间，我想从 cuda 代码中获取 ptx 和 sass 代码。设备源代码如下：
__global__ void testPtx(int *devBuff,float *devDummy,unsigned int *timeBuff){
unsigned int temp=0;
unsigned int start,end;
volatile unsigned int *tempPtr;
tempPtr = (volatile unsigned int *)&devBuff[0];
start = clock64();
temp=*tempPtr;
__threadfence();
end = clock64();
*devDummy=(float)(1.0/(float)(temp));
*timeBuff = (unsigned int)(end-start);
}
我能够从源代码生成 sass 和 ptx 文件。但是我对生成的文件感到困惑，并且有一些我将在此处描述的问题。
似乎有多种方法可以生成 ptx 和 sass 代码并阅读它们。生成 ptx 和 sass 代码的一种方法是在 nvcc 中使用 --keep 标志，它将生成包含源代码的 ptx 和 sass 的所有中间文件。另一种使用 -cubin 和 -ptx 选项生成文件的方法，同时使用 nvcc 编译源代码。可以使用 nvdisasm 和 cuobjdump 等 2 个二进制工具生成 sass 代码。反汇编 nvdisasm 只能使用 cubin 文件，而 cuobjdump 可以使用主机二进制文件。
我正在使用 nvcc is nvcc -O0 -o binfile -m64 -gencode arch=compute_60,code=sm_60 --verbose --resource-usage sourcefile.cu 生成主机二进制文件。我试图在编译阶段放置 -cubin 选项，但没有生成 cubin 文件（例如nvcc -O0 -o binfile -m64 -cubin -gencode arch=compute_60,code=sm_60 --verbose --resource-usage sourcefile.cu）。所以我使用nvcc生成cubin文件的方法是nvcc -cubin sourcefile.cu。但是，可以通过在主编译阶段（例如ptx）放置 -ptx 标志来生成 nvcc -O0 -o binfile -m64 -ptx -gencode arch=compute_60,code=sm_60 --verbose --resource-usage sourcefile.cu 文件。对于反汇编和提取 sass，生成的 cubin 文件可以与 nvdisasm 一起使用，主机二进制文件（binfile）可以使用 cuobjdump。但是，使用nvcc -cubin sourcefile.cu然后使用nvdisasm -c sourcefile.cubin生成的sass代码与使用cuobjdump工具（cuobjdump -sass binfile）生成的sass代码不同。从 sass 生成的 nvdisasm 代码提供 here，使用 cuobjdump 生成的代码提供 here。我很困惑为什么 sass 代码不同，如果我在这里做错了什么。我想同时使用这两种二进制工具（最好更多地使用 nvdisasm），但我想确保我生成的 sass 代码对应于源代码，而不是它的不同变体。另外我想知道我是否可以在编译源代码的同时生成 cubin 文件，而不是像我在这里提到的那样单独生成它。我的目标是生成包含所有编译标志（类似于 nvcc -O0 -o binfile -m64 -cubin -gencode arch=compute_60,code=sm_60 --verbose --resource-usage sourcefile.cu）的 cubin 文件，该文件应对应于源代码和 ptx 代码。对于我的工作，我在 ubuntu 18.04 上使用 pascal gpu。如果我遗漏了任何细节，或者我的帖子需要更多解释，请告诉我。谢谢。
解决方法
我在这里看到的唯一重要的事情是确保您的拱形设置匹配。这是我看到的：
$ cat t39.cu
__global__ void testPtx(int *devBuff,float *devDummy,unsigned int *timeBuff){
unsigned int temp=0;
unsigned int start,end;
volatile unsigned int *tempPtr;
tempPtr = (volatile unsigned int *)&devBuff[0];
start = clock64();
temp=*tempPtr;
__threadfence();
end = clock64();
*devDummy=(float)(1.0/(float)(temp));
*timeBuff = (unsigned int)(end-start);
}
$ nvcc -c t39.cu
$ cuobjdump -sass t39.o
Fatbin elf code:
================
arch = sm_52
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
code for sm_52
Function : _Z7testPtxPiPfPj
.headerflags @"EF_CUDA_SM52 EF_CUDA_PTX_SM(EF_CUDA_SM52)"
/* 0x001fc400fec007f6 */
/*0008*/ MOV R1,c[0x0][0x20] ; /* 0x4c98078000870001 */
/*0010*/ CS2R R0,SR_CLOCKLO ; /* 0x50c8000005070000 */
/*0018*/ MOV R2,c[0x0][0x140] ; /* 0x4c98078005070002 */
/* 0x001ffc00162007f2 */
/*0028*/ MOV R3,c[0x0][0x144] ; /* 0x4c98078005170003 */
/*0030*/ LDG.E.CV R3,[R2] ; /* 0xeed4e00000070203 */
/*0038*/ MEMBAR.GL.IVALLD ; /* 0xef98000000070101 */
/* 0x003fd820e3a00ff6 */
/*0048*/ CS2R R2,SR_CLOCKLO ; /* 0x50c8000005070002 */
/*0050*/ I2F.F32.U32 R3,R3 ; /* 0x5cb8000000370a03 */
/*0058*/ IADD32I R4,R3,0x1800000 ; /* 0x1c00180000070304 */
/* 0x001ff400fda007f6 */
/*0068*/ LOP32I.AND R4,R4,0x7f800000 ; /* 0x0407f80000070404 */
/*0070*/ ISETP.GT.U32.AND P0,PT,c[0x2][0x0],PT ; /* 0x4b68038800070407 */
/*0078*/ @P0 BRA 0xa0 ; /* 0xe24000000200000f */
/* 0x001ff400fe0007fd */
/*0088*/ CAL 0x120 ; /* 0xe260000009000040 */
/*0090*/ { MOV R7,R4 ; /* 0x5c98078000470007 */
/*0098*/ BRA 0xd0 }
/* 0xe24000000307000f */
/* 0x001fd801fec0071d */
/*00a8*/ MUFU.RCP R4,R3 ; /* 0x5080000000470304 */
/*00b0*/ FFMA R5,c[0x2][0x4] ; /* 0x5180020800170305 */
/*00b8*/ FADD.FTZ R5,-R5,-RZ ; /* 0x5c5930000ff70505 */
/* 0x001fc400fe2007f6 */
/*00c8*/ FFMA R7,R5,R4 ; /* 0x5980020000570407 */
/*00d0*/ IADD R6,-R0,R2 ; /* 0x5c12000000270006 */
/*00d8*/ MOV R2,c[0x0][0x148] ; /* 0x4c98078005270002 */
/* 0x001fc400fe0007f2 */
/*00e8*/ MOV R3,c[0x0][0x14c] ; /* 0x4c98078005370003 */
/*00f0*/ { MOV R4,c[0x0][0x150] ; /* 0x4c98078005470004 */
/*00f8*/ STG.E [R2],R7 }
/* 0xeedc200000070207 */
/* 0x001ffc00fe2007f2 */
/*0108*/ MOV R5,c[0x0][0x154] ; /* 0x4c98078005570005 */
/*0110*/ STG.E [R4],R6 ; /* 0xeedc200000070406 */
/*0118*/ EXIT ; /* 0xe30000000007000f */
/* 0x001fb400fec007f6 */
/*0128*/ SHL R4,0x1 ; /* 0x3848000000170304 */
/*0130*/ SHR.U32 R4,0x18 ; /* 0x3828000001870404 */
/*0138*/ ISETP.NE.U32.AND P0,RZ,PT ; /* 0x5b6a03800ff70407 */
/* 0x001fb400fec007fd */
/*0148*/ @P0 BRA 0x1c0 ; /* 0xe24000000700000f */
/*0150*/ SHL R4,0x1 ; /* 0x3848000000170304 */
/*0158*/ ISETP.NE.AND P0,PT ; /* 0x5b6b03800ff70407 */
/* 0x001fc801ffa00712 */
/*0168*/ @!P0 MUFU.RCP R4,R3 ; /* 0x5080000000480304 */
/*0170*/ @!P0 RET ; /* 0xe32000000008000f */
/*0178*/ FFMA R3,1.84467440737095516160e+19,RZ ; /* 0x32807fdf80070303 */
/* 0x001fd801fec0071d */
/*0188*/ MUFU.RCP R4,R3 ; /* 0x5080000000470304 */
/*0190*/ FFMA R5,c[0x2][0x4] ; /* 0x5180020800170305 */
/*0198*/ FADD.FTZ R5,-RZ ; /* 0x5c5930000ff70505 */
/* 0x001ffc00fe0007f6 */
/*01a8*/ FFMA R4,R4 ; /* 0x5980020000570404 */
/*01b0*/ { FFMA R4,RZ ; /* 0x32807fdf80070404 */
/*01b8*/ RET }
/* 0xe32000000007000f */
/* 0x001ff400fda007f6 */
/*01c8*/ IADD32I R5,-0xfd ; /* 0x1c0ffffff0370405 */
/*01d0*/ ISETP.GT.U32.AND P0,0x1,PT ; /* 0x3668038000170507 */
/*01d8*/ @P0 BRA 0x300 ; /* 0xe24000001200000f */
/* 0x001fd000fe2007f1 */
/*01e8*/ LOP32I.AND R11,0x7fffff ; /* 0x040007fffff7030b */
/*01f0*/ MOV32I R10,0x3 ; /* 0x010000000037f00a */
/*01f8*/ IADD32I R4,-0xfc ; /* 0x1c0ffffff0470404 */
/* 0x001c7400fe0007f2 */
/*0208*/ LOP32I.OR R6,R11,0x3f800000 ; /* 0x0423f80000070b06 */
/*0210*/ { SHL R10,R10,R5 ; /* 0x5c48000000570a0a */
/*0218*/ MUFU.RCP R7,R6 }
/* 0x5080000000470607 */
/* 0x381fc400fcc00ff6 */
/*0228*/ FFMA R8,R6,R7,c[0x2][0x4] ; /* 0x5180038800170608 */
/*0230*/ FADD.FTZ R8,-R8,-RZ ; /* 0x5c5930000ff70808 */
/*0238*/ FFMA.RM R9,R7.reuse,R8.reuse,R7.reuse ; /* 0x5988038000870709 */
/* 0x001fd440fe2007f5 */
/*0248*/ FFMA.RP R8,R8,R7 ; /* 0x5990038000870708 */
/*0250*/ LOP32I.AND R7,R9.reuse,0x7fffff ; /* 0x040007fffff70907 */
/*0258*/ FSET.NEU.FTZ.AND R8,R9,PT ; /* 0x588d038000870908 */
/* 0x001fc400fca007f1 */
/*0268*/ LOP32I.OR R7,0x800000 ; /* 0x0420080000070707 */
/*0270*/ IADD R8,RZ ; /* 0x5c1200000ff70808 */
/*0278*/ LOP.AND R10,R7 ; /* 0x5c47000000770a0a */
/* 0x001fd800fe8407f1 */
/*0288*/ LOP3.LUT.NZ P1,R5.reuse,0xf8 ; /* 0x5be103bf805708ff */
/*0290*/ SHR.U32 R4,R4 ; /* 0x5c28000000470704 */
/*0298*/ SHR.U32 R10,R5 ; /* 0x5c28000000570a0a */
/* 0x001f8400fd8207f1 */
/*02a8*/ LOP.AND.NZ P2,R10.reuse,0x2 ; /* 0x3842300000270aff */
/*02b0*/ LOP.AND.NZ P0,0x1 ; /* 0x3840300000170aff */
/*02b8*/ PSETP.OR.AND P1,P1,P2,PT ; /* 0x509003804107100f */
/* 0x001ff400fda007ec */
/*02c8*/ ISETP.EQ.U32.AND P2,PT ; /* 0x5b6403800ff70b17 */
/*02d0*/ PSETP.AND.AND P0,P0,PT ; /* 0x5090038020070007 */
/*02d8*/ @P0 IADD32I R4,0x1 ; /* 0x1c00000000100404 */
/* 0x001ffc00fe0007f6 */
/*02e8*/ @P2 SHL R4,0x1 ; /* 0x3848000000120404 */
/*02f0*/ { LOP3.LUT R4,c[0x2][0x8],0xf8 ; /* 0x02f8018800270404 */
/*02f8*/ RET }
/* 0xe32000000007000f */
/* 0x001ffc01ffe0071d */
/*0308*/ MUFU.RCP R4,R3 ; /* 0x5080000000470304 */
/*0310*/ RET ; /* 0xe32000000007000f */
/*0318*/ BRA 0x318 ; /* 0xe2400fffff87000f */
/* 0x001f8000fc0007e0 */
/*0328*/ NOP; /* 0x50b0000000070f00 */
/*0330*/ NOP; /* 0x50b0000000070f00 */
/*0338*/ NOP; /* 0x50b0000000070f00 */
..........
Fatbin ptx code:
================
arch = sm_52
code version = [7,2]
producer = <unknown>
host = linux
compile_size = 64bit
compressed
$ nvcc -cubin t39.cu
$ nvdisasm -c t39.cubin
.headerflags @"EF_CUDA_TEXMODE_UNIFIED EF_CUDA_64BIT_ADDRESS EF_CUDA_SM52 EF_CUDA_VIRTUAL_SM(EF_CUDA_SM52)"
.elftype @"ET_EXEC"
//--------------------- .text._Z7testPtxPiPfPj --------------------------
.section .text._Z7testPtxPiPfPj,"ax",@progbits
.sectioninfo @"SHI_REGISTERS=12"
.align 32
.global _Z7testPtxPiPfPj
.type _Z7testPtxPiPfPj,@function
.size _Z7testPtxPiPfPj,(.L_34 - _Z7testPtxPiPfPj)
.other _Z7testPtxPiPfPj,@"STO_CUDA_ENTRY STV_DEFAULT"
_Z7testPtxPiPfPj:
.text._Z7testPtxPiPfPj:
/*0008*/ MOV R1,c[0x0][0x20] ;
/*0010*/ CS2R R0,SR_CLOCKLO ;
/*0018*/ MOV R2,c[0x0][0x140] ;
/*0028*/ MOV R3,c[0x0][0x144] ;
/*0030*/ LDG.E.CV R3,[R2] ;
/*0038*/ MEMBAR.GL.IVALLD ;
/*0048*/ CS2R R2,SR_CLOCKLO ;
/*0050*/ I2F.F32.U32 R3,R3 ;
/*0058*/ IADD32I R4,0x1800000 ;
/*0068*/ LOP32I.AND R4,0x7f800000 ;
/*0070*/ ISETP.GT.U32.AND P0,PT ;
/*0078*/ @P0 BRA `(.L_1) ;
/*0088*/ CAL `($_Z7testPtxPiPfPj$__cuda_sm20_rcp_rn_f32_slowpath) ;
/*0090*/ { MOV R7,R4 ;
/*0098*/ BRA `(.L_2) }
.L_1:
/*00a8*/ MUFU.RCP R4,R3 ;
/*00b0*/ FFMA R5,c[0x2][0x4] ;
/*00b8*/ FADD.FTZ R5,-RZ ;
/*00c8*/ FFMA R7,R4 ;
.L_2:
/*00d0*/ IADD R6,R2 ;
/*00d8*/ MOV R2,c[0x0][0x148] ;
/*00e8*/ MOV R3,c[0x0][0x14c] ;
/*00f0*/ { MOV R4,c[0x0][0x150] ;
/*00f8*/ STG.E [R2],R7 }
/*0108*/ MOV R5,c[0x0][0x154] ;
/*0110*/ STG.E [R4],R6 ;
/*0118*/ EXIT ;
.weak $_Z7testPtxPiPfPj$__cuda_sm20_rcp_rn_f32_slowpath
.type $_Z7testPtxPiPfPj$__cuda_sm20_rcp_rn_f32_slowpath,@function
.size $_Z7testPtxPiPfPj$__cuda_sm20_rcp_rn_f32_slowpath,(.L_34 - $_Z7testPtxPiPfPj$__cuda_sm20_rcp_rn_f32_slowpath)
$_Z7testPtxPiPfPj$__cuda_sm20_rcp_rn_f32_slowpath:
/*0128*/ SHL R4,0x1 ;
/*0130*/ SHR.U32 R4,0x18 ;
/*0138*/ ISETP.NE.U32.AND P0,PT ;
/*0148*/ @P0 BRA `(.L_3) ;
/*0150*/ SHL R4,0x1 ;
/*0158*/ ISETP.NE.AND P0,PT ;
/*0168*/ @!P0 MUFU.RCP R4,R3 ;
/*0170*/ @!P0 RET ;
/*0178*/ FFMA R3,RZ ;
/*0188*/ MUFU.RCP R4,R3 ;
/*0190*/ FFMA R5,c[0x2][0x4] ;
/*0198*/ FADD.FTZ R5,-RZ ;
/*01a8*/ FFMA R4,R4 ;
/*01b0*/ { FFMA R4,RZ ;
/*01b8*/ RET }
.L_3:
/*01c8*/ IADD32I R5,-0xfd ;
/*01d0*/ ISETP.GT.U32.AND P0,PT ;
/*01d8*/ @P0 BRA `(.L_4) ;
/*01e8*/ LOP32I.AND R11,0x7fffff ;
/*01f0*/ MOV32I R10,0x3 ;
/*01f8*/ IADD32I R4,-0xfc ;
/*0208*/ LOP32I.OR R6,0x3f800000 ;
/*0210*/ { SHL R10,R5 ;
/*0218*/ MUFU.RCP R7,R6 }
/*0228*/ FFMA R8,c[0x2][0x4] ;
/*0230*/ FADD.FTZ R8,-RZ ;
/*0238*/ FFMA.RM R9,R7.reuse ;
/*0248*/ FFMA.RP R8,R7 ;
/*0250*/ LOP32I.AND R7,0x7fffff ;
/*0258*/ FSET.NEU.FTZ.AND R8,PT ;
/*0268*/ LOP32I.OR R7,0x800000 ;
/*0270*/ IADD R8,RZ ;
/*0278*/ LOP.AND R10,R7 ;
/*0288*/ LOP3.LUT.NZ P1,0xf8 ;
/*0290*/ SHR.U32 R4,R4 ;
/*0298*/ SHR.U32 R10,R5 ;
/*02a8*/ LOP.AND.NZ P2,0x2 ;
/*02b0*/ LOP.AND.NZ P0,0x1 ;
/*02b8*/ PSETP.OR.AND P1,PT ;
/*02c8*/ ISETP.EQ.U32.AND P2,PT ;
/*02d0*/ PSETP.AND.AND P0,PT ;
/*02d8*/ @P0 IADD32I R4,0x1 ;
/*02e8*/ @P2 SHL R4,0x1 ;
/*02f0*/ { LOP3.LUT R4,0xf8 ;
/*02f8*/ RET }
.L_4:
/*0308*/ MUFU.RCP R4,R3 ;
/*0310*/ RET ;
.L_5:
/*0318*/ BRA `(.L_5) ;
.L_34:
$
这两套SASS是一样的。