machineIR passes & mlir dialects from tileiras & cute dsl

Previous parts: 1, 2 & 3

Let's continue to dissect compilers from nvidia. This time, I decided to examine compilers that are guaranteed to contain their custom MLIR dialects - tileiras and CuTe DSL

tileiras

As usual, their official repository is completely useless - it does not contain code for conversions/lowering so you can't build your own version of this compiler. So I used exactly the same techniques and scripts for its analysis from previous part

Strings from constructors - contain nothing of value

Passes shows 5 custom missed in vanilla llvm:

1. nvptx-set-global-array-alignment (Find good alignment for statically sized global/shared arrays)
2. nvptx-set-local-array-alignment (Find good alignment for statically sized alloca arrays)
3. nvvm-peephole-optimizer (NVVM Peephole Optimizer)
4. lower-aggr-copies (Lower Aggregate Copies)
5. mergereturn (Unify function exit nodes)
   

As you can see this list differs from cicc and prompts some philosophical musings: just how many LLVM forks do they have, and are those dev-teams even aware of each other’s existence?

enumerate registered MLIR dialects

There is very good doc about MLIR capabilities from NPU compiler, however it does not contains answer to our question - how we can extract from LLVM binary all registered MLIR dialects

So I spent several days poking around megabytes of MLIR sources. It seems that perfect solution does not exist (due of heavy using of template black magic), but we can get practically good one - method MLIRContext::getOrLoadDialect:

• it can be identified by string "a dialect with namespace". Warning: there are another methods using this string like MLIRContext::getOrLoadDynamicDialect but fortunately they usually not linked with custom LLVM forks
• it's called from BytecodeDialect::load method and so we can catch all dialects allowed in input .bc files

So I wrote yet another dirty IDC script to extract MLIR dialects - total 30 unique

CuTe DSL

Official repository claims that you can build your own compiler with vanilla LLVM (I never verified if that is true). What they didn't said - in their python packages compiler based on their closed-source LLVM fork, in my case I dissected _cutlass_ir.cpython-313-x86_64-linux-gnu.so

First remarkable thing is its size - 144Mb, for example size of cicc is 78Mb and size of tileiras 89Mb. Highly unlikely that python bindings have size of ~60Mb, so lets check  

strings from constructors

We could remark something very familiar - string like AddWaitForCCOrPOffDeck & co. Also section .rodata has huge blobs for "fermi macros", so we can conclude that it was linked with PTXAS

machineIR passes

+ 54 new passes, wow! Also it has 2767 nvvm intrinsics vs 1998 in plain LLVM

MLIR dialects 

As you can see it contains 32, including our beloved cute & cute_nvgpu

 

Conclusion

It appears that CuTe DSL compiler based on the same LLVM fork as cicc and so produces fastest code, also it supports 32 MLIR dialects and scripting from python - just combo

comparison of nvcc vs clang nvptx backend, part 3

Parts 1 & 2

Little disclaimer: I am not llvm expert and never created my own backend - in best case I can write simple optimization pass, lower yet another senseless MLIR dialect or extract some info from IR (and prefer gcc bcs its code more clean & neat and not scattered on thousands of sub-dirs as if written by an orthodox Java devotee), so take everything described below with healthy dose of skepticism

Task of searching for modifications/patches in general is well known, the problem in our case is that binary files is just huge

• cicc from SDK 13.1 has size 80mb
• libLLVM.so.21.1 128mb
  

besides that cicc does not contain DWARF debug info, so using of tools like bindiff seems like a questionable idea. Let's check if can invent some dirty tricks to reduce volume of code for comparison

CL options

Let's start with something easy - as you could know in llvm all CL options represented as static vars of template class cl::opt<Type>, like

static cl::opt<bool>
NVVMReflectEnabled("nvvm-reflect-enable", cl::init(true), cl::Hidden,
                   cl::desc("NVVM reflection, enabled by default"));


We can guess that code to create them must be located in array of constructors - in section .ctors or .init_array (actually second can have any name - only section type INIT_ARRAY matters). What really important here is that constructor of these object must load some constant strings from .rodata section

