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

read "the anatomy of go"

As you can expect if book about language with deliberately castrated syntax (even python has something similar to ternary operator btw) has 864 pages - it detailed everything about it's runtime. The only couple of missed things are

1. integration with CGO
2. generic methods - but they were introduced in version 1.27 month ago

I personally especially appreciated chapter 7 about integration of GC into native code with multithreading - with detailed description of what happened between "stop the world" and end of sweeping phase

So if you doing occasionally RE of go binaries - this book is absolutely Must Read

And now my cynical thoughts about go language

Pros

Idea of usermode threads scheduling + polling of sockets is very good and allows processing of network traffic with minimal latency (although I have a vague doubt that it was stolen from Erlang)

functions are first-class citizens and this allows to do things like fp-go (sadly practically no one uses it anyway)

Cons

I personally think that biggest flaw in design of go is lack of normal macros allowing modification of AST. Having such macros lets you write your own syntax with blackjack and patter matching

io_uring support looks unorganic- there are lots of 3rd party libraries but no embedding in language itself like they done with sockets polling. I think this is very strange - the run-time already has almost all code for it in Windows IO Completion processing logic

go lang implements it's own codegen - no LLVM or other monsters. At one hand this allows very fast compilation, but on other - it means that they lack of SSE/AVX (or generally speaking - automatic loop vectorization)

Implementation of generics is weird. Сommon objection against C++ VTBLs is that indirect calls cause code cache miss. Ok, but golang went even further - they have two calls for generics:

1. first for interface dictionary loading
2. second for shapes
   

Poor stack management - especially for recursive functions. They manage size of stack for every function call and in case of stack enlargement

1. allocate new memory for stack
2. copy whole content from old stack to newly allocated memory
   

As you can expect this is not cheap operation, and there is good question - do you really need to have for each frame contiguous region of memory? Highly likely that better solution would be chain of more short stack frames each allocated in heap at least for cold paths/recursive calls

defered code executes at end of function - and so it can't be used as C++ RAII

PGO is very limited  - for example they have hardcoded table class_to_size consisting of 67 elements and describing the sizes of pools for allocated objects. The authors provide no explanation as to why they chose precisely these sizes, but PGO could collect statistics about allocations of you application and made new table tailored specially for you