sass disasm on perl

as an illustration of the use of the modules presented in my previous post I made yet another sass disasm - fully written on Perl. It is almost exact copy of my nvd - implemented just in 460 LoC, the only unsupported feature is registers tracking - bcs I still don't make perl binding for it. What it can do better than original nvdisasm:

• shows LUT operations
• shows instructions properties/predicates 
• shows relocs for each code section
• shows const bank params
  

and the most important thing - bcs it's based on Ced - you can patch any instruction from your script. Or customize output/save it somewhere like DB via Perl DBI/add your own passes to reveal some dirty nvidia secrets

like

Barriers

perl modules for CUBINs patching

After playing a bit with my ced I came to the conclusion that implemented DSL for editing is not enough - like it would be good to have subroutines to patch repeated/similar instructions, check that patched instruction is what I want, patch attributes/relocs etc

In other words, I need full-fledged PL. Although I've read books series "modern compiler implementation" from Andrew Appel and "crafting interpreters" I think making my own PL is overkill, so I made several XS modules to edit/patch CUBIN files for Perl. Why Perl?

• I am able to write on it almost all I want
• when I can't - I can always to develop my own module(s)
  
• yet I don't feel sick like from pseudo languages like python
• and it damn good and fast when you try to sketch out prototypes for things you have no idea how to make
  

 

ELF::FatBinary

source

for extracting/replacing CUBIN files from FatBinaries

see details here

Cubin::Ced 

source

In essence this is wrapper around Ced - it allows you to disasm/patch SASS instructions

Currently it don't support registers tracking

See doc in POD format 

Cubin::Attrs

source

Module to extract/patch attributes of CUBIN files + also relocs

doc in POD format

Sample

addresses of cuda kernel functions

 Quote from official document:

It is not allowed to take the address of a __device__ function in host code

I haven't been surprised for a long time that entire CUDA is made up of ridiculous restrictions. What if I told you that paranoid nvidia lies as usually and actually you can get addresses of kernel functions in your host code?

But first lets check what workarounds we can employ to have functions pointers. I don't know for what pedagogical purpose this code intentionally was written so poorly and does not free the allocated memory - and now millions of brainless artificial idiots will copy-paste it forever, so I made patched version. You can realize that attempt to read from early gathered with cudaMemcpyFromSymbol addresses will results error 1 (invalid argument)

Ok. but we could just return address of function directly from another kernel function, right? So I made quick & dirty hack

I brute-forced all combinations of cf1(__device__/__constant__) & variants of cudaMemcpyFromSymbol/cudaMemcpy - and with no luck

So it's time to run

cuda-gdb

practical ced usage: extracting sm machine ID

It's funny how paranoid nvidia trying to hide as much info from their customers as it can. One sample is so called "special registers" - even with PTX you can extract only limited set of them

So I played a bit with some undocumented SRs - namely with SR_MACHINE_ID_XX & SR_REGALLOC. I made legal loading of special registers and then patched those SASS instructions with my Ced

Lets see how those code looks in

PTX

      mov.u32         %r2, %tid.x;
      st.global.u32   [%rd2], %r2;
      mov.u32         %r3, %tid.y;
      st.global.u32   [%rd2+4], %r3;
      mov.u32         %r4, %tid.z;
      st.global.u32   [%rd2+8], %r4;
      // inline asm
      mov.u32 %r1, %smid;
      // inline asm
      st.global.u32   [%rd2+12], %r1;
      mov.u32         %r5, 21;
      st.global.u32   [%rd2+16], %r5;

As you can see - nothing special, just load in r2-r5 some values and store them in r12 holding address of function argument (obtained with cvta.to.global.u64)

tool to extract/replace files within CUDA FatBinaries

Patching of cubin files is good, but loading and running them requires lots of code and using of Driver API. It would be much more convenient to patch SASS directly in binaries produced by nvcc

Unfortunately evil nvidia as usually shows it's paranoia:

• cuobjdump can list & extract content but not replace. Also it is extremely buggy on old libraries like libcublas.so v7
• official fatbinary is too complex and rebuilds whole file from scratch
• format of fatbinary is undocumented
  

