parser of PTX instructions

A couple of facts to start things off

From official "Inline PTX Assembly in CUDA":

The compiler front end does not parse the asm() statement template string and does not know what it means or even whether it is valid PTX input

And second, less well-known one: order of instruction's attributes (except types of operand) is not important

The combination of these facts leads to stark conclusion - CUDA compiler front-ends totally ignore PTX inline asm and only PTXAS known how to parse them. For example cuKLEE does this wrong

So I made simple (and hopefully fast) parser of PTX instructions

Note: this is not full featured replacement of PTX parser. it is designed specifically to extract instruction attributes and determine the correct instruction form based on argument types and counts

For example for

cvt.bf16x2.e5m2x2.rn.relu.scaled::n2::ue8m0.satfinite d, a, scale-factor;

output will be something like

tail: d, a, scale-factor;
3 tail operands
--> cvt
 line 71: 01x E32Q16
--- types 2:
 bf16x2
 e5m2x2
--- attrs 4:
 1:5 satfinite
 2:7 scaled::n2::ue8m0
 1:0 relu
 3:3 rn 

identification of const bank0 params

Official documentation doesn't disclose ConstBank0 (c[0x0]) memory layout used at the SASS level for kernel arguments and special registers (like %gridid, %nctaid)

So I've spent last week trying to solve this deceptively simple problem. Names of params are documented in official doc - seems that this time ptxas can't add something new to this list

Unfortunately I was unable to find inside ptxas some nice looking tables for pile of SM.  What other approaches can we use? As usually the first thought is to do some brute-force.

Brute-forcing

Lets write in plain PTX dummy function trash with u32 return value - something like

.visible .func (.param .u32 func_retval0) trash
{
  .reg .u32       %r<3>;
  mov.u32    %r0, %gridid;
  mov.u32    %r1, %nctaid.x;
  add.u32    %r0, %r0, %r1;
  st.param.u32 [func_retval0+0], %r0;
  ret;
}
The final st.param is very important bcs otherwise compiler will just eliminate whole code. Instead of gridid & nctaid.x we can substitute pair of special registers, compile with ptxas to specific SM and then parse output of nvdisasm/nvd/whatever can disasm SASS

Surprisingly, this stupid method worked very well, however there are holes in params. So it's time to check

CUDA runtime

I extracted them in December and now we can parse output of nvdisasm to find not identified yet offsets. The funny part is that official nvdisasm failed on several files, like sm54.elf

nvdisasm error   : Could not establish the target of this branch operation

or on sm23.elf

nvdisasm error   : Wrong Anti dependency order in function 'vfprintf_internal'
nvdisasm         .         @P1 LD.E.CG.64 R14, [R4], P0
nvdisasm         .          -- Anti(PRED,0),0*,0 -->
nvdisasm         .         @!P1 LEA.HI.X P0, R7, R12, RZ, R13

So I was forced to use my own nvd

It turned out that the parameter space is divided into two parts - there are block of parameters at offset 0x1860 (holding for example starting PC of kernel) used by kernel launch logic and CnpXXX functions

So now we know lots of offsets and their sizes. However to identify semantics of many found offsets we need debugger

cuda-gdb rushes to the rescue

I made fake PTX for each SM, patched it with my ced and inspected in debugger values with command $_cuda_const_bank(0, offset). Actually this was the most boring part of work and I still didn't recognized some fields. Also I don't have expensive monsters like sm100+ so I extracted only params from Maxwell till Hopper

 

Results

Source

I also add this code to my XS perl module and nvd, so output looks like

/*58*/  XMAD R02,R17,c[0][0x8],R02 ?trans1;
 ; cb0 param %ntid_x


Names starting with '%' were extracted with just disasm of dummy trash function

 

Happy hacking!

RE of PTX grammar from ptxas, part 4

Parts 1, 2 & 3

First of all, it should be noted that the mask of instruction attributes has size 20 bytes, so I updated dump for them.

