Hacking by Breaking: Five Strategies to Reverse Engineer Opaque Compilers

A technical guide outlines five reverse-engineering strategies for diagnosing performance problems in closed-source ML compilers such as libtpu.so, proprietary MLIR accelerator plugins, and cuBLASLt routines, where a four-fold performance drop in a production LLM pipeline can otherwise take six months to resolve through vendor support. The strategies include feeding malformed inputs to trigger internal assertions and recover embedded .rodata strings via `strings -a` after `strip --strip-debug` or `strip --strip-all`, intercepting dynamic symbols through LD_PRELOAD shims on cuLaunchKernel and cublasLtMatmul, sniffing vtables in GDB, extracting in-memory IR bytecode, and disassembling PTX/SASS hardware kernels. The approach matters because vendors' support portals demand minimal standalone reproducers that proprietary model weights and legal restrictions often prevent engineers from sharing.

Hacking by Breaking: Five Strategies to Reverse Engineer Opaque Compilers You hit a four-fold performance drop in a production LLM pipeline. The profiler points straight at a closed-source binary blob: libtpu.so , a proprietary MLIR accelerator plugin, or a closed routine in cuBLASLt. You file a support ticket with the hardware vendor. Their portal demands a minimal standalone reproducer. But your computation graph uses proprietary model weights, your company’s legal team refuses to share the code, and the vendor’s engineering turnaround takes six months. You have no source code. You have a stripped .so file, a core dump, and an SLA to hit. The standard response is to treat the compiler as an inscrutable black box and resort to superstitious trial-and-error—tweaking batch sizes, inserting dummy operations, or padding tensor shapes in the dark. I hate guessing at compiler heuristics. Superstition has no place in systems engineering. Compilers cannot hide their true behavior. A compiler is a deterministic state machine that must allocate host memory, emit hardware-specific instructions, conform to operating system calling conventions, and enforce thousands of internal invariant assertions. Even when a vendor strips .symtab , disables debug logging, and packs its backend into a monolithic shared library, the binary remains physically bound to the machine it runs on. When an opaque compiler refuses to tell you how it works, you break it on purpose. You feed it malformed inputs to trigger its internal assertions, intercept its dynamic symbols, sniff its vtables in a debugger, extract its in-memory IR bytecodes, and disassemble its hardware kernels. php flowchart TD A "Host Python / C++ Frontend" -- |"Malformed Inputs Strategy 1 "| B "Compiler Pipeline Opaque .so " B -- |"Memory Scraper Strategy 5 "| C "In-Memory IR / Bytecode<br/ LLVM 0xDEC04342, MLIR ML\xefR " B -- |"DWARF CFI / .eh frame"| D "Crash Handler / Stderr<br/ Assertion Archaeology " B -- |"System V ABI / Registers Strategy 3 "| E "Debugger / GDB<br/ Vtable Sniffing & RTTI " B -- |"GOT/PLT Interception Strategy 2 "| F "LD PRELOAD Shim<br/ cuLaunchKernel, cublasLtMatmul " F -- |"Hardware Profiling & SASS Strategy 4 "| G "Accelerator Silicon<br/ PTX / SASS / LLM Reconstruction " Strategy 1: Deliberate Malformations and Assertion Archaeology Compilers are assert-driven software. A modern optimizing compiler like LLVM, XLA, or a proprietary device backend runs hundreds of transformation passes. Each pass relies on rigid preconditions: tensor dimensions must be non-zero, memory strides must be multiples of the vector register width, and induction variables must not overflow. When an opaque compiler is running normally, it runs silent. It swallows its intermediate representations, runs its internal heuristics, and emits machine code. You learn nothing. When you intentionally violate one of its invariants, that abstraction collapses. In C and C++ compiler codebases, assertions are built with macros: standard assert , Google’s CHECK and DCHECK in Abseil, or LLVM’s llvm unreachable and report fatal error . When one of these checks fails, the macro formats a diagnostic string containing the source file name FILE , line number LINE , failing condition, and function signature PRETTY FUNCTION , writes it directly to stderr , and calls abort . Vendors run strip --strip-debug or strip --strip-all on their release binaries to remove the symbol table .symtab and debug information .debug . But strip cannot touch string literals in .rodata . Those strings are program data, directly referenced by code pointers in the text segment. Every assertion message, file path, pass name, and internal type name remains embedded in the executable image. 1 2 3 4 5 6 7 8 bash Exposing .rodata strings from a stripped vendor ML compiler blob $ strings -a libproprietary compiler.so | grep -E "Pass|Assertion|Pattern" | head -n 8 /workspace/src/compiler/transforms/gemm fusion pass.cc:342 Assertion operand- layout .is strided row major failed in GemmFusionPass::RunOnBlock /workspace/src/compiler/ir/tile scheduler.cc:118 CHECK tile m % 16 == 0 failed: Tile M dimension must be aligned to 16 for Tensor Core MMA ProprietaryTritonEmitter::EmitAsyncCopyGlobalToShared MultiBufferingPass: pipeline depth must not exceed 4 stages To weaponize this, you craft malformed computation graphs designed to probe edge cases: 1. Non-contiguous strides: Instead of standard contiguous tensors, feed inputs with negative strides, prime-number leading dimensions, or transposed views without a contiguous copy. 2. Dimension boundaries: Feed tensors with shape 1, 0, 128 , batch sizes that are large prime numbers, or shapes that exceed 32-bit integer limits 2^31 - 1 elements . 3. Illegal dtype combinations: In modern mixed-precision pipelines, supply an FP8 E4M3 tensor to an operation that natively expects FP8 E5M2 or BF16 without inserting an explicit conversion cast. 4. Cyclic or illegal dependencies: In JIT runtimes that support custom fusion, construct an operation that consumes its own output across a broadcast boundary. When the binary aborts, it produces a stack trace. On Linux, the 64-bit System V ABI requires unwind information in the .eh frame and .eh frame hdr sections for C++ exception handling. This unwind table cannot be stripped without breaking language semantics. 1 2 3 4 5 6 7 8 9 FATAL Assertion tile k % 16 == 0 failed at /build/backend/gpu/mma emitter.cc:512 Function: proprietary::gpu::MmaEmitter::EmitWarpMma proprietary::ir::GemmTileNode Stack trace: 0 0x00007ff1a2b04192 in GI abort from /lib/x86 64-linux-gnu/libc.so.6 1 0x00007ff198124b80 in proprietary::logging::LogMessageFatal::~LogMessageFatal 2 0x00007ff198452c11 in proprietary::gpu::MmaEmitter::EmitWarpMma proprietary::ir::GemmTileNode 3 0x00007ff198458920 in proprietary::gpu::passes::TmaLoweringPass::RunOnBlock proprietary::ir::Block 4 0x00007ff198460144 in proprietary::gpu::PassManager::Run proprietary::ir::Module 5 0x00007ff19841029a in proprietary compile graph entry From a single deliberate failure, the architecture of the opaque compiler is laid bare: - We have the exact C++ namespaces proprietary::gpu::passes . - We know the pass ordering: TmaLoweringPass executes before MmaEmitter . - We learn that the hardware lowering uses Tensor Memory Accelerator TMA and Warp-Level Matrix Multiply MMA . - We extract the exact architectural constraint: the inner loop tiles the $K$ dimension in multiples of 16, and our non-standard reduction dimension broke the unrolled accumulator lowering. Strategy 2: Dynamic Interception via GOT/PLT and Driver Shims Proprietary ML runtimes and compilers are distributed as shared libraries .so on Linux, .dylib on macOS . These libraries do not execute in a vacuum. To interact with host memory, launch kernels, or invoke math routines, they must link against shared system libraries. This linkage creates a clean interception boundary. In the ELF format, position-independent code resolves external function calls through two tables: the Procedure Linkage Table PLT and the Global Offset Table GOT . By default, dynamic symbols use lazy binding: the PLT stub calls into the runtime dynamic linker ld-linux.so , which resolves the symbol address, writes it into the GOT, and jumps to the target code. Subsequent calls jump directly through the GOT. By setting the LD PRELOAD environment variable, we force the dynamic linker to bind our custom symbols ahead of any standard libraries. Instead of writing a generic shim for malloc or free , we target domain-specific runtime interfaces. If the opaque compiler invokes Nvidia’s cuBLAS or cuBLASLt for GEMMs, we can intercept cublasLtMatmul <sup 1 fn:1 </sup . If it generates and launches custom GPU kernels, it must ultimately call the CUDA Driver API: cuLaunchKernel or cuLaunchKernelEx <sup 2 fn:2 </sup . Even when a vendor statically links the high-level CUDA Runtime library libcudart static.a into their proprietary .so , they cannot statically link the CUDA Driver libcuda.so.1 . The driver talks directly to the kernel-mode driver /dev/nvidiactl and /dev/nvidia . Every kernel launch, without exception, must pass through cuLaunchKernel . Here is an interception shim that hooks cuLaunchKernel , unpacks its launch bounds, inspects the parameter buffer, and checks pointer alignment: 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 // shim cuda launch.c // Compile: gcc -O3 -shared -fPIC -o shim cuda launch.so shim cuda launch.c -ldl define GNU SOURCE include <stdio.h include <dlfcn.h include <stdint.h include <cuda.h // 1. Function pointer signature matching the CUDA Driver API. static CUresult real cuLaunchKernel CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void kernelParams, void extra = NULL; CUresult cuLaunchKernel CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void kernelParams, void extra { // 2. Resolve the real symbol from libcuda.so via RTLD NEXT on first invocation. if real cuLaunchKernel { real cuLaunchKernel = CUresult CUfunction, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, CUstream, void , void dlsym RTLD NEXT, "cuLaunchKernel" ; if real cuLaunchKernel { fprintf stderr, " SHIM ERROR Failed to resolve real cuLaunchKernel: %s\n", dlerror ; return CUDA ERROR UNKNOWN; } } // 3. Log execution parameters. fprintf stderr, "\n INTERCEPT cuLaunchKernel: Grid= %u,%u,%u Block= %u,%u,%u SharedMem=%u bytes\n", gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes ; // 4. Inspect kernel arguments if passed via the standard kernelParams array. if kernelParams { fprintf stderr, " Parameters:\n" ; for int i = 0; i < 6; ++i { void param ptr = kernelParams i ; if param ptr break; uintptr t val = uintptr t param ptr; // 5. Check if the parameter is a device memory address typically 47-bit on Linux x86-64 . if val 0x1000000000ULL { // Check pointer alignment: 16 bytes for float4/LDG.128, 128 bytes for L2 cache lines. unsigned int align = val & 0x7f == 0 ? 128 : val & 0xf == 0 ? 16 : 1 ; fprintf stderr, " Arg %d : Device Ptr 0x%lx Alignment: %u bytes \n", i, val, align ; } else { fprintf stderr, " Arg %d : Scalar / Offset 0x%lx %lu \n", i, val, val ; } } } // 6. Forward execution to the hardware driver. return real cuLaunchKernel f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra ; } Run your workload with the shim injected: 1 bash $ LD PRELOAD=./shim cuda launch.so python run inference.py The output immediately exposes the lowering decisions of the black box: - Grid and Block Dimensions: A grid of 64, 1, 1 with a block of 128, 1, 1 means the compiler mapped each tile to a single warp group 4 warps = 128 threads across 64 Streaming Multiprocessors. - Dynamic Shared Memory: Allocating exactly 98304 bytes 96 KB reveals that the compiler is targeting a multi-stage pipeline designed to fit within the L1/shared memory boundary of an NVIDIA Ampere or Ada Lovelace SM typically 100 KB or 128 KB configured capacity . - Device Pointer Alignment: Every pointer argument has a 128-byte alignment, confirming that the compiler relies on vector load instructions LDG.E.128 and expects data buffers to be aligned to hardware cache sector boundaries. The Dynamic Linking Pitfall: RTLD LOCAL and dlopen Shims Some opaque frameworks guard against LD PRELOAD by opening their backends at runtime using dlopen "backend.so", RTLD NOW | RTLD LOCAL . Symbols loaded with RTLD LOCAL are not added to the global symbol search list, meaning their internal cross-calls do not resolve against your preloaded library. To defeat this, preload dlopen itself: 1 2 3 4 5 6 7 8 9 js void dlopen const char filename, int flags { static void real dlopen const char , int = NULL; if real dlopen real dlopen = dlsym RTLD NEXT, "dlopen" ; // Force all dynamically loaded libraries into the global symbol table int forced flags = flags & ~RTLD LOCAL | RTLD GLOBAL; fprintf stderr, " SHIM Intercepted dlopen %s , forcing RTLD GLOBAL\n", filename ? filename : "self" ; return real dlopen filename, forced flags ; } By flipping RTLD LOCAL to RTLD GLOBAL , every dynamically loaded library is pulled into the global scope where your shims can intercept its PLT stubs. Strategy 3: ABI Archaeology, Vtable Sniffing, and Hardware PMUs When an opaque compiler calls internal functions across module boundaries, arguments are passed in physical hardware registers according to the Application Binary Interface ABI . On x86-64 Linux, the System V AMD64 ABI<sup 3 fn:3 </sup dictates: - First six integer or pointer arguments are passed in %rdi , %rsi , %rdx , %rcx , %r8 , %r9 . - Floating-point arguments are passed in %xmm0 through %xmm7 . - For C++ non-static member functions, %rdi holds the implicit this pointer. If you set a breakpoint at an unstripped entry point in GDB or attach to a JIT process during execution , inspecting %rdi gives you direct access to the compiler’s internal object model. Vtable Sniffing under the Itanium C++ ABI Almost all C++ compilers on Linux follow the Itanium C++ ABI<sup 4 fn:4 </sup . According to the specification, any polymorphic C++ class a class with at least one virtual function begins with an 8-byte pointer to its virtual method table vptr : 1 2 3 4 php Object in Memory: %rdi + 0x00 - points to vtable entry 0 first virtual method %rdi - 0x08 - points to std::type info for the class RTTI %rdi - 0x10 - offset-to-top distance from this pointer to start of most-derived object At index -1 8 bytes before the first virtual function address , the vtable stores a pointer to the class’s std::type info structure. Under the Itanium ABI, the std::type info structure contains a pointer to an internal null-terminated byte string holding the mangled type name at offset +0x08 : 1 2 3 4 5 // std::type info internal layout Itanium ABI struct type info { void vptr; // Virtual table for std::type info const char type name; // Mangled type name e.g., "N14xla compiler14TritonFusionOpE" }; This layout allows us to write a GDB Python script that sniffs any opaque object in %rdi , traverses its vtable backwards, extracts the mangled RTTI string, and demangles it using c++filt : 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 vtable sniffer.py Usage in GDB: source vtable sniffer.py; sniff $rdi import gdb import subprocess class SniffObjectCommand gdb.Command : """Sniff the true C++ runtime type of an object pointer in memory.""" def init self : super SniffObjectCommand, self . init "sniff", gdb.COMMAND USER def invoke self, arg, from tty : 1. Evaluate the pointer expression e.g., $rdi ptr = gdb.parse and eval arg ptr val = int ptr if ptr val == 0: print " SNIFF Null pointer." return try: 2. Dereference the object to get the vtable pointer first 8 bytes vptr = int gdb.parse and eval f" unsigned long {ptr val}" 3. Read the type info pointer at vptr - 8 bytes typeinfo ptr = int gdb.parse and eval f" unsigned long {vptr} - 8 " 4. Read the mangled name string pointer at typeinfo ptr + 8 bytes mangled name ptr = int gdb.parse and eval f" unsigned long {typeinfo ptr} + 8 " 5. Read the null-terminated string from target process memory mangled name = gdb.selected inferior .read memory mangled name ptr, 128 .tobytes .split b'\x00' 0 .decode 'utf-8' 6. Demangle the symbol using c++filt demangled = subprocess.check output "c++filt", "-t", mangled name .decode 'utf-8' .strip print f" SNIFF Object at {hex ptr val } has true type: {demangled} mangled: {mangled name} " except Exception as e: print f" SNIFF ERROR Failed to traverse vtable: {e}" SniffObjectCommand When you hit a breakpoint inside an opaque pass: 1 2 gdb sniff $rdi SNIFF Object at 0x7fffdb108a40 has true type: xla::gpu::TritonFusionAnalysis mangled: ZN3xla3gpu20TritonFusionAnalysisE Without seeing a single line of vendor source code, you have recovered the exact polymorphic class hierarchy, allowing you to deduce the data structure that represents the optimization state. Hardware Performance Counters and SASS Archaeology Host-side debugging tells you how the compiler organizes its data. To see how it schedules code for execution on physical hardware, we use vendor hardware performance counters via NVIDIA Nsight Compute ncu or AMD Omnitrace. When a compiler targets modern hardware like an NVIDIA Hopper H100 or Ada Lovelace GPU, it faces three critical physical constraints: 1. Register Pressure vs. Occupancy: An SM has 64K 32-bit registers. If the compiler allocates 255 registers per thread, it can schedule at most 256 threads 2 warps per SM, slashing theoretical occupancy to 12.5%. If it drops to 64 registers, occupancy can hit 50%. 2. Shared Memory Bank Conflicts: Shared memory has 32 banks, each 4 bytes wide. If 32 threads in a warp access the same bank at different addresses, the transaction serializes 32-way, turning a 1-cycle access into a 32-cycle pipeline stall. 3. Warp Scheduler Stalls: Is the kernel memory-bound waiting on long-scoreboard loads from HBM/DRAM or latency-bound waiting on short-scoreboard arithmetic dependencies ? Profile the generated kernel: 1 2 3 4 bash $ ncu --metrics smsp pipe shared cycles active.avg.pct of peak sustained elapsed,\ l1tex data bank conflicts pipe lsu mem shared op ld.sum,\ smsp warp issue stalled wait pct \ python run inference.py If l1tex data bank conflicts drops to zero on a non-transposed matrix multiplication, the compiler implemented either: - Shared memory padding: Padding the leading dimension of the tile allocation by 1 or 2 elements e.g., shared half tile 16 16 + 1 . - XOR swizzling: Applying a bitwise XOR between the row index and column bits when computing the shared memory offset offset = row ^ col 3 . By looking at hardware counters, you reverse the compiler’s heuristics directly from physical execution invariants. Strategy 4: Disassembly, ISA Semantics, and LLM Reconstruction When dynamic analysis is impractical, static disassembly is the alternative. But reading raw x86-64 assembly from a stripped compiler binary is grueling. Decades of legacy x86 cruft, register renaming, and compiler scheduling obfuscate high-level program intent. Accelerator ISAs are different. Unlike general-purpose CPUs, accelerator ISAs—like NVIDIA PTX Parallel Thread Execution <sup 5 fn:5 </sup , NVIDIA SASS Streaming Assembler , and AMD CDNA/RDNA assembly—are designed specifically for linear algebra and structured tensor memory operations. Their instructions operate directly on matrix tiles, warp groups, barrier synchronizations, and asynchronous copy pipelines. Because these instructions are semantically rich, reconstructing them into high-level source code is remarkably straightforward. Extracting Device Binaries NVIDIA device binaries are embedded inside host .so files within ELF sections named .nv fatbin or .nvFatBin . Extract them with NVIDIA’s binary utilities: 1 2 3 4 5 bash Dump the high-level PTX intermediate representation if preserved $ cuobjdump --dump-ptx libopaque backend.so kernels.ptx Disassemble the native machine instructions SASS from the compiled cubin $ nvdisasm -g -c libopaque backend.so kernels.sass Consider this disassembled PTX snippet extracted from an opaque matrix multiplication kernel: // Extracted from an opaque kernel targeting SM80+ Ampere/Hopper .visible .entry gemm kernel m16n16k16 .param .u64 A ptr, .param .u64 B ptr, .param .u64 C ptr, .param .u64 D ptr { .reg .b64 %rd<5 ; .reg .b32 %r<5 ; .reg .f32 %f<5 ; ld.param.u64 %rd1, A ptr ; ld.param.u64 %rd2, B ptr ; ld.param.u64 %rd3, C ptr ; ld.param.u64 %rd4, D ptr ; // 1. Asynchronous copy from global memory to shared memory bypassing registers cp.async.ca.shared.global %r1 , %rd1 , 16; cp.async.commit group; cp.async.wait group 0; bar.sync 0; // 2. Warp-Level Matrix Multiply-Accumulate WMMA wmma.load.a.sync.aligned.m16n16k16.f16.col %r1, %rd1 ; wmma.load.b.sync.aligned.m16n16k16.f16.row %r2, %rd2 ; wmma.load.c.sync.aligned.m16n16k16.f32.row %r3, %rd3 ; wmma.mma.sync.aligned.m16n16k16.al.f32.f16.f16.f32 %r4, %r1, %r2, %r3; wmma.store.d.sync.aligned.m16n16k16.f32.row %rd4 , %r4; ret; } Notice what is happening in this assembly: 1. cp.async.ca.shared.global : The kernel uses asynchronous copy instructions to stream 16-byte chunks directly from HBM into shared memory without burning register file bandwidth. 2. wmma.load. .m16n16k16 : The warp synchronizes to load $16 \times 16 \times 16$ tile fragments. Matrix A is column-major .col , Matrix B is row-major .row , and the inputs are FP16 .f16 . 3. wmma.mma.sync. : The warp-level tensor core instruction computes $D = A \times B + C$, accumulating into 32-bit floats .f32 . LLM Reconstruction Workflow Because accelerator instructions enforce strict algebraic contracts, large language models can decompile them with near-perfect fidelity. Feed the raw assembly to an LLM with the launch parameters obtained in Strategy 2: System Prompt: You are a senior systems and GPU compiler engineer. Decompile the following PTX snippet into equivalent, idiomatic CUDA C++ using WMMA primitives. Maintain all memory layouts, tile dimensions, and synchronization boundaries. The LLM reconstructs the clean CUDA source: 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 include <mma.h include <cuda fp16.h using namespace nvcuda; global void gemm kernel m16n16k16 const half restrict A, const half restrict B, const float restrict C, float restrict D { // 1. Declare WMMA fragments matching the m16n16k16 shape and datatypes wmma::fragment<wmma::matrix a, 16, 16, 16, half, wmma::col major frag a; wmma::fragment<wmma::matrix b, 16, 16, 16, half, wmma::row major frag b; wmma::fragment<wmma::accumulator, 16, 16, 16, float frag c; wmma::fragment<wmma::accumulator, 16, 16, 16, float frag d; // 2. Load operands from memory into fragment registers wmma::load matrix sync frag a, A, 16 ; wmma::load matrix sync frag b, B, 16 ; wmma::load matrix sync frag c, C, 16, wmma::mem row major ; // 3. Perform the warp-synchronous matrix multiply-accumulate wmma::mma sync frag d, frag a, frag b, frag c ; // 4. Store the resulting accumulator tile back to global memory wmma::store matrix sync D, frag d, 16, wmma::mem row major ; } The loop is closed: we took a stripped binary, intercepted its launch parameters via cuLaunchKernel , extracted the SASS/PTX, and reconstructed human-readable CUDA C++ that can be modified, optimized, and recompiled without the vendor’s proprietary library. Strategy 5: Dormant Registries and In-Memory Bytecode Sniffing Every compiler team leaves debug backdoors in production binaries. Building a complex compiler without intermediate print statements, pass dumps, and verification flags is impossible. When vendors ship production binaries, they often strip CLI flags from their public documentation or disable user-facing --verbose arguments. But ripping out the actual internal debug printing infrastructure is risky—it creates conditional compilation divergence between debug and release builds. Consequently, the diagnostic flags are usually left compiled into the binary, waiting to be turned on. The Global Option Registry in LLVM and MLIR Most modern compilers—including XLA, PyTorch Inductor backends, Triton, and proprietary vendor runtimes—are built on LLVM or MLIR. LLVM uses the llvm::cl::opt mechanism for command-line option registration<sup 6 fn:6 </sup . When an LLVM file defines a flag: 1 2 3 4 static llvm::cl::opt<bool PrintAfterAll "print-after-all", llvm::cl::desc "Print IR after each pass" , llvm::cl::init false ; This constructor registers the flag into a global static map llvm::cl::TopLevelSubCommand- OptionsMap at library load time. Even if the vendor’s top-level CLI strips --print-after-all , the flag registration still exists in memory. If the framework initializes LLVM using llvm::cl::ParseCommandLineOptions , it often reads the standard LLVM OPTIONS environment variable. Beyond generic LLVM options, frameworks expose their own environment variable escape hatches: - XLA: XLA FLAGS="--xla dump to=/tmp/xla dump --xla dump hlo as text" - Triton: TRITON PRINT IR=1 - PyTorch Inductor: TORCH COMPILE DEBUG=1 Patrick Toulmé used these exact dormant XLA flags to dump undocumented TPU VLIW machine instructions directly from JAX<sup 7 fn:7 </sup . By passing undocumented backend flags, the closed runtime happily wrote thousands of lines of low-level instruction bundles to /tmp/ , exposing hardware capabilities that the vendor had never publicly documented. Scraping Bytecode from Process Memory What if the vendor explicitly compiled out all text-dumping logic? They still cannot avoid constructing the Intermediate Representation IR in memory. A compiler cannot optimize a computation graph without holding the graph in memory. In modern architectures, that IR is represented as binary bytecode: LLVM Bitcode or MLIR Bytecode<sup 8 fn:8 </sup . These bytecodes begin with well-known magic numbers: - LLVM Bitcode: Begins with the 4-byte ASCII signature BC\xc0\xde little-endian unsigned 32-bit integer 0xDEC04342 . - MLIR Bytecode: Begins with the 4-byte signature ML\xefR 0x4D 0x4C 0xEF 0x52 . If you attach a debugger to the compiler process while it is running an optimization pass, or scan the process’s anonymous memory pages /proc/<pid /mem , you can search for these magic signatures: 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 dump in memory ir.py Scans an active compiler process memory space for LLVM bitcode or MLIR bytecode payloads. import re import sys LLVM BC MAGIC = b"\xde\xc0\x43\x42" 0xDEC04342 MLIR BC MAGIC = b"\x4d\x4c\xef\x52" ML\xefR def scan process memory pid : maps path = f"/proc/{pid}/maps" mem path = f"/proc/{pid}/mem" 1. Parse readable memory regions from /proc/<pid /maps regions = with open maps path, 'r' as f: for line in f: parts = line.strip .split addr range, perms = parts 0 , parts 1 if 'r' in perms and 'heap' in line or 'anon' in line or len parts == 5 : start, end = int x, 16 for x in addr range.split '-' regions.append start, end print f" Scanning {len regions } memory regions in PID {pid}..." 2. Open /proc/<pid /mem to search for bytecode signatures with open mem path, 'rb', 0 as mem: for start, end in regions: try: mem.seek start chunk = mem.read end - start except IOError, OSError : continue Check for LLVM bitcode for match in re.finditer re.escape LLVM BC MAGIC , chunk : offset = match.start payload addr = start + offset print f" + Found LLVM Bitcode magic at virtual address {hex payload addr }" Dump 64KB following the magic header data = chunk offset:offset + 65536 filename = f"dump {pid} llvm {hex payload addr }.bc" with open filename, "wb" as out: out.write data print f" Dumped bitcode to {filename}" Check for MLIR bytecode for match in re.finditer re.escape MLIR BC MAGIC , chunk : offset = match.start payload addr = start + offset print f" + Found MLIR Bytecode magic at virtual address {hex payload addr }" data = chunk offset:offset + 65536 filename = f"dump {pid} mlir {hex payload addr }.mlirbc" with open filename, "wb" as out: out.write data print f" Dumped MLIR bytecode to {filename}" if name == " main ": if len sys.argv = 2: print "Usage: python dump in memory ir.py <PID " sys.exit 1 scan process memory int sys.argv 1 Once you dump the memory buffer to disk, you run standard open-source tools: 1 2 3 4 5 bash Disassemble the dumped LLVM bitcode into human-readable LLVM IR $ llvm-dis dump 1234 llvm 0x7fffdb108000.bc -o dump.ll Or translate the dumped MLIR bytecode into high-level text format $ mlir-translate --mlir-to-text dump 1234 mlir 0x7fffdb205000.mlirbc -o dump.mlir The proprietary optimization pipeline vanishes. You are no longer dealing with a stripped binary; you have the complete, annotated Single Static Assignment SSA form of the program, complete with loop bounds, memory access indices, vector widths, and custom compiler attributes. The Systems Takeaway Vendors distribute closed-source compilers and proprietary backend blobs to protect intellectual property and lock developers into specific hardware ecosystems. That opacity is an illusion. A compiler cannot execute without anchoring itself to physical reality. It must conform to host calling conventions. It must link against shared operating system libraries. It must structure its memory to avoid hardware stalls. And it must emit machine instructions that obey the physical constraints of silicon registers and memory buses. When an opaque compiler misbehaves: 1. Don’t guess. Force internal assertions by supplying deliberately malformed bounds and strides. 2. Intercept its dynamic calls to cuLaunchKernel or cublasLtMatmul using LD PRELOAD to capture exact execution shapes and memory alignments. 3. Sniff polymorphic class hierarchies through Itanium ABI vtables and correlate them with hardware performance counters. 4. Extract stripped accelerator binaries and use the semantic density of matrix instructions to decompile them with an LLM. 5. Search process memory for LLVM 0xDEC04342 or MLIR ML\xefR bytecode signatures to recover the raw intermediate representation. Treat opaque compilers not as black boxes, but as distributed systems waiting to be instrumented. References Disclaimer: Researched and drafted with AI assistance Gemini 3.1 Pro, Gemini 3.8 Flash . Direction, technical judgment, and final edits are mine; every claim is traceable to the sources cited above. 1. NVIDIA cuBLASLt Library Documentation: High-level APIs for flexible GEMM executions, epilogues, and matrix descriptors. Link https://docs.nvidia.com/cuda/cublas/index.html cublaslt-matmul ↩︎ fnref:1 2. CUDA Driver API Reference - Kernel Execution: cuLaunchKernel and execution parameter specifications. Link https://docs.nvidia.com/cuda/cuda-driver-api/group CUDA EXEC.html ↩︎ fnref:2 3. System V Application Binary Interface AMD64 Architecture Processor Supplement: Defines standard register usage, calling conventions, and stack layout for x86-64 UNIX systems. Link https://gitlab.com/x86-psabis/x86-64-ABI ↩︎ fnref:3 4. Itanium C++ ABI Specification: The industry-standard specification for virtual table layouts, RTTI structures, and exception unwinding on Linux x86-64. Link https://itanium-cxx-abi.github.io/cxx-abi/abi.html ↩︎ fnref:4 5. NVIDIA PTX ISA Reference Manual: Documentation on Warp-Level Matrix Multiply-Accumulate WMMA/MMA instructions and asynchronous copy primitives. Link https://docs.nvidia.com/cuda/parallel-thread-execution/index.html ↩︎ fnref:5 6. LLVM CommandLine 2.0 Reference Manual: Implementation details of llvm::cl::opt and global option registration. Link https://llvm.org/docs/CommandLine.html ↩︎ fnref:6 7. Patrick Toulmé, “From JAX to VLIW: Tracing a Computation”: Using undocumented XLA flags to extract TPU machine instructions and VLIW execution bundles. Link https://patricktoulme.substack.com/p/from-jax-to-vliw-tracing-a-computation ↩︎ fnref:7 8. MLIR Bytecode Format Specification: Binary encoding format and magic signature definitions for MLIR bytecode. Link https://mlir.llvm.org/docs/BytecodeFormat/ ↩︎ fnref:8 CC BY 4.0 https://creativecommons.org/licenses/by/4.0/ by the author.