Well, the last problem already was solved. So I made utility to work with fatbinaries. You can

• list files with -v option
• extract file at some index: -i idx -o output.filename
• replace file at some index: -i idx -r replace.filename

Perl binding

Being lazy I prefer to use perl scripts to automate as much as possible, so I also made perl XS module ELF::FatBinary. Having also module ELF::Reader this allows more fine filtering of ELF files - like if file contains section/symbol with some specific name etc. See simple example how it might look like

Limitation

The tool can replace files inside fatbinary only in-place, so

1. compressed fatbinaries not supported
2. size of files must be the same
   

Some results

ced: sed-like cubin editor

Unfortunately, the only sass assembler I know of has several drawbacks:

• it's inactive last couple of years. I dropped email to his author and he didn't replied. Hope he is well
• it don't support modern sm architectures sm1xx
• it's matmul solver sometimes produces wrong instructions
• and it don't support many EIATTRS 

The last problem is not related with CuAssembler itself - it is more general: seems that nvdisasm produces output which cannot be used to assembly cubin files

Also we still don't know format of some sections like SHT_CUDA_RELOCINFO. All this makes task of rebuilding cubin files very hard

However do we really need to rebuild cubin files? In my experience 99.9% of desired patches are just set/remove some instructions attributes like register reusing/caching policy/wait groups for USCHED_INFO etc - just boring tuning to squeeze out the last couple of percent of productivity

So the flow of thought was something like

• it would be good to make plugin for hex-editor to disasm sass instruction at some known offset and show GUI where I could patch some fields
• I am talentless at creating GUI - so perhaps it would be better to dump instructions fields in text form and then just edit it
• hey - if you can parse this text representation and patch it back to sass - you don't need hex-editor at all - you could just use sed-like tool to patch instructions via script

and so being lazy and impatient I wrote such tool - it's called ced. Name similarity to sed is not coincidence - it allows you run text script to patch or replace some sass instructions inside cubin files

sass instructions: LUT operations

I was asked yesterday why I didn't transformed sample from my previous record


iadd r8, r2, r8 ; r8 = r2 + r8
iadd r8, r8, r8 ; r8 = r8 + r8
iadd r8, r8, ur4 ; r8 = r8 + ur4

to more simple

imad r8, r8, 2, ur4 ; r8 = r8 * 2 + ur4

While this is technically correct the problem here - ISA is non-orthogonal. You can use my ina to check available forms of IMAD for universal registers - and suddenly we will discover that it has only 2 forms

1. @Pg IMAD E:wide E:fmt E:Rd E:Pu E:Ra E:reuse_src_a E:Rb E:reuse_src_b -E:URc
2. @Pg IMAD E:wide E:fmt E:Rd E:Pu E:Ra E:reuse_src_a E:URb -E:Rc E:reuse_src_c
   

And no forms with imm value for Ra/Rb. So you can generate only something like:

imad r8, r8, rXX, ur4

And for UIMAD with imm values we have forms with universal registers only:

1. @UPg UIMAD E:wide E:fmt E:X E:URd E:UPu E:URa ,Sb ~E:URc !E:UPp
2. @UPg UIMAD E:wide E:fmt E:URd E:UPu E:URa ,Sb -E:URc
3. etc
   

But all this is just kids games compared to LUT operations. In short - you can have 255 combinations of logical operations over 3 operands driven by index. nvdisasm shows them like:

LOP3.LUT R0, R3, R0, RZ, 0x30, !PT 

Very informative, yeah. So I employed sympy to generate table of simplified expressions - however I am too old and lazy to write python scripts. So pretty obvious solution:

• make perl script to enumerate all possible combinations and generate python script
• which in turn generates string table
• and then sed add quotes and commas
  

And now my disasm shows much clearer output:
LOP3.LUT PT,R0,R3,R0,RZ, 0x30,!PT &req={5}; LUT 30: a & ~b
So here a = R3, b = R0 and result R0 = R3 & ~R0