structure for this attributes descriptor has size 0xd8 bytes and some fields:

• mask at offset 0
• name of instruction at 0xC8
• index at 0xD0
  

Instructions selecting first by name and then right form by operand types. This means that while the order of the attributes does not matter, the relative order of the operand types is important - leftmost is type of operand 0, next is type of operand 1 and so on

 

Names of numerical pseudo-instructions

in part 3 I pointed out that there are 473 names consisting only of numbers, like "1030557441". Grigory Evko suggested that this is adler32 hash from builtin function names, so I found huge function for instruction 0xc6 (_gen_proto) returning 1078 prototypes like

.weak .func (.reg .f32 %fv1) __cuda_sm20_div_rz_f32 (.reg .f32 %fa1, .reg .f32 %fa2)


and then intersected them by hash - so now we know all real names

EBNF grammar

You can see it here

To build run iptx.pl -e

The last two columns are operand suffix & encoding 

 

How complete it is?

That's good question. If we accept that attributes descriptors contain full list of attributes for each instruction then 20 bytes masks has 121 non-zero bits:

FD FF FF FF F1 FF FF 9F F9 FF E7 CF FF F3 DF FF FF 00 00 00

I was able to identify 114 of them - this is 94%

Also currently I extracted 119 tables with attributes names and only 11 are still not connected (check them with iptx.pl -t)

On other hand in function for attributes processing there are 3 switch tables with 139, 140 & 173 cases (last one has ~90% of entries with error "Unexpected instruction types specified")

RE of PTX grammar from ptxas, part 3

Parts 1 & 2

Pseudo instructions

Surprise-surprise - some PTX instructions not mapped directly to underlying SASS 1:1. Instead they generate lots of another PTX code. I already extracted their decrypted bodies, so it's time to describe how they connected to specific PTX pseudo instructions

 

There is function somewhere deep inside ptxas which register lots of handlers for dumping real PTX for pseudo instructions. Code for registration of single item looks like

  mov     rdi, [rbx+250h] ; dictionary of pseudo-instructions
  lea     rdx, emit_multimem_ld_reduce ; handler
  lea     rsi, aMultimemLdRedu         ; "multimem.ld_reduce" - pseudo instruction name
  call    reg_sm_cb



There are 587 such handlers - although 473 have strange names like "1030557441". I don't know what they mean - highly likely that this is product of another encryption somewhere inside parser - at least each such string has exactly 1 reference

Lets look inside some handler

  call    get_pool
  mov     rdi, [rax+18h]
  mov     esi, 0C350h ; 50000₁bytes - they don't skimp on matches
  call    alloc_buf
  test    rax, rax
  mov     r12, rax ; r12 holds address of string buffer
  jz      loc_5626FC1E9D78 ; die in alloc_failed
loc_5626FC1E9733:            ; CODE XREF: emit_multimem_ld_reduce+67D↓j
  lea     rdx, [r13+1A5E95h] ; whut ?
  lea     rsi, aS_11         ; "%s"
  mov     rdi, r12           ; s
  xor     eax, eax
  call    _sprintf ; note that even not snprintf - security above all!
  lea     rdx, [r13+1A5E98h] ; whut again ?
  movsxd  rdi, eax ; store in rdi length of written string
  lea     rsi, aS_11         ; "%s"
  mov     rbx, rdi
  xor     eax, eax
  add     rdi, r12           ; s
  call    _sprintf 


 

Debugger showed that R13 holds address of those decrypted string pool in memory. 

Just assess the level of paranoia - there is huge encrypted blob with strings 1.8Mb. Then they wrote 587 functions where each string from those blob can be used only by offset - 21042 unique offsets! Nvidia definitely didn't want us to see its dirty secrets.

 

So I wrote some code to extract all emitters, then all string offsets - see result. Now it would be good to link offsets from each emitter with real string, right?

 

Nothing is simpler - yet another Perl XS module to load memory mapped file + small perl script - and finally we can see this

