nvdis disassembler

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi, I wrote a little extension to Marcin KoÅ›cielnicki’s great tool NV50DIS. This tool, NVDIS, reads the ELF files produced by PTXAS. Here’s an example showing how to use it:

cat basic.ptx

.version 2.1
.target sm_20
.entry InstBasic (
.param .u32 __results
)
{
.reg .u32 %r<2>;

ld.param.u32 %r0,[__results];
mov.u32 %r1, 1;
st.global.b8 [%r0], %r1;
exit;
}

ptxas -arch sm_20 basic.ptx

nvdis elf.o

ELF File…
00000000: 2800440400005de4 mov b32 $r1 c1[0x100]
00000008: 2800400080001de4 mov b32 $r0 c0[0x20]
00000010: 1800000004009de2 mov b32 $r2 0x1
00000018: 9000000000009c05 st u8 wb g[$r0+0] $r2
00000020: 8000000000001de7 exit

cat basic10.ptx

.version 1.4
.target sm_10
.entry InstBasic (
.param .u32 __results
)
{
.reg .u32 %r<2>;

ld.param.u32 %r0,[__results];
mov.u32 %r1, 1;
st.global.b8 [%r0], %r1;
exit;
}

ptxas basic10.ptx

nvdis elf.o

ELF File…
00000000: 0423c7801000c801 mov b32 $r0 b32 s[0x10]
00000008: 0000000310010009 mov b16 $r1l 0x1
00000010: a0000781d00e0005 st u8 g14[$r0] $r1
exit

To build NVDIS, use “make nvdis”, or use Microsoft Visual Studio 2010 with nv50dis.sln.

