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

Nvidia is probably using them for unfair competition. Couple of most fat is _mma.warpgroup & _mma. They have lots of unique attributes not presented in other instructions. To reduce noise during analysis I add -w option to skip them

order of attributes is not important

Big surprise. Officially mma.sync has order of attributes

• aligned
• shape like m8n8k4 and friends
  
• row
• col
• dtype 

something like

mma.sync.aligned.m8n8k4.row.col.f16.f16.f16.f16

but this construction is also perfectly valid:

mma.sync.col.aligned.row.m8n8k4.f16.f16.f16.f16


 

I think that at least one of my hypothesis was correct - parser just collects any valid attributes without caring about their relative order

extracted data is not complete

lets check for example instruction tcgen05.fence - officially it has couple of possible values for attribute. However if we check how it looks in extracted attributes:

grep tcgen05.fence ptx_ops2.txt
160 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 tcgen05.fence


we can observer it has zero mask for all attributes

 

Another sample - istypep - again has almost zero attributes mask:

grep istypep ptx_ops2.txt
181 20 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 istypep    PU    O


here 20 in first byte is just operand type prefix

 

current status 

As of today

• identified 60 attributes
• covers 134 unique instructions (50%)
• covers 734 forms out of 1090 total (67%)