Lexer brute-force

RE of PTX grammar from ptxas, part 2

Part 1

PTX instructions that cicc cannot generate

While reverse-engineering Nvidia's compilation pipeline, I extracted the set of PTX instructions that cicc (the CUDA C++ frontend) is capable of emitting. The next logical step is to intersect them with full set of instructions accepted by ptxas - so we could get instructions which cicc just unable to produce. To do this I add to iptx.pl new option -U and got file ptx_not_in_cicc.txt with 114 unique names

PTX in total has only 268 unique names - so 114 is 42.5%. Notable missing instructions include:

• cctl for cache control
• lop3 - yeah, I saw them many times in SASS, so it generated by ptxas during optimization passes
• r2p
• 11 variants of tcgen05.*
• mad24/mul24
  
• all video instructions like vadd/vmad/vset etc
  

 

This gap is large enough to be surprising and leads me to conclusion that official LLVM MLIR dialects for cuda are totally incomplete

MLIR was initially a very dubious idea IMHO - what if we have some unscrupulous HW vendor who prefers to hide many details of it's hardware? And even worse - when multiple MLIR dialects are involved (like gpu, nvgpu, nvvm, linalg etc), at least one of them has to maintain accurate mappings between all of them. This leads to exponential explosion of complexity - you can expect items from each of used dialects while doing optimization, and also creates surface area for bugs.

some instructions are totally undocumented

RE of PTX grammar from ptxas

Disclaimer

Highly likely that author is an illiterate, inattentive, and incompetent lazy person with a poor imagination - therefore his hypotheses may be questionable, ideas delusional and his analysis simply incorrect. Also maybe I still haven't mastered ida pro in 28 years so extracted data can be incomplete/have missed parts. As always all code on perl and therefore offends the aesthetic feelings of believers

 

Prior works

• Official PTX ISA. We all know than nvidia is evil and paranoid, so this document also incomplete and maliciously conceals information. Proofs are somewhere below in this text
• ANTLR ptx grammar - very outdated, based on cuda-waste parser from 2010
• infamous zluda. It's enough to look at their AST to understand that they support at best a third of the instructions
• nvopen-tools by Grigory Evko. AI generated slop, but at least we can borrow from chapter 7 format of instructions and decoding scheme for arguments

So as you can see there is no machine readable grammar for modern PTX, Why this is important at all? Well, according to "Official guide to inline PTX"

The compiler front end does not parse the asm() statement template string and does not know what it means or even whether it is valid PTX input

Therefore you can successfully compile your buggy code to PTX and suddenly got mysterious errors during dynamic loading over JIT. Plus I always suspected that nvidia hides as much information from us as possible

 

So I started with some disassembly of ptxas version V10.1.243 from sdk 13.1 looking for PTX instruction names (encrypted btw)

 

Data extracting

Instruction attributes dynamically filled in two places

• in huge function at 0xC2341C - extracted data
• in array of functions located at 0x2971260 - data merged with previous chunk

Please don't ask me why there are 2 separate places. More importantly that code from both looks uniform

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)

ptx instructions emitting by nvidia compiler. part 2

Part 1 described v10

And today let's check cicc v12. The first thing that catches your eye is its size - almost 76Mb! And it also contains at least 5 different decryptors - Nvidia really wants to hide something from its grateful clients

Why it is so fat?

Bcs it contains at least 4 code generators: for arm32, aarch64, x86 & nvptx

+ at least 27 llvm bytecode blobs (signature 0x42 0x43 0xc0 0xde) - they contains mostly bodies of intrinsic functions like nvvm_mulq/nvvm_divq but on some llvm-dis just crashes:

