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

new ptx instructions from cuda sdk 13.4

File with extracted instructions, attributes masks and operands - 1669 items (vs 1420 in version from 13.0). Being lazy person I also wrote script to compare them (in Perl, of course). 21 newly added instructions:

• applypriority.async.bulk (1 forms)
• applypriority.async.bulk.tensor (12 forms)
• clmad.hi (2 forms)
• clmad.lo (2 forms)
• fabric.submit (1 forms)
• fabric.try_atom.async (13 forms)
• fabric.try_get.async (2 forms)
• fabric.try_get.async.tensor (10 forms)
• fabric.try_pullred.async (5 forms)
• fabric.try_put.async (8 forms)
• fabric.try_put.async.tensor (20 forms)
• fabric.try_red.async (16 forms)
• fabric.try_red.async.tensor (20 forms)
• fabric.wait (1 forms)
• mbarrier.check_layout (1 forms)
• multimem.red.async (1 forms)
• multimem.st.async (3 forms)
• spcompress (1 forms)
• spdecompress (1 forms)
• tcgen05.ld.red.spcompress (1 forms)
• tcgen05.ld.spcompress (1 forms)

It seems that attribute mask has the same size 20 bytes:
270 uniq ins
FD FF FF FF F1 FF FF 9F F9 FF E7 CF FF F3 DF FF FF 00 00 00
length of mask 121
291 uniq ins
FD FF FF FF FB FF FF FF E7 FF 9F 3F FF FF 7F FF FF FF 29 00
length of mask 138

As you can see they added 138 - 121 = 17 new attributes. But worst of all, the attribute indices have changed too. Lets compare for example first form for ld instruction:
old:
60 00 00 00 A0 EC 03 00 00 00 00 00 01 00 00 00 00 00 00 00
new
C0 00 00 00 A0 EC 03 00 00 00 00 00 04 00 00 80 00 00 00 00

Also I made option -f to compare new forms - 127 new forms, like
perl ip_cmp.pl -f ptx_ops3.txt
--- 249 lines
add: 11 old forms, 17 new
 line 751: add    ddd    Q32Q32
 line 752: add    dcd    Q32R16
 line 753: add    012    H64H32H64
 line 754: add    012    H64E32H64
 line 755: add    012    H32H64H64
 line 756: add    012    E32H64H64

etc

inefficient uniform registers allocation in ptxas

While poking around RT kernels (extracted from libcuda.so) I stumbled upon very remarkable sample - sm159.elf. First let's check what tools were used to compile it - dump section .note.nv.tkinfo (here it has index 5):

readelf -x 5 sm159.elf

Hex dump of section '.note.nv.tkinfo':
  0x00000000 0c000000 9c000000 d0070000 4e564944 ............NVID
  0x00000010 49412043 6f727000 02000000 00000000 IA Corp.........
  0x00000020 01000000 07000000 35000000 54000000 ........5...T...
  0x00000030 00707478 61730043 75646120 636f6d70 .ptxas.Cuda comp
  0x00000040 696c6174 696f6e20 746f6f6c 732c2072 ilation tools, r
  0x00000050 656c6561 73652031 332e312c 20563133 elease 13.1, V13
  0x00000060 2e312e30 00427569 6c642063 7564615f .1.0.Build cuda_
  0x00000070 612e6375 64615f61 2e333630 34343438 a.cuda_a.3604448
  0x00000080 365f3000 2d6d6178 72726567 636f756e 6_0.-maxrregcoun
  0x00000090 74203332 202d6172 63682073 6d5f3132 t 32 -arch sm_12
  0x000000a0 31202d64 6c636d20 6367202d 6d203634 1 -dlcm cg -m 64
  0x000000b0 20000000                             ... 
As you can see it was compiled with regular ptxas from SDK 13.1 (and not with nvasm_internal). This file contains huge function _ZN6RefTMA12TMACommand_tC1ERKNS_19ptx_inst_parametersE - first branch happens at offset 0xc70 and first back-edge at 37c0. Let's take a look at this piece of code:
    /*01b0*/      LDC.64 R28, c[0x0][0x358] ; move base ptr into R28-29
    /*01c0*/      R2UR UR4, R28.reuse ; assign UR4
    /*01d0*/      R2UR UR5, R29.reuse ;
    /*01e0*/      R2UR UR6, R28 ;
    /*01f0*/      R2UR UR7, R29 ;
    /*0200*/      R2UR UR8, R28.reuse ; assign the same value from R28-29
    /*0210*/      R2UR UR9, R29.reuse ;
    /*0220*/      R2UR UR10, R28.reuse ;
    /*0230*/      R2UR UR11, R29.reuse ;
    /*0240*/      ST.E desc[UR4][R16.64+0x18], R3 ; use UR4 for desc operand
    /*0250*/      R2UR UR12, R28.reuse ;
    /*0260*/      R2UR UR13, R29 ;
    /*0270*/      ST.E desc[UR6][R16.64+0x38], R3 ;
    /*0280*/      R2UR UR70, R28 ;
    /*0290*/      R2UR UR71, R29.reuse ;
    /*02a0*/      ST.E desc[UR8][R16.64+0x58], R3 ;
    /*02b0*/      R2UR UR68, R28.reuse ;
    /*02c0*/      R2UR UR69, R29.reuse ;
    /*02d0*/      ST.E desc[UR10][R16.64+0x78], R3 ;
    /*02e0*/      R2UR UR66, R28.reuse ;
    /*02f0*/      R2UR UR67, R29 ;
    /*0300*/      ST.E desc[UR12][R16.64+0x118], R5 ;

What noteworthy things can we examine here?

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%)