So I wrote dirty hack to disasm array of constructors and dump all used string to file. Results:

• strings from cicc - 3743
• strings from llvm - 7102

I also wrote simple script for comparison, so just run

perl cstr.pl clang-cstr.txt cicc-cstr.txt

What interesting we can note in its output?

• print-schedule - Print 'sched: [latency:throughput]' in .s output
• wmma-memory-space-opt
• dump-ir-before-memory-space-opt & dump-ir-after-memory-space-opt
• enable-bfi64 - Enable generation of 64-bit BFI instructions
• nvptx-add-scalar-move-for-vector-load, nvptx-lower-global-ctor-dtor & nvptx-use-max-local-array-alignment
  

This method retrieves all the strings used in constructors, and for this reason, it is not very reliable, so check every interesting string in disasm

Passes 

comparison of nvcc vs clang nvptx backend, part 2

Part 1

First of all, I have to say that digging through LLVM source code is simply disgusting - its code tree is monstrously huge and parts from nvptx scattered across the most unexpected places. So to extract needed tables I just stuck libLLVM.so.21.1 (btw size 135Mb) into IDA Pro and reused the same IDC scripts for cicc, just removed strings decryption. Typical open-source, yeah

Mlir

cicc does not have MLIR at all. This fact probably deserves a larger font:

cicc does not have MLIR

The proof is left as an exercise for the reader very simple - just grep strings like mlir/linalg/affine etc in cicc binary. Let's assume that we have some way to load LLVM bc into cicc - even in this case, we wouldn't be able to force it to generate more efficient PTX for code containing MLIR

Btw this is not first case when serious amount of real code generation was done without mlir - for example vectorization for RISC-V too. And somewhere on the side, you’re bound to hear that argument from clueless managers: "if even Oprah Microsoft uses a monorepo, you must be happy using one too" 

Builtins

• from cicc - 870
• from clang -  715

Notably, neither Nvidia nor LLVM documented the list of supported builtins. "doc is also quite good in my opinion", of course

Intrinsics

• from cicc - 3183
• from clang - 1998
  

I think this is one of the secrets behind why Nvidia's code is faster - they just capable of generating 1.5 times as many PTX instructions. What happens when NVPTX backend can't compress group of operations into single PTX? Sure it generates several - like ld/cvt/something or something/cvt/st. And ptxas isn't enough smart to detect such patterns and apply more optimal instructions

 

Conclusion

Category	cicc	Clang (NVPTX)
Builtins	870	715
Intrinsics	3183	1998

If you use MLIR - then you cursed forever must use NVPTX backend only and your code will always be 15-20% slower

comparison of code from nvcc vs clang

Lets continue our methodical dissection of PTX/SASS code generated by different compilers

Last week I implemented on CUDA count-sketch algorithm. It requires d = 2 * t + 1 calculating of 2 different family of hashes. As you can see d is always odd and don't fit on warpsize 32, so I used reverse approach - warp processing 32 items from stream and each thread calculates it's own hashes in loop. The code is terrible, consists of nothing but dirty hacks and optimized for specific structures, so that’s why I’m ashamed to show it in its entirety. Nsight shows that code for hashes calculating is where most of the time is spent (~80%), so this code will be our test data. As you can see it's just plain boring integer arith - no fashionable mma or even warp reduce operations

For tests I used clang-21 and nvcc/ptxas from CUDA SDK 13.1. Output data located in 2 sub-dirs:

clang-21, options for clang --target=nvptx64-nvidia-cuda --cuda-device-only -fomit-frame-pointer -O --offload-arch=sm_XX -S

nvcc, see details how to dump LLVM bc here. Btw cicc from sdk 13.1 is based on llvm21 too, like
strings ../../../13.1/cicc/cicc | grep ^LLVM
LLVM21.0.0


As final stage the same ptxas was used to produce object files with same set of options -c -arch sm_XX

Results

ptx generated by LLVM NVPTX