#0  0x000055abcc64743d in llvm::Intrinsic::getIntrinsicInfoTableEntries (id=0, T=...) at /home/redp/disc/src/llvm-project/llvm/lib/IR/Function.cpp:1339
1339      unsigned TableVal = IIT_Table[id-1];
>>> where
#0  0x000055abcc64743d in llvm::Intrinsic::getIntrinsicInfoTableEntries (id=0, T=...) at /home/redp/disc/src/llvm-project/llvm/lib/IR/Function.cpp:1339
#1  0x000055abcc5fe41f in UpgradeIntrinsicFunction1 (F=0x55abce63cf18, NewFn=@0x7ffc60414e70: 0x0) at /home/redp/disc/src/llvm-project/llvm/include/llvm/IR/Function.h:204
#2  0x000055abcc60111a in llvm::UpgradeIntrinsicFunction (F=F@entry=0x55abce63cf18, NewFn=@0x7ffc60414e70: 0x0) at /home/redp/disc/src/llvm-project/llvm/lib/IR/AutoUpgrade.cpp:1226
#3  0x000055abcc584778 in (anonymous namespace)::BitcodeReader::globalCleanup (this=0x55abce608e30) at /home/redp/disc/src/llvm-project/llvm/lib/Bitcode/Reader/BitcodeReader.cpp:3696
#4  0x000055abcc5856cc in (anonymous namespace)::BitcodeReader::parseModule (this=<optimized out>, ResumeBit=<optimized out>, ShouldLazyLoadMetadata=<optimized out>, Callbacks=...) at /home/redp/disc/src/llvm-project/llvm/lib/Bitcode/Reader/BitcodeReader.cpp:4385
#5  0x000055abcc5959ca in (anonymous namespace)::BitcodeReader::parseBitcodeInto (Callbacks=..., IsImporting=false, ShouldLazyLoadMetadata=false, M=0x55abce5f3d80, this=0x55abce608e30) at /usr/include/c++/9/bits/std_function.h:564
#6  llvm::BitcodeModule::getModuleImpl (this=<optimized out>, Context=..., MaterializeAll=<optimized out>, ShouldLazyLoadMetadata=<optimized out>, IsImporting=<optimized out>, Callbacks=...) at /home/redp/disc/src/llvm-project/llvm/lib/Bitcode/Reader/BitcodeReader.cpp:7981
#7  0x000055abcc596070 in llvm::BitcodeModule::getLazyModule (this=0x7ffc60415d10, Context=..., ShouldLazyLoadMetadata=<optimized out>, IsImporting=<optimized out>, Callbacks=...) at /usr/include/c++/9/bits/std_function.h:263
#8  0x000055abcc550e49 in main (argc=<optimized out>, argv=<optimized out>) at /home/redp/disc/src/llvm-project/llvm/include/llvm/Support/CommandLine.h:1399

>>> p id
$1 = 0


At least they should check that index can become negative, no? Who would doubt that llvm is very reliable and secure

So if they process llvm ByteCode then they also must link half of llvm run-time to do it, but they also use 

ptx instructions emitting by nvidia compiler

I recently became curious what exactly ptx instructions can produce nvidia compiler - like if it uses something totally undocumented or vice versa - some official ptx instructions are never generated during compilation

The first thing is where those compiler located - no, it's not nvcc. Real compiler is cicc from packet cuda-nvvm. cicc from v10 has size 21Mb. The strings utility shows many interesting things, like

Portions Copyright (c) 1988-2016 Edison Design Group, Inc.
Portions Copyright (c) 2007-2016 University of Illinois at Urbana-Champaign.
Based on Edison Design Group C/C++ Front End

So they use front-end from Edison Design Group and llvm as back-end

Then I extracted several tables:

1. compiler errors
2. list of built-in functions with prototypes
3. list of llvm attributes - as you can see they are mostly correspond to nvvm LLVM dialect
4. and finally what I looked for - list with internal instruction names & their bodies to place into PTX file

As you can see mapping is very straightforward - for example for instruction BFE_S32rii (index 0x27) generating PTX bfe.s32

Results

I found only minor PTX instructions not presented in their official documentations: suq.xxx - perhaps should mean surface qword or something like this