Показаны сообщения с ярлыком llvm. Показать все сообщения
Показаны сообщения с ярлыком llvm. Показать все сообщения

суббота, 19 сентября 2026 г.

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

вторник, 31 марта 2026 г.

dumping llvm bitcode from cicc

requires building of hijacked .so with appropriate LLVM version. I am too lazy for this
 
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

среда, 18 марта 2026 г.

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

пятница, 8 января 2016 г.

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)

вторник, 1 октября 2013 г.

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" ?