Caveats: This works on Windows with CUDA 3.2 PTXAS ELF generated files, using Cygwin GCC or MSVC 2010 to build NVDIS. I don’t know about other platforms and CUDA versions. The source is provided. This program looks at the ELF file version to switch between the different disassemblers in nv50-back.c or nvc0-back.c. And, it is a real hack of “readelf.c”. I replaced the file IO in the readelf code with an “unsigned char file_buffer”, which contains the entire contents of the elf.o file. I couldn’t find a great readelf.c to start with–many, many different versions out there, but none that seem to produce the nice output you get in the Cygwin’s readelf utility. The NV50DIS files used in this version were grabbed from the NV50DIS GIT repository (http://0x04.net/cgit/index.cgi/nv50dis/) a few days ago, and I think are already old.

Ken D.
nv50dis.zip (992 KB)

2

Hi, I wrote a little extension to Marcin KoÅ›cielnicki’s great tool NV50DIS. This tool, NVDIS, reads the ELF files produced by PTXAS. Here’s an example showing how to use it:

cat basic.ptx

.version 2.1
.target sm_20
.entry InstBasic (
.param .u32 __results
)
{
.reg .u32 %r<2>;

ld.param.u32 %r0,[__results];
mov.u32 %r1, 1;
st.global.b8 [%r0], %r1;
exit;
}

ptxas -arch sm_20 basic.ptx

nvdis elf.o

ELF File…
00000000: 2800440400005de4 mov b32 $r1 c1[0x100]
00000008: 2800400080001de4 mov b32 $r0 c0[0x20]
00000010: 1800000004009de2 mov b32 $r2 0x1
00000018: 9000000000009c05 st u8 wb g[$r0+0] $r2
00000020: 8000000000001de7 exit

cat basic10.ptx

.version 1.4
.target sm_10
.entry InstBasic (
.param .u32 __results
)
{
.reg .u32 %r<2>;

ld.param.u32 %r0,[__results];
mov.u32 %r1, 1;
st.global.b8 [%r0], %r1;
exit;
}

ptxas basic10.ptx

nvdis elf.o

ELF File…
00000000: 0423c7801000c801 mov b32 $r0 b32 s[0x10]
00000008: 0000000310010009 mov b16 $r1l 0x1
00000010: a0000781d00e0005 st u8 g14[$r0] $r1
exit

To build NVDIS, use “make nvdis”, or use Microsoft Visual Studio 2010 with nv50dis.sln.

Caveats: This works on Windows with CUDA 3.2 PTXAS ELF generated files, using Cygwin GCC or MSVC 2010 to build NVDIS. I don’t know about other platforms and CUDA versions. The source is provided. This program looks at the ELF file version to switch between the different disassemblers in nv50-back.c or nvc0-back.c. And, it is a real hack of “readelf.c”. I replaced the file IO in the readelf code with an “unsigned char file_buffer”, which contains the entire contents of the elf.o file. I couldn’t find a great readelf.c to start with–many, many different versions out there, but none that seem to produce the nice output you get in the Cygwin’s readelf utility. The NV50DIS files used in this version were grabbed from the NV50DIS GIT repository (http://0x04.net/cgit/index.cgi/nv50dis/) a few days ago, and I think are already old.

Ken D.

3

Hi Ken,

2 things:

davide

4

Hi Ken,

2 things:

davide

5

Hi Davide,

NVDIS is a disassembler for machine code for the GPU. In a large overview here, the CUDA C compiler produces PTX, which is then converted to machine code by PTXAS. This machine code (and/or the PTX source) gets stored in your CUDA program executable. NVDIS reads the file produced by PTXAS, not the executable. But, it can be very easily modified to do that. If you want, you can run PTXAS on one of the PTX files produced by the CUDA C compiler (use -keep), then disassemble it with this tool. Essentially, disassembly is useful in understanding how PTXAS works. The tool works with the CUDA v3.x Toolkit, and not the older toolkits because the input must be an ELF file.

The extension over NV50DIS/NVC0DIS is two fold: an ELF file reader front end, and an automatic way to select the disassembler. NV50DIS/NVC0DIS disassemble machine code fine, but what from what input file? Since I use PTXAS all the time, I wanted the disassemblers to read the elf.o file directly. Also, depending on how I call PTXAS (with -arch), I didn’t want to manually type “NV50DIS foo” or “NVC0DIS bar” for different targets. The NVDIS tool combines the two disassemblers, and selects the disassembler base on the target automatically. CUOBJDUMP will probably do this at some point, but doesn’t at the moment.

The NVDIS tool is a stand alone tool, not part of NV50DIS. It uses some of the C source code of NV50DIS, with some minor changes. Since it’s not part of that project, NVDIS can get out of sync.

You’re right on the link. The server seems down. Unfortunately, I have no control over that.

Ken

6

Hi Davide,

NVDIS is a disassembler for machine code for the GPU. In a large overview here, the CUDA C compiler produces PTX, which is then converted to machine code by PTXAS. This machine code (and/or the PTX source) gets stored in your CUDA program executable. NVDIS reads the file produced by PTXAS, not the executable. But, it can be very easily modified to do that. If you want, you can run PTXAS on one of the PTX files produced by the CUDA C compiler (use -keep), then disassemble it with this tool. Essentially, disassembly is useful in understanding how PTXAS works. The tool works with the CUDA v3.x Toolkit, and not the older toolkits because the input must be an ELF file.

The extension over NV50DIS/NVC0DIS is two fold: an ELF file reader front end, and an automatic way to select the disassembler. NV50DIS/NVC0DIS disassemble machine code fine, but what from what input file? Since I use PTXAS all the time, I wanted the disassemblers to read the elf.o file directly. Also, depending on how I call PTXAS (with -arch), I didn’t want to manually type “NV50DIS foo” or “NVC0DIS bar” for different targets. The NVDIS tool combines the two disassemblers, and selects the disassembler base on the target automatically. CUOBJDUMP will probably do this at some point, but doesn’t at the moment.

The NVDIS tool is a stand alone tool, not part of NV50DIS. It uses some of the C source code of NV50DIS, with some minor changes. Since it’s not part of that project, NVDIS can get out of sync.

You’re right on the link. The server seems down. Unfortunately, I have no control over that.

Ken

7

I can’t get this to build under either Visual Studio 2010 or Cygwin. Could you please provide either a binary or more detailed instructions?

8

I can’t get this to build under either Visual Studio 2010 or Cygwin. Could you please provide either a binary or more detailed instructions?

9

Attached is a zip file with both the VS2010 build (nv50dis/Debug/nvdis.exe) and the Cygwin build (nv50dis/nvdis.exe). I hope that works. If not, please send me a message. --Ken
nv50dis.zip (1.44 MB)

10

Attached is a zip file with both the VS2010 build (nv50dis/Debug/nvdis.exe) and the Cygwin build (nv50dis/nvdis.exe). I hope that works. If not, please send me a message. --Ken

11

Thanks, that helps. I’ve also worked out why I couldn’t build it - there’s some sort of problem with the release build.

12

Thanks, that helps. I’ve also worked out why I couldn’t build it - there’s some sort of problem with the release build.