I decided to complete saga with ptx instructions which can be generated with popular compilers and check clang with NVPTX target

Gathering from LLVM source files with PTX instructions turned out to be a challenging task. The Usual Suspects NVPTXInstrInfo.td & NVPTXIntrinsics.td lay in directory llvm/lib/Target/NVPTX, but they don't contain operation definitions. For unknown reason file NVVMOps.td located in mlir/include/mlir/Dialect/LLVMIR/ - don't ask me why

Then I add new option -N to my long-suffering iptx.pl to collect PTX instructions. As you can see it's not true parser - it just extracting any strings and then try to find matches in set of PTX instructions. Dirty & ugly hack but since we lookup on small dictionary with only 270 instruction it works in near-linear time. Another problem is that strings with instruction names can be modified like append operands etc. Well, since cicc itself based on LLVM - if some instruction marked as missed in clang it's highly likely also will be marked as missed in cicc too, so the obtained sets are sufficient for the comparison task

Results

Check two files:

1. ptx-cicc.txt - set of PTX instructions which cicc can't produce. As you can see coverage 57.8%
2. ptx-nv.txt - set of PTX instructions which clang can't produce for NVPTX. Coverage 58.9%

And instructions missed in clang (but presented in cicc) are:

• _ldsm
• _mma
• discard
• dp2a hi/lo
• madc hi/lo
• p2r
• trap

9 / (270 * 0.58) = 5.7%, which correlates reasonably well with the number of PTX instructions from the cicc that use undocumented attributes (5-7%)

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

dumping llvm bitcode from cicc

method 1

requires building of hijacked .so with appropriate LLVM version. I am too lazy for this

 

method 2

cool, but does not work - cicc claims on bad arguments. I've tried many combinations with no luck

 

But hey - we are under linux and can make many hacks, for example check what arguments genuine nvcc passing to cicc. For this I ran nvcc -dc -keep under strace:

strace -o c.strace -s 512 -f --trace=/^exec nvcc ...


Arguments:

• -s NUM - maximum string size, bcs arguments can be very long - I set this parameter to 512
• -f - trace child processes
• and finally --trace - since I don't know which exactly syscall used to launch processes I used regex syntax for all calls starting with exec

Lets check output file c.strace and see launches of

• gcc/cc1plus
• cicc
• ptxas
• fatbinary
• bin2c
• cudafe++
• etc
  

After some trials right combination of arguments for cicc is
NVVMCCWIZ=553282 cicc --nv_arch compute_XX --device-c -keep 1.cpp1.ii
ls -l *.bc
-rw-rw-r-- 1 redp redp 8072 mar 31 13:25 1.lgenfe.bc
-rw-rw-r-- 1 redp redp 9988 mar 31 13:25 1.lnk.bc
-rw-rw-r-- 1 redp redp 6500 mar 31 13:25 1.opt.bc

lgenfe.bc - bitcode from front-end

opt.bc - bitcode after all optimization passes

to disassembly we can now just use llvm-dis-21:

  %1 = tail call i32 asm sideeffect "activemask.b32 $0;", "=r"() #3, !dbg !11
  %2 = tail call { i32, i1 } @llvm.nvvm.shfl.sync.i32(i32 %1, i32 3, i32 %val, i32 16, i32 31) #3, !dbg !17
  %3 = extractvalue { i32, i1 } %2, 0, !dbg !17


read a couple of books about compilers

LLVM Compiler for RISC-V Architecture

Describes details of risc-v vectorization support in llvm. It should be noted that the implementation of vector operations in risc-v was done later than in Intel and sve in arm64 - they took into account many flaws (like made explicit masks for vector operations) and were implemented in a much more convenient way from the programmer's point of view

On other hand any HW vendor can add it's own ISA subset and support of this custom processors in compiler can become very segmented and pure nightmare

 

Also I want to note that support of risc-v vectors in LLVM carefully avoids MLIR (IMHO second most overrated thing after LLM) - to do this they even had to patch their holy cow tablegen

 

Drawbacks:

• there is no introduction about LLVM IR/risc-v specific IR, so long IR listings are very hard to follow
• author don't give link to source code implementing some algo. Fortunately elixir indexed whole LLVM source tree
  

4/5

Dive into Deep Learning Compiler

As far as I know, this is the only book describing AI/ML compilers so far. Also TVM looks very promising - unlike monsters like XLA/iree it is compact and observable for mere mortals

Drawbacks:

• book is not completed - last two chapter about NN & deployment are just "place holder"
• it's unclear why for matrix multiplication on CUDA they didn't get cublas as base case
• and openblas for cpu version

Despite this, considering that the book is freely downloadable, my rating is 4 out of 5

CFG with LLVM

On holydays I read book "LLVM Cookbook" (not very good - lots of meaningless copy-pasted code blocks are annoying) and played a bit with fresh llvm-3.7.1 (was released 5 january)

So I decided to check whether it is possible to implement MS CFG in llvm. I have two news - good and bad, as usually

Good: yes, you can easy add instrumentation in llvm - just add some plugin for IR derived from FunctionPass and add call to your guard_check_icall before each VTBL call (or even on any ptr call). I think it will take one day for any CS-student

Bad: you need integration with MS linker and it seems that support of CFG in COFF files is totally undocumented. LLVM itself cannot make load_config and even more - their definition of coff_load_configuration in include\llvm\Object\COFF.h has no fields for CFG (like GuardCFCheckFunctionPointer and GuardCFFunctionTable)

llvm 3.3 - wtf ?

was built under xp 64bit with visual studio 2010
And I got:
30>  Testing Time: 106.52s
30>  ********************
30>  Failing Tests (32):
30>      LLVM :: CodeGen/ARM/lsr-scale-addr-mode.ll
30>      LLVM :: CodeGen/X86/rodata-relocs.ll
30>      LLVM :: Linker/2003-08-24-InheritPtrSize.ll
30>      LLVM :: Linker/2008-03-05-AliasReference2.ll
30>      LLVM :: Linker/2008-07-06-AliasFnDecl2.ll
30>      LLVM :: Linker/2008-07-06-AliasWeakDest2.ll
30>      LLVM :: Linker/2009-09-03-mdnode2.ll
30>      LLVM :: Linker/2011-08-04-DebugLoc2.ll
30>      LLVM :: Linker/2011-08-04-Metadata2.ll
30>      LLVM :: Linker/2011-08-18-unique-class-type2.ll
30>      LLVM :: Linker/2011-08-18-unique-debug-type2.ll
30>      LLVM :: Linker/2011-08-22-ResolveAlias2.ll
30>      LLVM :: Linker/DbgDeclare2.ll
30>      LLVM :: Linker/available_externally_b.ll
30>      LLVM :: Linker/linkmdnode2.ll
30>      LLVM :: Linker/linknamedmdnode2.ll
30>      LLVM :: Linker/metadata-b.ll
30>      LLVM :: Linker/module-flags-1-b.ll
30>      LLVM :: Linker/module-flags-2-b.ll
30>      LLVM :: Linker/module-flags-3-b.ll
30>      LLVM :: Linker/module-flags-4-b.ll
30>      LLVM :: Linker/module-flags-5-b.ll
30>      LLVM :: Linker/module-flags-6-b.ll
30>      LLVM :: Linker/module-flags-7-b.ll
30>      LLVM :: Linker/module-flags-8-b.ll
30>      LLVM :: Linker/partial-type-refinement-link.ll
30>      LLVM :: Linker/testlink2.ll
30>      LLVM :: Linker/unnamed-addr1-b.ll
30>      LLVM :: Linker/visibility2.ll
30>      LLVM :: MC/MachO/gen-dwarf-producer.s
30>      LLVM :: Transforms/ArgumentPromotion/byval-2.ll
30>      LLVM :: Transforms/LoopSimplify/indirectbr.ll
30>
30>    Expected Passes    : 8274
30>    Expected Failures  : 52
30>    Unsupported Tests  : 264
30>    Unexpected Failures: 32
Is it "normal" ?