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 

Instruction cvt has potential 40 forms

Parser selected single form for operand types "E32Q16", format of operands "01x" - first two is some variables and 3rd must be hexadecimal number (in this case for scale-factor)

Instruction has 4 attributes, left column indicates the bit index for corresponding attribute in mask

If we swap some attributes like

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

- we will get exactly the same result

How parser works 

The original parser in PTXAS employs (f)lex/yacc. I decided that this is overkill - when we know attributes masks for each instructions we could not only produce EBNF grammar but also generate all necessary lookup tables. So I added to my iptx.pl option -C to produce ops.inc with all pre built boring stuff

First thing parser does is split input string into (optional) predicate, body of instruction and tail containing operands. You can extract them with methods pred/body/tail of PTXParser class

Then it splits body by dots and put every sub-string into map - when happens some match with instruction or attribute name - they are removing from this tokens map

After that I try to find instruction name. The problem here is that instruction itself can contain dots - like cp.reduce.async.bulk.tensor, so I lookup first 5 tokens in descending order of the number of tokens. This ensures longest-prefix matching (e.g., matching cvt.pack before cvt).

Now when we know instruction name - we have list of forms. and next logical step is to select appropriate. There are 2 filters:

1. By number of operand. First naive implementation was just split operands by commas, however official PTX documentation has wonderful example:

ld.global.L2::evict_last.L1::evict_last.v4.u64 { %reg0, %reg1, %reg2, %reg3}, [addr]; 

Actually there is only 2 operands. So parsing is little bit more complex. For this reason, this filter is applied optionally - only with -t option

2. By operand types (method try_type). Their descriptors can have 3 forms:

• T [ 4 | 8 | 16 ] - 3 possible widths for type T
• T16 - single possible width
• just T - ignore width (see function cmp_letter) 

Unfortunately in some rare cases multiple forms could be selected. So at final step to extract attributes I build union set of their masks with OR (method fill_attrs). That's to avoid case when some attribute A1 belongs only to Form1 and another A2 to Form2 - if we build intersection of masks then neither will be parsed. And finally recognized attributes stored in std::multimap ParseRes::attrs with mask index

Multimap is used because some attributes can repeats - for example for matrix there are layout A & layout B, so corresponding layout attribute can have up to 2 values

Residual tokens: any unparsed attributes can be retrieved via method rem_attrs

That's all, folks. At each step I use pre built lookup-tables, so overall complexity of algorithm is O(N) where N is number of tokens