ptx asm in llvm ir

I asked 7 month ago in r/llvm question how I can insert PTX asm right in LLVM IR and got exactly zero answers. So, as usual, I had to figure it out on my own (depressing little song "No Help is Coming" is playing in the background)

How inline PTX looks like in text form:
%7 = call i32 asm sideeffect "madc.hi.cc.u32 $0,$1,$2,$3;", "=r,r,r,r"(i32 %.sroa.018.0.extract.trunc, i32 %.sroa.282.0.extract.trunc, i32 0) #5, !srcloc !9 

So basically it is just call result-type asm with some arguments in parentheses (note that type of result $0 is i32 and it described after keyword call). If you need result of PTX instruction - just assign it to some variable. Official documentation says that #5 is attributes list - somewhere below it defined as
attributes #5 = { nounwind }
and !9 is metadata - is this case for debug info srcloc:
!9 = !{i32 46731}


Well, that was easy part of story - and now Something Completely Different (tm). LLVM IR is strictly typed (I would say - redundantly strictly), so types first time defined for each argument - like i32 for $1, $2 and $3. Second time - in string called operand constraint codes - in my case it is "=r,r,r,r". And official documentation blatantly lies about them. Let's check some source code - method getRegForInlineAsmConstraint in NVPTXISelLowering.cpp. As you can see it accepts following codes:

• b - 1bit, predicates
• c & h - 16bit, like (.b16 / .u16 / .s16)
• r & f - 32bit, like (.b32 / .u32 / .s32) and .f32 for f
• l, N, d - 64bit, (.b64 / .u64 / .s64) & .f64 for d
• q - 128bit since sm70+
• 0 - meaning is still unknown
  

Symbol '=' is so called Constraint Modifier:

• = Write-only output operand (overwrites previous contents)
• + Read-write operand (input and output tied to the same register)
• & Early-clobber operand (modified before inputs are consumed)
• ~ Clobber list marker (tells LLVM a register or memory/flags are modified implicitly
  

Yet another unpleasant discovery - you can freely swap order of operands - for example this variant is exactly the same as above one:
call i32 asm sideeffect "madc.hi.cc.u32 $0,$2,$1,$3;", "=r,r,r,r"(i32 %.sroa.282.0.extract.trunc, i32 %.sroa.018.0.extract.trunc, i32 0)

This makes the task of parsing & comparison of PTX instructions non-trivial - especially in complex cases like
%1 = call { i32, i32, i32, i32 } asm sideeffect "tex.grad.1d.v4.u32.f32 {$0, $1, $2, $3}, [$4, {$5}], {$6}, {$7};", "=r,=r,=r,=r,l,f,f,f"(i64 %tmp5, float %tmp6, float %tmp7, float %tmp8)

PTX from cicc

Once you understand how inline PTX is represented in LLVM IR, the next step is examining how nvidia’s own internal toolchain leverages it.

While doing some RE of nvidia's llvm-based back-end I dumped inline PTX instructions. Now when I have PTX parser the next logic step is try to parse PTX from cicc and for example try find some undocumented instruction/attributes (which nvidia uses for unfair competitive advantage). So I added to my parser option -r to dump instructions with unrecognized attributes, and also wrote little perl script to collect them. Then run whole pipe like

../ptx.parse/tp -r < ptx.txt | perl ../ptx.parse/ra.pl

And try to guess what happened? Yes - nvidia uses ~5-7% of instructions with undocumented attributes

This does not mean that cicc will produce every time PTX with 5-7% of instructions which clang can't. This means that it just generates more effective code in some cases. For example there is instruction mma.sync.aligned.m16n8k4.row.col.f32.tf32.tf32.f32

However mma has following 43 forms
 line 84:F16F16F16F16
 line 85:F32F16F16F32
 line 86:F32F32
 line 87:I32I8I8I32
 line 88:I32I32
 line 89:F16F16F16F16
 line 90:F32F16F16F16
 line 91:F32F16F16F32
 line 92:F16F16F16F32
 line 93:I32I32
 line 94:I32I8I8I32
 line 95:I32I32
 line 96:F64F64F64F64
 line 97:F32F32
 line 98:F32Q8Q8F32
 line 99:F16Q8Q8F16
 line 100:F32Q8Q8F32
 line 101:F16R4Q8F16
 line 102:F32R4Q8F32
 line 103:F16Q8R4F16
 line 104:F32Q8R4F32
 line 105:F16R4R4F16
 line 106:F32R4R4F32
 line 107:F16R4Q8F16
 line 108:F32R4Q8F32
 line 109:F16Q8R4F16
 line 110:F32Q8R4F32
 line 111:F16R4R4F16
 line 112:F32R4R4F32
 line 113:F16Q8Q8F16
 line 114:F32R4R4F32Q8
 line 115:F32R4R4F32Q8
 line 116:F32Q8Q8F32Q8
 line 117:F32R4Q8F32Q8
 line 118:F32Q8R4F32Q8
 line 119:F32Q8Q8F32Q8
 line 120:F32R4Q8F32Q8
 line 121:F32Q8R4F32Q8
 line 122:F16F16F16F16
 line 123:F32F16F16F32
 line 124:F32F32
 line 125:I32I8I8I32
 line 126:I32I32

and if we grep only first and last F32 we will get
 line 85:F32F16F16F32
 line 86:F32F32
 line 91:F32F16F16F32
 line 92:F16F16F16F32
 line 97:F32F32
 line 98:F32Q8Q8F32
 line 100:F32Q8Q8F32
 line 102:F32R4Q8F32
 line 104:F32Q8R4F32
 line 106:F32R4R4F32
 line 108:F32R4Q8F32
 line 110:F32Q8R4F32
 line 112:F32R4R4F32
 line 123:F32F16F16F32
 line 124:F32F32

as you can see there are no forms with 4 F32 or F32T32T32F32. How this can work? Well, mma is pseudo-instruction and they processed specially by PTXAS

What conclusion can be drawn from all that has been said? If your compiler/toolchain generate LLVM IR - try to compile it with cicc to get better quality of output code