https://maknee.github.io/blog/2026/NVIDIA-TileIR-Internals-from-CuTile-to-MLIR-LLVM-to-SASS/ Henry Zhu Toggle navigation * about * blog (current) * ramblings * publications * projects * subscribe * NVIDIA TileIR Internals: from CuTile to MLIR/LLVM to SASS Created in January 30, 2026 2026 * TileIR * TileIR In this post, we'll dig deep into how TileIR works, from how it generates instructions to analyzing its different passes. We'll trace how a Mixture-of-Experts (MoE) kernel written in CuTile gets compiled down through cuda_tile - nv_tileaa - nv_tileas - NVVM - LLVM - SASS. Here's what to expect: * What is CuTile? -- The tile-centric programming model * Running Example -- An MoE kernel we'll trace through every stage * The Dialects -- From cuda_tile through nv_tileaa and nv_tileas to NVVM/LLVM * The Passes -- TileIR passes: what they do and when they run Based on CUDA 13.1. Some details are undocumented and may change in future releases. What is CuTile? [cutile] CuTile separates user responsibility (splitting work into blocks and tiles) from system responsibility (mapping to threads) (Image source: GPU MODE) CuTile is NVIDIA's new "tile-centric" programming model for modern NVIDIA GPUs. This abstraction is powerful: CuTile lets the programmer think in terms of tiles rather than threads, while the compiler handles the complexity of coordinating hundreds of threads across fragmented data. A single CuTile line ct.mma(a, b, acc) could get transformed to many tensor core instructions. What is TileIR? TileIR is NVIDIA's MLIR-based compiler infrastructure that powers CuTile. It progressively lowers your high-level tensor operations through multiple MLIR dialects and NVIDIA specific tools: [pipeline_overview] TileIR compilation pipeline: Python - SASS The user-facing tool is tileirasLike ptxas but for TileIR. Yes, NVIDIA named it "tile-ir-as" (tile IR assembler)., which orchestrates this entire pipeline. --------------------------------------------------------------------- Running Example: MoE Kernel Throughout this post, we'll trace this MoE (Mixture of Experts) kernel through every compilation stage. This is code from NVIDIA's cutile-python samplesThere's also a C++ API: NVIDIA/cuda-tile. Operations like ct.gather, ct.mma, cuda_tile.load_view_tko documented in TileIR docs.: @ct.kernel def fused_moe_kernel( A, # Input tokens, shape (batch, K) B, # Expert weights, shape (num_experts, N, K) C, # Output tensor, shape (num_tokens * topk, N) topk_weights, # Router weights for each token-expert pair sorted_token_ids, # Token indices sorted by expert assignment sorted_expert_ids, # Expert index for each TILE_M num_token_replicas: int, mul_routed_weight: ConstBool, TILE_M: ConstInt, TILE_N: ConstInt, TILE_K: ConstInt, ): M = sorted_token_ids.shape[0] N = B.shape[1] K = B.shape[2] GROUP_SIZE_M = 8 bid_m, bid_n = swizzle_2d(M, N, TILE_M, TILE_N, GROUP_SIZE_M) # - cuda_tile.get_tile_block_id # Gather token indices for this block token_id_indices = bid_m * TILE_M + ct.arange(TILE_M, dtype=ct.int32) token_ids = ct.gather(sorted_token_ids, token_id_indices) # - cuda_tile.load_view_tko a_row_indices = token_ids // num_token_replicas expert_id = ct.load(sorted_expert_ids, index=bid_m, shape=()) # - cuda_tile.load_ptr_tko # Initialize accumulator accumulator = ct.full((TILE_M, TILE_N), 0.0, dtype=ct.float32) # - cuda_tile.constant for k in range(0, ct.cdiv(K, TILE_K)): # - cuda_tile.for # Load A tile (gathered by token indices) a_col_indices = k * TILE_K + ct.arange(TILE_K, dtype=ct.int32) a = ct.gather(A, (a_row_indices[:, None], a_col_indices[None, :])) # - cuda_tile.load_view_tko # Load B tile (expert weights) b = ct.load(B, (expert_id, k, bid_n), shape=(1, TILE_K, TILE_N), order=(0, 2, 1)).reshape((TILE_K, TILE_N)) # - cuda_tile.load_ptr_tko accumulator = ct.mma(a, b, accumulator) # - cuda_tile.mmaf - THE COMPUTE! if mul_routed_weight: moe_weight = ct.gather(topk_weights, token_ids) accumulator = accumulator * moe_weight[:, None] # - cuda_tile.mulf # Scatter results back to output c_col_indices = bid_n * TILE_N + ct.arange(TILE_N, dtype=ct.int32) accumulator = ct.astype(accumulator, C.dtype) # - cuda_tile.ftof ct.scatter(C, (token_ids[:, None], c_col_indices[None, :]), accumulator) # - cuda_tile.store_ptr_tko The three key operations we'll trace: Python cuda_tile What it does ct.gather(A, load_view_tko Gather tokens by expert assignment indices) (indirect load) ct.load(B, ...) load_ptr_tko Load expert weights (direct load) ct.mma(a, b, mmaf Matrix multiply-accumulate on tensor acc) cores Watch how these transform through nv_tileaa, nv_tileas and finally to SASS instructions. --------------------------------------------------------------------- Compiling with tileiras The tileiras command-line tool is the ahead-of-time compiler that transforms .cutile bytecode into GPU binaries. tileiras --gpu-name sm_120 MoE.cutile -o moe.cubin Undocumented Environment Variables These TileIR-specific environment variables affect compilation: Variable Description TILEIR_ALWAYS_SWIZZLE Force swizzle mode TILEIR_PREFER_TMA_FOR_LOAD_STORE Prefer TMA for all load/store operations TILEIR_DELAY_TMA_STORE_WAIT Delay TMA store wait (optimization for overlapping compute) Interesting undocumented CLI options The --print-before-all flag dumps LLVM IR before each compilation pass. $ tileiras --print-before-all --gpu-name=sm_120 MoE.cutile -o moe.cubin 2>&1 *** IR Dump Before Add __emutls_[vt]. variables for emultated TLS model (lower-emutls) *** ; ModuleID = 'LLVMDialectModule' source_filename = "LLVMDialectModule" target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128" @__CUDA_TILEIR_FUNC_NAME_0 = internal constant [17 x i8] c"fused_moe_kernel\00" ... All LLVM passes dumped (27 unique passes) *** IR Dump Before Add __emutls_[vt]. variables for emultated TLS model (lower-emutls) *** *** IR Dump Before Canonicalize natural loops (loop-simplify) *** *** IR Dump Before CodeGen Prepare (codegenprepare) *** *** IR Dump Before Constant Hoisting (consthoist) *** *** IR Dump Before Exception handling preparation (dwarf-eh-prepare) *** *** IR Dump Before Expand Atomic instructions (atomic-expand) *** *** IR Dump Before Expand fp (expand-fp) *** *** IR Dump Before Expand indirectbr instructions (indirectbr-expand) *** *** IR Dump Before Expand large div/rem (expand-large-div-rem) *** *** IR Dump Before Expand memcmp() to load/stores (expand-memcmp) *** *** IR Dump Before Expand reduction intrinsics (expand-reductions) *** *** IR Dump Before Instrument function entry/exit with calls to e.g. mcount() (post-inline-ee-instrument) *** *** IR Dump Before Interleaved Access Pass (interleaved-access) *** *** IR Dump Before Lower AMX intrinsics (lower-amx-intrinsics) *** *** IR Dump Before Lower AMX type for load/store (lower-amx-type) *** *** IR Dump Before Lower Garbage Collection Instructions (gc-lowering) *** *** IR Dump Before Merge contiguous icmps into a memcmp (mergeicmps) *** *** IR Dump Before ObjC ARC contraction (objc-arc-contract) *** *** IR Dump Before Partially inline calls to library functions (partially-inline-libcalls) *** *** IR Dump Before Pre-ISel Intrinsic Lowering (pre-isel-intrinsic-lowering) *** *** IR Dump Before Prepare callbr (callbrprepare) *** *** IR Dump Before Remove unreachable blocks from the CFG (unreachableblockelim) *** *** IR Dump Before Replace intrinsics with calls to vector library (replace-with-veclib) *** *** IR Dump Before Safe Stack instrumentation pass (safe-stack) *** *** IR Dump Before Scalarize Masked Memory Intrinsics (scalarize-masked-mem-intrin) *** *** IR Dump Before Shadow Stack GC Lowering (shadow-stack-gc-lowering) *** *** IR Dump Before X86 Partial Reduction (x86-partial-reduction) *** --------------------------------------------------------------------- Pipeline Overview [pipeline_overview] TileIR compilation pipeline: Python - SASS TileIR takes your CuTile Python code through a series of progressive lowerings: Stage Format Description Python CuTile API High-level tensor operations (make_tensor_view; mmaf) .cutile Bytecode Serialized representation of the kernel cuda_tile MLIR High-level tensor ops; architecture-independent Dialect nv_tileaa MLIR Tile-level ops; explicit memory references Dialect nv_tileas MLIR Scheduled ops; async pipelines Dialect LLVM/NVVM LLVM IR Standard LLVM with NVIDIA intrinsics PTX Assembly Virtual GPU assembly SASS Machine Native GPU instructions (sm_120) Code Each stage removes abstraction and adds architecture-specific detail. By the time we reach SASS, every memory access pattern, tensor core instruction, and synchronization barrier is explicit. --------------------------------------------------------------------- The Dialects TileIR uses three main MLIR dialects to represent computations at different abstraction levels. Let's trace our MoE kernel through each one: Python cuda_tile nv_tileaa nv_tileas SASS ct.gather load_view_tko tileaa.load_view tileas.utcpglobalmem UTCPMULTI / (A, idx) LDG ct.load load_ptr_tko tileaa.load_tko tileas.tcgen05_ld TCGEN05.LD.S (B, ...) ct.mma(a, mmaf tileaa.mmaf_tko tileas.tcgen05_mma TCGEN05.MMA b, c) cuda_tile: High-Level Tensor Operations [cuda_tile_dialect] cuda_tile dialect operations The cuda_tile dialect is closest to your Python code. Operations work on abstract tensor views without worrying about memory layout or hardware details. Key operations: * make_tensor_view - Create a view into a tensor with shape and strides * get_tile_block_id - Get the current thread block's position in the grid * load_view_tko / store_view_tko - Load/store tiles with token-based ordering * mmaf - Matrix multiply-accumulate (targets tensor cores) * for / continue - Loop constructs for K-dimension iteration MoE in cuda_tile Recall our MoE kernel above. Here's how the key operations map to cuda_tile IR: Python - cuda_tile mapping: Python (CuTile) cuda_tile IR Purpose ct.gather() load_view_tko Gather elements by indices ct.load() load_ptr_tko Load contiguous tile from memory ct.mma() mmaf Matrix multiply-accumulate (tensor cores) ct.scatter() store_ptr_tko Scatter elements to output ct.full() constant Initialize accumulator for k in range for/continue K-dimension iteration loop () ct.astype() ftof Type conversion (F32 - output dtype) Expand to see cuda_tile IR from MoE kernel key sections // cuda_tile dialect - MoE kernel %1 = "cuda_tile.constant"() : () -> (ct.view) // TILE_M %2 = "cuda_tile.constant"() : () -> (ct.view) // TILE_N %3 = "cuda_tile.constant"() : () -> (ct.view) // TILE_K %4 = "cuda_tile.assume"(%arg0) : (ct.view) -> (ct.view) %5 = "cuda_tile.assume"(%arg1) : (ct.view) -> (ct.view) %10 = "cuda_tile.make_tensor_view"(%4, %5, %6, %7, %8, %9) : (ct.view, ct.view, ct.view, ct.view, ct.view, ct.view) -> (ct.token) %11 = "cuda_tile.make_tensor_view"(%arg2, %arg3) : (ct.view, ct.view) -> (ct.token) %12 = "cuda_tile.make_token"() : () -> (ct.ptr) %20, %21, %22 = "cuda_tile.get_tile_block_id"() : () -> (ct.view, ct.view, ct.view) %23 = "cuda_tile.divi"(%4, %1) : (ct.view, ct.view) -> (ct.view) // M / TILE_M %24 = "cuda_tile.muli"(%1, %23) : (ct.view, ct.view) -> (ct.view) %25 = "cuda_tile.divi"(%20, %24) : (ct.view, ct.view) -> (ct.view) %30 = "cuda_tile.remi"(%20, %25) : (ct.view, ct.view) -> (ct.view) // expert routing %31 = "cuda_tile.cmpi"(%30, %1) : (ct.view, ct.view) -> (ct.view) %32 = "cuda_tile.select"(%31, %30, %25) : (ct.view, ct.view, ct.view) -> (ct.view) %40 = "cuda_tile.iota"() : () -> (ct.view) %41 = "cuda_tile.reshape"(%24) : (ct.view) -> (ct.view) %42 = "cuda_tile.broadcast"(%41) : (ct.view) -> (ct.view) %43 = "cuda_tile.addi"(%42, %40) : (ct.view, ct.view) -> (ct.view) %44 = "cuda_tile.offset"(%42, %43) : (ct.view, ct.view) -> (ct.view) %50, %51 = "cuda_tile.load_ptr_tko"(%44, %31, %42, %12) // ct.load() : (ct.view, ct.view, ct.view, ct.ptr) -> (ct.view, ct.ptr) %52 = "cuda_tile.make_partition_view"(%10) : (ct.token) -> (ct.part) %53, %54 = "cuda_tile.load_view_tko"(%52, %43, %12) // ct.gather() : (ct.part, ct.view, ct.ptr) -> (ct.view, ct.ptr) %60 = "cuda_tile.for"(%1, %23, %3, %arg4) {1 regions} // K-loop : (ct.view, ct.view, ct.view, ct.view) -> (ct.view) %61 = "cuda_tile.muli"(%iter, %3) : (ct.view, ct.view) -> (ct.view) %62 = "cuda_tile.broadcast"(%61) : (ct.view) -> (ct.view) %63, %64 = "cuda_tile.load_ptr_tko"(%62, %31, %42, %12) : (ct.view, ct.view, ct.view, ct.ptr) -> (ct.view, ct.ptr) %65, %66 = "cuda_tile.load_view_tko"(%52, %62, %12) : (ct.part, ct.view, ct.ptr) -> (ct.view, ct.ptr) %67 = "cuda_tile.mmaf"(%63, %65, %acc) // ct.mma() : (ct.view, ct.view, ct.view) -> (ct.view) "cuda_tile.continue"(%67) : (ct.view) -> () %70 = "cuda_tile.ftof"(%60) : (ct.view) -> (ct.view) // ct.astype() %71 = "cuda_tile.store_ptr_tko"(%44, %70, %31, %12) // ct.scatter() : (ct.view, ct.view, ct.view, ct.ptr) -> (ct.ptr) "cuda_tile.return"() nv_tileaa [nv_tileaa_dialect] nv_tileaa dialect operations The nv_tileaa dialect lowers tensor views to concrete memory references. This is where we start seeing explicit memory operations. Key changes from cuda_tile: * make_tensor_view - make_memref (explicit memory references) * get_tile_block_id - get_program_id (program-centric naming) * mmaf - dot (more explicit accumulation) * Explicit tiled_load / tiled_store with memory tokens * New ops: splat, broadcast, addptr for memory address calculations Expand to see nv_tileaa IR from MoE kernel key sections // nv_tileaa dialect - MoE kernel // Tile-level ops (architecture-independent) "nv_tileaa.func"() {nv_tileaa.kernel_spec} {1 regions} // Input validation %1 = "nv_tileaa.assume"(%arg0) : (aa.memref) -> (aa.memref) %2 = "nv_tileaa.assume"(%arg1) : (iN) -> (iN) %3 = "nv_tileaa.assume"(%2) : (iN) -> (iN) // Splat: scalar - tensor (for broadcasting) %10 = "nv_tileaa.splat"(%3) : (iN) -> (tensor<...>) %11 = "nv_tileaa.splat"(%2) : (iN) -> (tensor<...>) // Memory reference creation (lowered from make_tensor_view) %20 = "nv_tileaa.make_memref"(%1, %2, %3, %4, %5, %6) : (aa.memref, iN, iN, iN, iN, iN) -> (aa.btile) %21 = "nv_tileaa.make_memref"(%1, %2) : (aa.memref, iN) -> (aa.btile) %22 = "nv_tileaa.create_mem_token"() : () -> (aa.ptr) // Program indexing %30 = "nv_tileaa.get_program_id"() : () -> (iN) %31 = "nv_tileaa.splat"(%30) : (iN) -> (tensor<...>) %32 = "nv_tileaa.make_range"(%c0, %c128) : (iN, iN) -> (tensor<...>) %33 = "nv_tileaa.extract"(%32) : (tensor<...>) -> (iN) // Pointer arithmetic %40 = "nv_tileaa.splat"(%1) : (aa.memref) -> (tensor<...>) %41 = "nv_tileaa.addptr"(%40, %33) : (tensor<...>, tensor<...>) -> (tensor<...>) // Masked loads %50, %51 = "nv_tileaa.load"(%41, %mask, %c0, %22) : (tensor<...>, tensor<...>, tensor<...>, aa.ptr) -> (tensor<...>, aa.ptr) // Tiled memory operations %60 = "nv_tileaa.block_tile"(%20) : (aa.btile) -> (aa.mtoken) %61 = "nv_tileaa.extract"(%32) : (tensor<...>) -> (iN) %62, %63 = "nv_tileaa.tiled_load"(%60, %61, %22) : (aa.mtoken, iN, aa.ptr) -> (tensor<...>, aa.ptr) %64 = "nv_tileaa.view"(%62) : (tensor<...>) -> (tensor<...>) // Shape manipulation %70 = "nv_tileaa.expand_dims"(%33) : (tensor<...>) -> (tensor<...>) %71 = "nv_tileaa.broadcast"(%70) : (tensor<...>) -> (tensor<...>) // DOT OPERATION (lowered from cuda_tile.mmaf) %80 = "nv_tileaa.dot"(%50, %64, %acc) : (tensor<...>, tensor<...>, tensor<...>) -> (tensor<...>) // Output %90 = "nv_tileaa.fp_to_fp"(%80) : (tensor<...>) -> (tensor<...>) %91 = "nv_tileaa.store"(%41, %90, %mask, %22) : (tensor<...>, tensor<...>, tensor<...>, aa.ptr) -> (aa.ptr) "nv_tileaa.return"() Key transformations from cuda_tile - nv_tileaa: cuda_tile nv_tileaa Change make_tensor_view make_memref Abstract view - concrete memory ref get_tile_block_id get_program_id Tile-centric - program-centric naming mmaf dot High-level MMA - explicit dot product load_view_tko tiled_load + view Decomposed into separate ops ct.view types tensor<...> Abstract - explicit tensor shapes ct.token aa.btile; Memory tokens more specific aa.mtoken Pass #12 observation: The 32 fp_to_fp operations suggest this MoE kernel produces 32 output tiles that need precision conversion from F32 accumulator to the output dtype. nv_tileas [nv_tileas_tcgen05] nv_tileas dialect with tcgen05 operations The nv_tileas dialect is where architecture-specific code generation happens. This dialect introduces: Async Pipeline Operations: * async.pipeline.create - Create a software pipeline for overlapping compute/memory * producer_acquire / producer_commit - Acquire/release pipeline stages * consumer_wait / consumer_release - Synchronize consumers with producers Tensor Memory Operations: * tcgen05.alloc - Allocate dedicated tensor memory * tmem_load / tmem_store - Access tensor memory Tensor Core Operations: * tcgen05.mma - Matrix Multiply-Accumulate * block_scaled_mma - Block-scaled MMA for mixed precision * mma.fence - Memory fence for MMA operations Expand to see nv_tileas IR from MoE kernel key sections // nv_tileas dialect - MoE kernel // Tile-level Scheduled Assembly // Layout conversion and view operations %1, %2 = "nv_tileas.load"(%ptr, %mask, %c0, %token) : (tensor<...>, tensor<...>, tensor<...>, aa.ptr) -> (tensor<...>, aa.ptr) %3, %4 = "nv_tileas.tiled_load"(%btile, %idx, %token) : (aa.mtoken, iN, aa.ptr) -> (tensor<...>, aa.ptr) %5 = "nv_tileas.view"(%3) : (tensor<...>) -> (tensor<...>) // Convert layout for tensor cores %10 = "nv_tileas.convert_layout"(%bcast) : (tensor<...>) -> (tensor<...>) %11 = "nv_tileas.convert_layout"(%5) : (tensor<...>) -> (tensor<...>) %12 = "nv_tileas.convert_layout"(%1) : (tensor<...>) -> (tensor<...>) // DOT with input allowances %20 = "nv_tileas.dot"(%10, %11, %12, %c1) : (tensor<...>, tensor<...>, tensor<...>, iN) -> (tensor<...>) // TMA descriptor %25 = "nv_tileas.make_tiled_tma_desc"(%memref) {tmaIdx=0} : (aa.btile) -> (!tma.desc) // ASYNC PIPELINE (producer-consumer model) // Pipeline and iterator creation %30 = "nv_tileas.async.pipeline.create_pipeline"() : () -> (!pipeline) %31 = "nv_tileas.async.pipeline.create_pipeline"() : () -> (!pipeline) %32 = "nv_tileas.async.pipeline.create_iterator"(%30) : (!pipeline) -> (!iter) %33 = "nv_tileas.async.pipeline.create_iterator"(%31) : (!pipeline) -> (!iter) // Agent switch (4 regions for producer/consumer roles) "nv_tileas.async.pipeline.agent_switch"(%arg0, %30, %32, %31, %33) {4 regions} : (aa.memref, !pipeline, !iter, !pipeline, !iter) -> () // Tensor allocation (double-buffering) %40 = "nv_tileas.alloc_tensor"() : () -> (tensor<128x64xbf16>) %41 = "nv_tileas.alloc_tensor"() : () -> (tensor<64x128xbf16>) // Slice operations %50 = "nv_tileas.extract_slice"(%40, %c0) : (tensor<...>, iN) -> (tensor<...>) %51 = "nv_tileas.insert_slice"(%data, %40, %c0, %c64) : (tensor<...>, tensor<...>, iN, iN) -> (tensor<...>) // PRODUCER: acquire - write - commit %60 = "nv_tileas.async.pipeline.producer_acquire"(%30, %32) : (!pipeline, !iter) -> (!stage) %61 = "nv_tileas.async.pipeline.producer_write"(%60, %30) {1 regions} : (!stage, !pipeline) -> (!stage) %62 = "nv_tileas.async.load"(%51, %ptr, %mask, %c16) : (tensor<...>, tensor<...>, tensor<...>, tensor<...>) -> (!async) "nv_tileas.async.pipeline.yield"(%62) : (!async) -> () "nv_tileas.async.pipeline.producer_commit"(%30, %61) : (!pipeline, !stage) -> () // CONSUMER: wait - read - release %70 = "nv_tileas.async.pipeline.consumer_wait"(%31, %33) : (!pipeline, !iter) -> (!stage) %71, %72 = "nv_tileas.async.pipeline.consumer_read"(%70, %31) {1 regions} : (!stage, !pipeline) -> (!stage, tensor<...>) %73 = "nv_tileas.copy"(%buf) : (tensor<...>) -> (tensor<...>) "nv_tileas.async.pipeline.yield"(%73) : (tensor<...>) -> () "nv_tileas.async.pipeline.consumer_release"(%31, %71) : (!pipeline, !stage) -> () // Matrix multiply (100+ ops for tiled GEMM) %80 = "nv_tileas.dot"(%50, %72, %acc, %c1) : (tensor<...>, tensor<...>, tensor<...>, iN) -> (tensor<...>) %81 = "nv_tileas.dot"(%50, %72, %80, %c1) : (tensor<...>, tensor<...>, tensor<...>, iN) -> (tensor<...>) // TMA load %90 = "nv_tileas.async.tiled_tma_load"(%btile, %buf, %25, %idx, %c0, %c64) : (aa.mtoken, tensor<...>, !tma.desc, iN, iN, iN) -> (!async) // Output %100 = "nv_tileas.insert_slice"(%result, %41, %c0, %c0) : (tensor<...>, tensor<...>, iN, iN) -> (tensor<...>) %101 = "nv_tileas.view"(%100) : (tensor<...>) -> (tensor<...>) %102 = "nv_tileas.convert_layout"(%101) : (tensor<...>) -> (tensor<...>) NVVM + LLVM After nv_tileas, the compiler lowers to NVVM (NVIDIA's LLVM dialect) and then to standard LLVM IR. Key NVVM intrinsics: * @llvm.nvvm.mma.sync.* - Tensor core matrix multiply * @llvm.nvvm.ldmatrix.* - Load matrix fragments from shared memory * @llvm.nvvm.cp.async.* - Asynchronous memory copy * @llvm.nvvm.bar.warp.sync - Warp-level synchronization * @llvm.nvvm.tcgen05.* - Tensor core intrinsics Expand to see NVVM/LLVM IR key sections ; Thread ID and warp-level operations %233 = call range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() %234 = icmp eq i32 %233, 0 %235 = ashr i32 %233, 5 %236 = call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 -1, i32 %235, i32 0, i32 31) %237 = call { i32, i1 } @llvm.nvvm.elect.sync(i32 -1) ; Mbarrier initialization (async pipeline synchronization) call void @llvm.nvvm.mbarrier.init.shared( ptr addrspace(3) getelementptr inbounds nuw (i8, ptr addrspace(3) @global_smem, i64 82000), i32 %241) call void @llvm.nvvm.mbarrier.init.shared( ptr addrspace(3) getelementptr inbounds nuw (i8, ptr addrspace(3) @global_smem, i64 82008), i32 %241) ; Cluster-wide fence and barrier call void asm sideeffect "fence.mbarrier_init.release.cluster;", "n"(i32 0) call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0) ; Async copy from global to shared memory (cp.async) %1478 = select i1 %1459, i32 16, i32 0 call void @llvm.nvvm.cp.async.cg.shared.global.16.s( ptr addrspace(3) %1477, ptr addrspace(1) %1451, i32 %1478) call void @llvm.nvvm.cp.async.cg.shared.global.16.s( ptr addrspace(3) %1485, ptr addrspace(1) %1452, i32 %1486) ; Signal mbarrier arrival after async copy call void @llvm.nvvm.cp.async.mbarrier.arrive.noinc.shared(ptr addrspace(3) %1535) ; TCGEN05 tensor core intrinsics ; Allocate tensor memory %tmem = call i32 @llvm.nvvm.tcgen05.alloc(i32 65536) ; Load data into tensor memory call void @llvm.nvvm.tcgen05.ld(i32 %tmem, ptr addrspace(3) %smem_ptr, i32 %size) ; Execute TCGEN05 MMA (128x256x64 tile) call void @llvm.nvvm.tcgen05.mma(i32 %tmem_a, i32 %tmem_b, i32 %tmem_c) ; Fence and wait for tensor core completion call void @llvm.nvvm.tcgen05.fence() call void @llvm.nvvm.tcgen05.wait() SASS The final output is SASS. Key SASS instructions: * HMMA.16816.F32.BF16 - Half-precision matrix multiply-accumulate * TCGEN05.MMA - Tensor core MMA * TCGEN05.LD.S - Tensor memory load * UTCPMULTI / LDG - Global memory loads * SYNCS.EXCH - Async synchronization exchange * FENCE.VIEW.ASYNC.S - Async memory fence Expand to see SASS key sections ; SASS - MoE kernel (fused_moe_kernel) ; Target: sm_120a ; Thread ID and CTA setup /*0020*/ S2R R0, SR_TID.X ; ; Get thread ID /*0060*/ S2UR UR8, SR_CgaCtaId ; ; Get CTA ID (uniform reg) ; Async fence and mbarrier sync (cluster sync) /*0110*/ FENCE.VIEW.ASYNC.S ; /*0120*/ SYNCS.EXCH.64 URZ, [UR8+0x14050], UR4 ; /*0130*/ SYNCS.EXCH.64 URZ, [UR8+0x14058], UR4 ; /*0140*/ SYNCS.EXCH.64 URZ, [UR8+0x14060], UR6 ; ; ... (data loading, address calculation) ... ; Tensor core HMMA - 16x8x16 BF16-F32 matrix multiply ; R156 = A matrix fragment (reused across 7 HMMAs) ; R124,R120,R116,R112,R108,R104,R100 = B matrix fragments ; R200,R204,R64,R60,R56,R52,R48 = accumulator tiles /*4a00*/ HMMA.16816.F32.BF16 R200, R156, R124, R200 ; /*4a10*/ HMMA.16816.F32.BF16 R204, R156, R120, R204 ; /*4a20*/ HMMA.16816.F32.BF16 R64, R156, R116, R64 ; /*4a30*/ HMMA.16816.F32.BF16 R60, R156, R112, R60 ; /*4a40*/ HMMA.16816.F32.BF16 R56, R156, R108, R56 ; /*4a50*/ HMMA.16816.F32.BF16 R52, R156, R104, R52 ; /*4a60*/ HMMA.16816.F32.BF16 R48, R156, R100, R48 ; ; Second A fragment (R148) with different B fragments /*4a70*/ HMMA.16816.F32.BF16 R200, R148, R126, R200 ; /*4a80*/ HMMA.16816.F32.BF16 R204, R148, R122, R204 ; /*4a90*/ HMMA.16816.F32.BF16 R64, R148, R118, R64 ; --------------------------------------------------------------------- The TileIR passes TileIR runs multiple passes to transform your code. The passes are grouped by the scope they operate on: [pass_flow] TileIR pass pipeline [pass_glossary] Detailed pass pipeline: cuda_tile.entry - nv_tileaa.func (x12) - builtin.module - gpu.module --------------------------------------------------------------------- Pass 1: cuda_tile.entry Entry point canonicalization--validates kernel structure, emits compile-time constants for tile sizes/strides, propagates input constraints via assume operations, creates tensor views, and establishes memory ordering via make_token. --------------------------------------------------------------------- Pass 2: nv_tileaa.func (x12 iterations) Iterative lowering from cuda_tile to nv_tileaa. First iteration converts make_tensor_view - make_memref, get_tile_block_id - get_program_id, mmaf - dot, decomposes load_view_tko into block_tile + tiled_load + view. Subsequent iterations perform refinement and optimization. Final iteration emits precision conversions (fp_to_fp), adds kernel metadata, and prepares for async pipeline lowering. --------------------------------------------------------------------- Pass 3: builtin.module Module-level transforms and nv_tileas emission--creates async pipeline operations, software pipelines for overlapping compute/memory, producer-consumer synchronization, TMA descriptors, and double buffers. --------------------------------------------------------------------- Pass 4: gpu.module Final lowering to NVVM/LLVM--converts nv_tileas.dot - nvvm.mma.sync, lowers async ops to barrier/fence instructions, converts memory ops to NVVM intrinsics (ldmatrix, cp.async, mbarrier.*), and emits address space annotations. Complete Pass Catalog Below is a catalog of passes that run within the TileIR pipeline. Conversion Passes Pass Name Source Target Description Frontend: convert-cudatile-to-tileaa cuda_tile nv_tileaa CuTile DSL to TileAA abstract assembly Middle-end: convert-tileaa-to-tileas nv_tileaa nv_tileas Abstract to scheduled assembly convert-nv-tileas-to-llvm nv_tileas llvm Backend: TileAS to LLVM IR Convert tile convert-nv-tile-func-to-llvm nv_tile llvm function ops to LLVM convert-gpu-to-nvvm gpu nvvm GPU dialect to NVVM intrinsics Structured convert-scf-to-cf scf cf control flow to basic blocks nv-tile-ir-convert-target-to-nvvm nv_tile nvvm Target-specific ops to NVVM Async pipeline convert-pipeline-to-nvvm pipeline nvvm ops to NVVM barriers Arithmetic convert-arith-to-llvm arith llvm operations to LLVM convert-cf-to-llvm cf llvm Control flow to LLVM Generic convert-to-llvm * llvm catch-all LLVM conversion convert-math-to-llvm math llvm Math operations to LLVM convert-nvvm-to-llvm nvvm llvm NVVM intrinsics to LLVM Undefined convert-ub-to-llvm ub llvm behavior ops to LLVM convert-vector-to-llvm vector llvm Vector ops to LLVM convert-debuginfo-to-llvm debug llvm Debug info to LLVM metadata TileAS Optimization Passes Pass Name Description tileas-assign-dot-layouts Assign optimal data layouts for dot (MMA) operations tileas-assign-pipeline-layouts Assign layouts for async pipeline stages tileas-assign-load-store-layouts Assign layouts for memory operations tileas-attach-tma-desc-args Attach TMA descriptor arguments to kernel signature tileas-dynamic-persistent Enable dynamic persistent kernel execution tileas-insert-OCG-knobs Insert Online Code Generation tuning knobs tileas-legalize-tmem-copy Legalize tensor memory copy operations tileas-plan-cta Plan CTA (thread block) configuration tileas-remove-buffer-alias Remove buffer aliasing for optimization tileas-remove-dead-args Dead argument elimination tileas-remove-layout-conversions Remove unnecessary layout conversions tileas-resolve-agent-boundary Resolve warp specialization agent boundaries tileas-slicing Tensor slicing for pipelining tileas-materialize-async Materialize async load/store/dot operations tileas-materialize-convert-layout Materialize layout conversion copy atoms tileas-materialize-schedule Materialize schedule to warp-specialized IR tileas-unroll-register-loops Unroll loops at register level tileas-unspecialized-pipeline Handle non-warp-specialized pipelines tileas-optimize-alloc-tensor Optimize tensor allocation placement tileas-optimize-reduce Optimize reduction operations tileas-recompute-for-scheduling Recompute values for better scheduling tileas-legalize-fma-dot Legalize FMA in dot products tileas-legalize-reduce Legalize reduction operations tileas-slice-and-fuse Slice and fuse operations for locality tileas-refine-atom-by-resource Refine copy atoms based on resource constraints tileas-generate-schedule Generate execution schedule (Serial or CostBased) tileas-prepare-for-scheduling Prepare IR for scheduling pass tileas-optimize-dot-accumulation Optimize dot product accumulation lower-tma-load-store-to-async Lower TMA ops to async variants tileas-print-decomposed-tv-layout Debug: print decomposed tensor view layouts --------------------------------------------------------------------- Conversion Patterns Registered The TileAA-TileAS conversion registers 20+ patterns: TileAAToTileASTiledLoadOpPattern // Tiled load conversion TileAAToTileASDotOpPattern // Dot product conversion TileAAToTileASExtractOpPattern // Extraction conversion TileAAToTileASBroadcastOpPattern // Broadcast conversion TileAAToTileASGatherLoadOpPattern // Gather load conversion TileAAToTileASScatterStoreOpPattern // Scatter store conversion TileAAToTileASExpandDimsOpPattern // Dimension expansion TileAAToTileASExtractSliceOpPattern // Slice extraction TileAAToTileASGenerateOpPattern // Generate conversion TileAAToTileASLoadOpPattern // Load conversion TileAAToTileASPermuteOpPattern // Permute conversion TileAAToTileASReduceOpPattern // Reduce conversion TileAAToTileASScanOpPattern // Scan conversion TileAAToTileASStoreOpPattern // Store conversion TileAAToTileASTiledAtomicRMWOpPattern // Atomic RMW conversion TileAAToTileASTiledStoreOpPattern // Tiled store conversion TileAAToTileASViewOpPattern // View conversion TileAAToTileASYieldOpPattern // Yield conversion --------------------------------------------------------------------- Conclusion TileIR is a sophisticated MLIR-based compiler that progressively lowers high-level tensor operations to optimized GPU machine code. It's an interesting piece of software that combines MLIR and the rest of NVIDIA's toolchain to make the tile abstraction work. Resources: * CuTile Python * CUDA Tile * NVIDIA TileIR Documentation --------------------------------------------------------------------- Appendix: TileIR Passes Reference This appendix documents the TileIR-specific passes in the compilation pipeline. Passes are organized into categories: Conversion and TileAS Optimization Conversion Passes (16) Conversion passes transform IR between MLIR dialects. convert-cudatile-to-tileaa Converts high-level cuda_tile dialect to nv_tileaa. Key transformations: * cuda_tile.mmaf - nv_tileaa.dot * cuda_tile.load_view_tko - nv_tileaa.tiled_load * cuda_tile.store_ptr_tko - nv_tileaa.tiled_store * cuda_tile.for - scf.for + nv_tileaa.yield void ConvertCudaTileToTileAA::runOnOperation() { ModuleOp module = getOperation(); ConversionTarget target(getContext()); target.addLegalDialect(); target.addIllegalDialect(); RewritePatternSet patterns(&getContext()); // Register 20+ conversion patterns patterns.add(...); patterns.add(...); patterns.add(...); applyPartialConversion(module, target, std::move(patterns)); } --------------------------------------------------------------------- convert-tileaa-to-tileas Main middle-end conversion: nv_tileaa - nv_tileas (Tile Assembly). Key transformations: * nv_tileaa.tiled_load - nv_tileas.async_load + pipeline ops * nv_tileaa.dot - nv_tileas.dot with layout annotations * Inserts shared memory allocations void ConvertTileAAToTileAS::runOnOperation() { FuncOp funcOp = getOperation(); // Walk all tileaa operations funcOp.walk([&](nv_tileaa::TiledLoadOp loadOp) { // Create async copy with TMA descriptor auto asyncCopy = builder.create(...); // Allocate shared memory buffer auto smemAlloc = builder.create(...); }); funcOp.walk([&](nv_tileaa::DotOp dotOp) { // Convert to tileas.dot with layout attributes auto tiledDot = builder.create(...); tiledDot->setAttr("lhs_layout", selectMMALayout(...)); }); } --------------------------------------------------------------------- convert-nv-tileas-to-llvm Backend code generation: nv_tileas - LLVM IR with NVVM intrinsics. Key transformations: * tileas.tcgen05_mma - @llvm.nvvm.tcgen05.mma.* * tileas.tcgen05_ld - @llvm.nvvm.tcgen05.ld.* * tileas.async_copy - @llvm.nvvm.cp.async.* * Barrier ops - @llvm.nvvm.barrier.* void ConvertTileASToLLVM::runOnOperation() { ModuleOp module = getOperation(); ConversionTarget target(getContext()); target.addLegalDialect(); RewritePatternSet patterns(&getContext()); // MMA operations patterns.add([](tcgen05::MMAOp op) { // Generate NVVM MMA intrinsic return builder.create(...); }); // Memory operations with TMA patterns.add([](tcgen05::LoadOp op) { return builder.create(...); }); } --------------------------------------------------------------------- convert-gpu-to-nvvm Converts GPU dialect operations to NVVM intrinsics. GPU Op NVVM Intrinsic gpu.thread_id nvvm.read.ptx.sreg.tid.* gpu.block_id nvvm.read.ptx.sreg.ctaid.* gpu.block_dim nvvm.read.ptx.sreg.ntid.* gpu.barrier nvvm.barrier0 --------------------------------------------------------------------- convert-pipeline-to-nvvm Converts async pipeline operations to NVVM barrier intrinsics. Pipeline Op NVVM Op pipeline.producer_acquire nvvm.mbarrier.arrive.* pipeline.producer_commit nvvm.mbarrier.arrive.* + phase pipeline.consumer_wait nvvm.mbarrier.wait.* pipeline.consumer_release nvvm.mbarrier.arrive.* --------------------------------------------------------------------- TileAS Optimization Passes (30) TileAS passes optimize and schedule tile operations. tileas-assign-dot-layouts Assigns MMA-compatible layouts to dot product operands. void AssignDotLayouts::runOnOperation() { FuncOp funcOp = getOperation(); funcOp.walk([&](DotOp dotOp) { auto lhsType = dotOp.getLhs().getType(); auto rhsType = dotOp.getRhs().getType(); // Select MMA shape based on types MMAShape mmaShape = selectMMAShape(lhsType, rhsType); // Assign layouts for operands Layout lhsLayout = computeLhsLayout(mmaShape, lhsType); Layout rhsLayout = computeRhsLayout(mmaShape, rhsType); dotOp->setAttr("lhs_layout", lhsLayout); dotOp->setAttr("rhs_layout", rhsLayout); }); } MMA shapes: m16n8k16, m16n16k16, m64n256k64 --------------------------------------------------------------------- tileas-assign-load-store-layouts Optimizes memory access patterns for coalesced loads. void AssignLoadStoreLayouts::runOnOperation() { FuncOp funcOp = getOperation(); funcOp.walk([&](LoadOp loadOp) { auto tensorType = loadOp.getResult().getType(); // Check for TMA opportunity if (canUseTMA(loadOp)) { Layout tmaLayout = computeTMALayout(tensorType); loadOp->setAttr("layout", tmaLayout); loadOp->setAttr("use_tma", true); } else { // Vectorized load layout Layout vecLayout = computeVectorizedLayout(tensorType); loadOp->setAttr("layout", vecLayout); } }); } --------------------------------------------------------------------- tileas-assign-pipeline-layouts Assigns layouts for async pipeline buffers. void AssignPipelineLayouts::runOnOperation() { FuncOp funcOp = getOperation(); funcOp.walk([&](PipelineOp pipelineOp) { for (auto& stage : pipelineOp.getStages()) { // Assign shared memory layouts for buffers for (auto buffer : stage.getBuffers()) { Layout smemLayout = computeSwizzledLayout(buffer); buffer->setAttr("layout", smemLayout); } } }); } --------------------------------------------------------------------- tileas-generate-schedule Generates execution schedule using cost-based or serial scheduler. void GenerateSchedule::runOnOperation() { FuncOp funcOp = getOperation(); // Build dependency graph DependencyGraph depGraph(funcOp); // Select scheduler based on options Scheduler* scheduler; if (useCostBasedScheduler) { scheduler = new CostBasedScheduler(depGraph); } else { scheduler = new SerialScheduler(depGraph); } // Generate schedule Schedule schedule = scheduler->generateSchedule(); // Apply schedule to IR applySchedule(funcOp, schedule); } Scheduler types: * Serial: Topological order * CostBased: Latency-aware with heuristics --------------------------------------------------------------------- tileas-materialize-schedule Materializes abstract schedule into warp-specialized IR. void MaterializeSchedule::runOnOperation() { FuncOp funcOp = getOperation(); Schedule schedule = getSchedule(funcOp); if (schedule.getStrategy() == Strategy::WarpSpecialize) { // Split into producer/consumer auto [producerOps, consumerOps] = partitionOps(funcOp, schedule); // Create agent regions createAgentRegion(producerOps, AgentRole::Producer); createAgentRegion(consumerOps, AgentRole::Consumer); // Insert synchronization insertBarriers(funcOp, schedule); } } --------------------------------------------------------------------- tileas-materialize-async Creates async pipeline structure with multi-buffering. void MaterializeAsync::runOnOperation() { FuncOp funcOp = getOperation(); int numStages = getOption("num-stages"); funcOp.walk([&](scf::ForOp forOp) { if (canPipeline(forOp)) { // Create N buffers for N-stage pipeline SmallVector buffers; for (int i = 0; i < numStages; i++) { buffers.push_back(allocateBuffer(forOp)); } // Transform loop body emitPrologue(forOp, buffers); emitSteadyState(forOp, buffers); emitEpilogue(forOp, buffers); } }); } --------------------------------------------------------------------- tileas-materialize-convert-layout Expands layout conversions to actual data movement. void MaterializeConvertLayout::runOnOperation() { FuncOp funcOp = getOperation(); funcOp.walk([&](ConvertLayoutOp convertOp) { auto srcLayout = getLayout(convertOp.getSource()); auto dstLayout = getLayout(convertOp.getResult()); // Generate shuffle or shared memory path if (canUseShuffles(srcLayout, dstLayout)) { emitShuffleConversion(convertOp); } else { emitSharedMemoryConversion(convertOp); } }); } --------------------------------------------------------------------- tileas-attach-tma-desc-args Injects TMA descriptor arguments into kernel signatures. void AttachTMADescArgs::runOnOperation() { FuncOp funcOp = getOperation(); SmallVector tmaOps; funcOp.walk([&](Operation* op) { if (usesTMA(op)) tmaOps.push_back(op); }); for (auto& tmaOp : tmaOps) { // Create TMA descriptor type auto descType = TMADescriptorType::get( tmaOp.getShape(), tmaOp.getElementType(), tmaOp.getSwizzle() ); // Add to function arguments funcOp.insertArgument(descType, "tma_desc"); } } --------------------------------------------------------------------- tileas-slicing Slices tensors for pipelined execution. void TileASSlicing::runOnOperation() { FuncOp funcOp = getOperation(); funcOp.walk([&](LoadOp loadOp) { auto tensorType = loadOp.getResult().getType(); int sliceDim = getSliceDimension(loadOp); int sliceSize = computeSliceSize(tensorType, sliceDim); // Replace single load with sliced loads SmallVector slices; for (int i = 0; i < numSlices; i++) { auto slice = builder.create( loadOp.getSource(), sliceDim, i * sliceSize, sliceSize ); slices.push_back(slice); } }); } --------------------------------------------------------------------- tileas-plan-cta Plans CTA (thread block) configuration. void PlanCTA::runOnOperation() { FuncOp funcOp = getOperation(); // Analyze resource requirements int smemRequired = analyzeSharedMemory(funcOp); int regsRequired = analyzeRegisters(funcOp); // Compute optimal CTA shape CTAConfig config = computeCTAConfig( smemRequired, regsRequired, targetOccupancy ); funcOp->setAttr("cta_shape", config.toAttribute()); } --------------------------------------------------------------------- tileas-resolve-agent-boundary Resolves data flow across warp specialization boundaries. void ResolveAgentBoundary::runOnOperation() { FuncOp funcOp = getOperation(); funcOp.walk([&](AgentSwitchOp switchOp) { // Identify values crossing boundary SmallVector crossingValues; for (Value v : switchOp.getOperands()) { if (crossesBoundary(v, switchOp)) { crossingValues.push_back(v); } } // Insert shared memory communication for (Value v : crossingValues) { insertSharedMemoryTransfer(v, switchOp); } }); } --------------------------------------------------------------------- tileas-remove-buffer-alias Removes buffer aliasing using fixed-point iteration. void RemoveBufferAlias::runOnOperation() { FuncOp funcOp = getOperation(); bool changed = true; while (changed) { changed = false; funcOp.walk([&](AllocTensorOp allocOp) { for (auto& use : allocOp.getResult().getUses()) { if (isAliasingUse(use)) { createNonAliasingBuffer(use); changed = true; } } }); } } --------------------------------------------------------------------- tileas-remove-dead-args Removes unused arguments from region operations. --------------------------------------------------------------------- tileas-remove-layout-conversions Eliminates redundant layout conversions. void RemoveLayoutConversions::runOnOperation() { FuncOp funcOp = getOperation(); funcOp.walk([&](ConvertLayoutOp convertOp) { auto srcLayout = getLayout(convertOp.getSource()); auto dstLayout = getLayout(convertOp.getResult()); // Remove identity conversions if (srcLayout == dstLayout) { convertOp.replaceAllUsesWith(convertOp.getSource()); convertOp.erase(); } }); } --------------------------------------------------------------------- tileas-optimize-alloc-tensor Optimizes tensor allocations through reuse and elimination. void OptimizeAllocTensor::runOnOperation() { FuncOp funcOp = getOperation(); LivenessAnalysis liveness(funcOp); SmallVector allocs; funcOp.walk([&](AllocTensorOp op) { allocs.push_back(op); }); for (auto& alloc : allocs) { // Find reusable buffer if (auto reusable = findReusableBuffer(alloc, liveness)) { reuseBuffer(alloc, reusable); } } } --------------------------------------------------------------------- tileas-optimize-reduce Optimizes reduction operations with warp shuffle or shared memory. void OptimizeReduce::runOnOperation() { FuncOp funcOp = getOperation(); funcOp.walk([&](ReduceOp reduceOp) { int reductionSize = getReductionSize(reduceOp); if (reductionSize <= 32) { setAtom(reduceOp, "warp_shuffle"); } else if (reductionSize <= 1024) { setAtom(reduceOp, "shared_memory"); } else { setAtom(reduceOp, "multi_stage"); } }); } --------------------------------------------------------------------- tileas-optimize-dot-accumulation Optimizes MMA accumulation patterns for better register utilization. void OptimizeDotAccumulation::runOnOperation() { FuncOp funcOp = getOperation(); funcOp.walk([&](DotOp dotOp) { auto accumPattern = analyzeAccumulationPattern(dotOp); switch (accumPattern) { case AccumPattern::SimpleLoop: optimizeSimpleAccumulation(dotOp); break; case AccumPattern::SplitK: optimizeSplitKAccumulation(dotOp); break; case AccumPattern::StreamK: optimizeStreamKAccumulation(dotOp); break; } }); } --------------------------------------------------------------------- tileas-recompute-for-scheduling Trades recomputation for reduced register pressure. void TileASRecomputeForScheduling::runOnOperation() { FuncOp funcOp = getOperation(); RegisterPressureAnalysis regPressure(funcOp); funcOp.walk([&](Operation* op) { for (Value result : op->getResults()) { if (shouldRecompute(result, regPressure)) { markForRecomputation(result); } } }); applyRecomputations(funcOp); } bool shouldRecompute(Value v, RegisterPressureAnalysis& rpa) { // Recompute if value is cheap but keeping it live causes spills int computeCost = estimateComputeCost(v.getDefiningOp()); int spillCost = rpa.estimateSpillCost(v); return computeCost < spillCost; } --------------------------------------------------------------------- tileas-legalize-fma-dot Ensures FMA operations match hardware capabilities. void LegalizeFmaDot::runOnOperation() { FuncOp funcOp = getOperation(); funcOp.walk([&](DotOp dotOp) { if (hasFmaAccumulation(dotOp)) { legalizeFma(dotOp); } }); } void legalizeFma(DotOp dotOp) { auto accType = dotOp.getAccumulator().getType(); if (!isLegalAccumulatorType(accType)) { auto legalType = getLegalAccumulatorType(accType); insertAccumulatorConversion(dotOp, legalType); } if (isMixedPrecision(dotOp)) { legalizeMixedPrecisionFma(dotOp); } } --------------------------------------------------------------------- tileas-legalize-reduce Ensures reductions use supported types and sizes. void LegalizeReduce::runOnOperation() { FuncOp funcOp = getOperation(); funcOp.walk([&](ReduceOp reduceOp) { if (!isLegalReduction(reduceOp)) { legalizeReduction(reduceOp); } }); } void legalizeReduction(ReduceOp reduceOp) { auto inputType = reduceOp.getInput().getType(); auto reductionKind = reduceOp.getReductionKind(); if (!isSupportedElementType(inputType.getElementType())) { insertTypeConversion(reduceOp); } if (!isSupportedReductionSize(inputType, reduceOp.getReductionDim())) { splitReduction(reduceOp); } } --------------------------------------------------------------------- tileas-legalize-tmem-copy Legalizes tensor memory (tmem) copy operations. Tensor memory is dedicated storage for tensor core operands. void TileASLegalizeTmemCopy::runOnOperation() { FuncOp funcOp = getOperation(); funcOp.walk([&](Operation* op) { if (auto copyOp = dyn_cast(op)) { if (involvesTmem(copyOp)) { legalizeTmemCopy(copyOp); } } }); } void legalizeTmemCopy(CopyOp copyOp) { auto srcLayout = getLayout(copyOp.getSource()); auto dstLayout = getLayout(copyOp.getDest()); // Infer register layout from tmem layout auto regLayout = inferRegisterLayoutFromTmem(srcLayout); // Insert necessary layout conversions if (needsConversion(srcLayout, regLayout)) { insertLayoutConversion(copyOp, srcLayout, regLayout); } } --------------------------------------------------------------------- tileas-slice-and-fuse Applies loop tiling (slicing) and fusion for improved data locality. void SliceAndFuse::runOnOperation() { FuncOp funcOp = getOperation(); SmallVector fusionGroups; collectFusionCandidates(funcOp, fusionGroups); for (auto& group : fusionGroups) { auto sliceSize = computeOptimalSliceSize(group); sliceOperations(group, sliceSize); fuseOperations(group); } } void fuseOperations(FusionGroup& group) { // Create fused loop nest // - Single loop iterating over slices // - Multiple operations per slice iteration auto fusedLoop = createFusedLoop(group); for (auto* op : group.getOperations()) { moveIntoFusedLoop(op, fusedLoop); } } --------------------------------------------------------------------- tileas-refine-atom-by-resource Adjusts operation granularity ("atom") based on available hardware resources. void RefineAtomByResource::runOnOperation() { FuncOp funcOp = getOperation(); auto resources = getTargetResources(funcOp); funcOp.walk([&](Operation* op) { if (hasAtomAttribute(op)) { refineAtom(op, resources); } }); } void refineAtom(Operation* op, ResourceConstraints& resources) { auto currentAtom = getAtom(op); int smemRequired = estimateSmemUsage(op, currentAtom); int regsRequired = estimateRegUsage(op, currentAtom); // Refine if over resource limits (SM120: 228KB smem, 65536 regs) if (smemRequired > resources.maxSmem || regsRequired > resources.maxRegs) { auto refinedAtom = findSmallerAtom(op, resources); setAtom(op, refinedAtom); } } --------------------------------------------------------------------- tileas-prepare-for-scheduling Normalizes IR and annotates operation latencies for the scheduler. void PrepareForScheduling::runOnOperation() { FuncOp funcOp = getOperation(); normalizeLoops(funcOp); insertSchedulingAnchors(funcOp); annotateLatencies(funcOp); identifyBarriers(funcOp); } void annotateLatencies(FuncOp funcOp) { funcOp.walk([&](Operation* op) { int latency = estimateLatency(op); op->setAttr("sched.latency", builder.getI64IntegerAttr(latency)); }); } --------------------------------------------------------------------- tileas-unroll-register-loops Unrolls loops that access register-resident tensors (required since GPU registers cannot be dynamically indexed). void TileASUnrollRegisterLoops::runOnOperation() { FuncOp funcOp = getOperation(); funcOp.walk([&](scf::ForOp forOp) { if (accessesRegisterTensors(forOp)) { if (!canAvoidUnroll(forOp)) { // Must unroll - register tensors require static indexing unrollLoop(forOp); } } }); } bool accessesRegisterTensors(scf::ForOp forOp) { bool accessesRegs = false; forOp.walk([&](Operation* op) { for (Value operand : op->getOperands()) { if (isRegisterTensor(operand)) { accessesRegs = true; } } }); return accessesRegs; } --------------------------------------------------------------------- tileas-unspecialized-pipeline Implements software pipelining without warp specialization (all warps do both load and compute). void TileASUnspecializedPipeline::runOnOperation() { FuncOp funcOp = getOperation(); int numStages = getOption("unspecialized-pipeline-num-stages"); funcOp.walk([&](scf::ForOp forOp) { if (canPipeline(forOp)) { applySoftwarePipelining(forOp, numStages); } }); } void applySoftwarePipelining(scf::ForOp forOp, int numStages) { emitPrologue(forOp, numStages); // Pre-load data for first N iterations emitSteadyState(forOp, numStages); // Overlap load(i+N) with compute(i) emitEpilogue(forOp, numStages); // Drain remaining computations } --------------------------------------------------------------------- tileas-dynamic-persistent Transforms kernels into dynamic persistent kernels that process work items from a queue. void TileASDynamicPersistent::runOnOperation() { FuncOp funcOp = getOperation(); if (funcOp->hasAttr("dynamic_persistent")) { emitWarning("Kernel is already dynamic persistent"); return; } transformToPersistent(funcOp); funcOp->setAttr("dynamic_persistent", builder.getUnitAttr()); } void transformToPersistent(FuncOp funcOp) { // Insert outer loop that fetches work items: // while (workAvailable()) { // workItem = fetchWork(); // processWorkItem(workItem); // signalCompletion(); // } } --------------------------------------------------------------------- tileas-insert-OCG-knobs Inserts OCG (Optimizing Code Generator) hints for the PTXAS backend. void TileASInsertOCGKnobs::runOnOperation() { FuncOp funcOp = getOperation(); funcOp.walk([&](Operation* op) { if (auto loopOp = dyn_cast(op)) { insertOCGDirectives(loopOp); } if (auto mmaOp = dyn_cast(op)) { insertMMAOptimizationHints(mmaOp); } }); } void insertOCGDirectives(Operation* op) { op->setAttr("ocgEnterDirectives", buildOCGDirectives(op, /*enter=*/true)); op->setAttr("ocgLeaveDirectives", buildOCGDirectives(op, /*enter=*/false)); } --------------------------------------------------------------------- Appendix: IR Dumps This appendix contains the IR dumps from the MoE kernel compilation. Some of the IR below uses %0 placeholders. cuda_tile IR // cuda_tile dialect operations // High-level tensor operations from CuTile Python API // === Pass #1 scope=cuda_tile.entry === "cuda_tile.module"() {1 regions} "cuda_tile.entry"() {1 regions} %0 = "cuda_tile.constant"() : () -> (ct.view) %0 = "cuda_tile.constant"() : () -> (ct.view) %0 = "cuda_tile.constant"() : () -> (ct.view) %0 = "cuda_tile.constant"() : () -> (ct.view) %0 = "cuda_tile.constant"() : () -> (ct.view) %0 = "cuda_tile.constant"() : () -> (ct.view) %0 = "cuda_tile.constant"() : () -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.make_tensor_view"(%cuda_tile.assume, %cuda_tile.assume, %cuda_tile.assume, %cuda_tile.assume, %cuda_tile.assume, %cuda_tile.assume) : (ct.view, ct.view, ct.view, ct.view, ct.view, ct.view) -> (ct.token) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.assume"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.make_tensor_view"(%cuda_tile.assume, %cuda_tile.assume) : (ct.view, ct.view) -> (ct.token) %0 = "cuda_tile.make_token"() : () -> (ct.ptr) %0, %1, %2 = "cuda_tile.get_tile_block_id"() : () -> (ct.view, ct.view, ct.view) %0 = "cuda_tile.divi"(%cuda_tile.assume, %cuda_tile.constant) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.divi"(%cuda_tile.assume, %cuda_tile.constant) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.muli"(%cuda_tile.constant, %cuda_tile.divi) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.divi"(%cuda_tile.get_tile_block_id, %cuda_tile.muli) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.muli"(%cuda_tile.divi, %cuda_tile.constant) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.subi"(%cuda_tile.divi, %cuda_tile.muli) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.mini"(%cuda_tile.subi, %cuda_tile.constant) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.remi"(%cuda_tile.get_tile_block_id, %cuda_tile.mini) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.cmpi"(%cuda_tile.remi, %cuda_tile.constant) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.cmpi"(%cuda_tile.mini, %cuda_tile.constant) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.xori"(%cuda_tile.cmpi, %cuda_tile.cmpi) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.cmpi"(%cuda_tile.remi, %cuda_tile.constant) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.andi"(%cuda_tile.xori, %cuda_tile.cmpi) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.addi"(%cuda_tile.remi, %cuda_tile.mini) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.select"(%cuda_tile.andi, %cuda_tile.addi, %cuda_tile.remi) : (ct.view, ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.addi"(%cuda_tile.muli, %cuda_tile.select) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.remi"(%cuda_tile.get_tile_block_id, %cuda_tile.muli) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.cmpi"(%cuda_tile.remi, %cuda_tile.constant) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.cmpi"(%cuda_tile.muli, %cuda_tile.constant) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.xori"(%cuda_tile.cmpi, %cuda_tile.cmpi) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.cmpi"(%cuda_tile.remi, %cuda_tile.constant) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.andi"(%cuda_tile.xori, %cuda_tile.cmpi) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.addi"(%cuda_tile.remi, %cuda_tile.muli) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.select"(%cuda_tile.andi, %cuda_tile.addi, %cuda_tile.remi) : (ct.view, ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.divi"(%cuda_tile.select, %cuda_tile.mini) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.muli"(%cuda_tile.addi, %cuda_tile.constant) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.iota"() : () -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.muli) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.addi"(%cuda_tile.broadcast, %cuda_tile.iota) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.exti"(%cuda_tile.addi) : (ct.view) -> (ct.view) %0 = "cuda_tile.exti"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.exti) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.cmpi"(%cuda_tile.exti, %cuda_tile.broadcast) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.offset"(%cuda_tile.broadcast, %cuda_tile.exti) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.constant) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0, %1 = "cuda_tile.load_ptr_tko"(%cuda_tile.offset, %cuda_tile.cmpi, %cuda_tile.broadcast, %cuda_tile.make_token) : (ct.view, ct.view, ct.view, ct.ptr) -> (ct.view, ct.ptr) %0 = "cuda_tile.reshape"(%arg) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.divi"(%cuda_tile.load_ptr_tko, %cuda_tile.broadcast) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.make_partition_view"(%cuda_tile.make_tensor_view) : (ct.token) -> (ct.part) %0, %1 = "cuda_tile.load_view_tko"(%cuda_tile.make_partition_view, %cuda_tile.addi, %cuda_tile.make_token) : (ct.part, ct.view, ct.ptr) -> (ct.view, ct.ptr) %0 = "cuda_tile.reshape"(%cuda_tile.load_view_tko) : (ct.view) -> (ct.view) %0 = "cuda_tile.divi"(%cuda_tile.assume, %cuda_tile.constant) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.iota"() : () -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.divi) : (ct.view) -> (ct.view) %0 = "cuda_tile.exti"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.exti) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.exti"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.exti) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.exti"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.exti) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.constant) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.for"(%cuda_tile.constant, %cuda_tile.divi, %cuda_tile.constant, %cuda_tile.constant) {1 regions} : (ct.view, ct.view, ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.muli"(%arg, %cuda_tile.constant) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.muli) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.addi"(%cuda_tile.broadcast, %cuda_tile.iota) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.addi) : (ct.view) -> (ct.view) %0 = "cuda_tile.exti"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.exti) : (ct.view) -> (ct.view) %0 = "cuda_tile.cmpi"(%cuda_tile.broadcast, %cuda_tile.broadcast) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.muli"(%cuda_tile.broadcast, %cuda_tile.broadcast) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.exti"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.exti) : (ct.view) -> (ct.view) %0 = "cuda_tile.cmpi"(%cuda_tile.broadcast, %cuda_tile.broadcast) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.andi"(%cuda_tile.cmpi, %cuda_tile.cmpi) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.addi"(%cuda_tile.muli, %cuda_tile.broadcast) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.offset"(%cuda_tile.broadcast, %cuda_tile.addi) : (ct.view, ct.view) -> (ct.view) %0, %1 = "cuda_tile.load_ptr_tko"(%cuda_tile.offset, %cuda_tile.andi, %cuda_tile.broadcast, %cuda_tile.make_token) : (ct.view, ct.view, ct.view, ct.ptr) -> (ct.view, ct.ptr) %0 = "cuda_tile.make_partition_view"(%cuda_tile.make_tensor_view) : (ct.token) -> (ct.part) %0, %1 = "cuda_tile.load_view_tko"(%cuda_tile.make_partition_view, %cuda_tile.reshape, %arg, %cuda_tile.divi, %cuda_tile.make_token) : (ct.part, ct.view, ct.view, ct.view, ct.ptr) -> (ct.view, ct.ptr) %0 = "cuda_tile.reshape"(%cuda_tile.load_view_tko) : (ct.view) -> (ct.view) %0 = "cuda_tile.mmaf"(%cuda_tile.load_ptr_tko, %cuda_tile.reshape, %arg) : (ct.view, ct.view, ct.view) -> (ct.view) "cuda_tile.continue"(%cuda_tile.mmaf) : (ct.view) -> () %0 = "cuda_tile.muli"(%cuda_tile.divi, %cuda_tile.constant) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.iota"() : () -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.muli) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.addi"(%cuda_tile.broadcast, %cuda_tile.iota) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.ftof"(%cuda_tile.for) : (ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.load_ptr_tko) : (ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.addi) : (ct.view) -> (ct.view) %0 = "cuda_tile.exti"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.exti) : (ct.view) -> (ct.view) %0 = "cuda_tile.exti"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.exti) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.cmpi"(%cuda_tile.broadcast, %cuda_tile.broadcast) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.exti"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.exti) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.muli"(%cuda_tile.broadcast, %cuda_tile.broadcast) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.exti"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.exti) : (ct.view) -> (ct.view) %0 = "cuda_tile.exti"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.exti) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.cmpi"(%cuda_tile.broadcast, %cuda_tile.broadcast) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.andi"(%cuda_tile.cmpi, %cuda_tile.cmpi) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.addi"(%cuda_tile.muli, %cuda_tile.broadcast) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.reshape"(%cuda_tile.assume) : (ct.view) -> (ct.view) %0 = "cuda_tile.broadcast"(%cuda_tile.reshape) : (ct.view) -> (ct.view) %0 = "cuda_tile.offset"(%cuda_tile.broadcast, %cuda_tile.addi) : (ct.view, ct.view) -> (ct.view) %0 = "cuda_tile.store_ptr_tko"(%cuda_tile.offset, %cuda_tile.ftof, %cuda_tile.andi, %cuda_tile.make_token) : (ct.view, ct.view, ct.view, ct.ptr) -> (ct.ptr) "cuda_tile.return"() nv_tileaa IR // nv_tileaa dialect operations // Tile-level ops (architecture-independent) // === Pass #1 scope=nv_tileaa.func === "nv_tileaa.func"() {nv_tileaa.kernel_spec} {1 regions} %0 = "nv_tileaa.assume"(%arg) : (aa.memref) -> (aa.memref) %0 = "nv_tileaa.assume"(%arg) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%nv_tileaa.assume) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%arg) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%nv_tileaa.assume) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%arg) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%nv_tileaa.assume) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%arg) : (aa.memref) -> (aa.memref) %0 = "nv_tileaa.assume"(%arg) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%nv_tileaa.assume) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%arg) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%nv_tileaa.assume) : (iN) -> (iN) %0 = "nv_tileaa.splat"(%nv_tileaa.assume) : (iN) -> (tensor<...>) %0 = "nv_tileaa.assume"(%arg) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%nv_tileaa.assume) : (iN) -> (iN) %0 = "nv_tileaa.splat"(%nv_tileaa.assume) : (iN) -> (tensor<...>) %0 = "nv_tileaa.assume"(%arg) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%nv_tileaa.assume) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%arg) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%nv_tileaa.assume) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%arg) : (aa.memref) -> (aa.memref) %0 = "nv_tileaa.assume"(%arg) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%nv_tileaa.assume) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%arg) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%nv_tileaa.assume) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%arg) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%nv_tileaa.assume) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%arg) : (aa.memref) -> (aa.memref) %0 = "nv_tileaa.assume"(%arg) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%nv_tileaa.assume) : (iN) -> (iN) %0 = "nv_tileaa.splat"(%nv_tileaa.assume) : (iN) -> (tensor<...>) %0 = "nv_tileaa.assume"(%arg) : (aa.memref) -> (aa.memref) %0 = "nv_tileaa.assume"(%arg) : (iN) -> (iN) %0 = "nv_tileaa.assume"(%nv_tileaa.assume) : (iN) -> (iN) %0 = "nv_tileaa.make_memref"(%nv_tileaa.assume, %nv_tileaa.assume, %nv_tileaa.assume, %nv_tileaa.assume, %nv_tileaa.assume, %nv_tileaa.assume) : (aa.memref, iN, iN, iN, iN, iN) -> (aa.btile) %0 = "nv_tileaa.make_memref"(%nv_tileaa.assume, %nv_tileaa.assume) : (aa.memref, iN) -> (aa.btile) %0 = "nv_tileaa.create_mem_token"() : () -> (aa.ptr) %0 = "nv_tileaa.get_program_id"() : () -> (iN) %0 = "nv_tileaa.splat"(%nv_tileaa.get_program_id) : (iN) -> (tensor<...>) %0 = "nv_tileaa.make_range"(%arith.constant, %arith.constant) : (iN, iN) -> (tensor<...>) %0 = "nv_tileaa.extract"(%arith.muli) : (tensor<...>) -> (iN) %0 = "nv_tileaa.splat"(%nv_tileaa.extract) : (iN) -> (tensor<...>) %0 = "nv_tileaa.extract"(%arith.extsi) : (tensor<...>) -> (iN) %0 = "nv_tileaa.splat"(%nv_tileaa.extract) : (iN) -> (tensor<...>) %0 = "nv_tileaa.splat"(%nv_tileaa.assume) : (aa.memref) -> (tensor<...>) %0 = "nv_tileaa.addptr"(%nv_tileaa.splat, %arith.extsi) : (tensor<...>, tensor<...>) -> (tensor<...>) %0, %1 = "nv_tileaa.load"(%nv_tileaa.addptr, %arith.cmpi, %arith.constant, %nv_tileaa.create_mem_token) : (tensor<...>, tensor<...>, tensor<...>, aa.ptr) -> (tensor<...>, aa.ptr) %0 = "nv_tileaa.splat"(%arg) : (iN) -> (tensor<...>) %0 = "nv_tileaa.block_tile"(%nv_tileaa.make_memref) : (aa.btile) -> (aa.mtoken) %0 = "nv_tileaa.extract"(%arith.addi) : (tensor<...>) -> (iN) %0, %1 = "nv_tileaa.tiled_load"(%nv_tileaa.block_tile, %nv_tileaa.extract, %nv_tileaa.create_mem_token) : (aa.mtoken, iN, aa.ptr) -> (tensor<...>, aa.ptr) %0 = "nv_tileaa.view"(%nv_tileaa.tiled_load) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileaa.make_range"(%arith.constant, %arith.constant) : (iN, iN) -> (tensor<...>) %0 = "nv_tileaa.expand_dims"(%arith.floordivsi) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileaa.splat"(%arith.extsi) : (iN) -> (tensor<...>) %0 = "nv_tileaa.splat"(%arith.extsi) : (iN) -> (tensor<...>) %0 = "nv_tileaa.splat"(%arith.extsi) : (iN) -> (tensor<...>) %0 = "nv_tileaa.splat"(%nv_tileaa.assume) : (aa.memref) -> (tensor<...>) %0 = "nv_tileaa.extract"(%arith.ceildivsi) : (tensor<...>) -> (iN) %0 = "nv_tileaa.splat"(%arg) : (iN) -> (tensor<...>) %0 = "nv_tileaa.extract"(%arith.muli) : (tensor<...>) -> (iN) %0 = "nv_tileaa.splat"(%nv_tileaa.extract) : (iN) -> (tensor<...>) %0 = "nv_tileaa.expand_dims"(%arith.addi) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileaa.broadcast"(%arith.extsi) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileaa.broadcast"(%arith.extsi) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileaa.addptr"(%nv_tileaa.splat, %arith.addi) : (tensor<...>, tensor<...>) -> (tensor<...>) %0, %1 = "nv_tileaa.load"(%nv_tileaa.addptr, %arith.andi, %arith.constant, %nv_tileaa.create_mem_token) : (tensor<...>, tensor<...>, tensor<...>, aa.ptr) -> (tensor<...>, aa.ptr) %0 = "nv_tileaa.block_tile"(%nv_tileaa.make_memref) : (aa.btile) -> (aa.mtoken) %0 = "nv_tileaa.extract"(%nv_tileas.convert_layout) : (tensor<...>) -> (iN) %0 = "nv_tileaa.extract"(%arith.floordivsi) : (tensor<...>) -> (iN) %0, %1 = "nv_tileaa.tiled_load"(%nv_tileaa.block_tile, %nv_tileaa.extract, %arg, %nv_tileaa.extract, %nv_tileaa.create_mem_token) : (aa.mtoken, iN, iN, iN, aa.ptr) -> (tensor<...>, aa.ptr) %0 = "nv_tileaa.view"(%nv_tileaa.tiled_load) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileaa.dot"(%nv_tileaa.load, %nv_tileaa.view, %arg) : (tensor<...>, tensor<...>, tensor<...>) -> (tensor<...>) %0 = "nv_tileaa.make_range"(%arith.constant, %arith.constant) : (iN, iN) -> (tensor<...>) %0 = "nv_tileaa.extract"(%arith.muli) : (tensor<...>) -> (iN) %0 = "nv_tileaa.splat"(%nv_tileaa.extract) : (iN) -> (tensor<...>) %0 = "nv_tileaa.fp_to_fp"(%scf.for) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileaa.expand_dims"(%nv_tileaa.load) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileaa.expand_dims"(%arith.addi) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileaa.broadcast"(%arith.extsi) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileaa.splat"(%arith.extsi) : (iN) -> (tensor<...>) %0 = "nv_tileaa.splat"(%arith.extsi) : (iN) -> (tensor<...>) %0 = "nv_tileaa.broadcast"(%arith.extsi) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileaa.splat"(%arith.extsi) : (iN) -> (tensor<...>) %0 = "nv_tileaa.splat"(%nv_tileaa.assume) : (aa.memref) -> (tensor<...>) %0 = "nv_tileaa.addptr"(%nv_tileaa.splat, %arith.addi) : (tensor<...>, tensor<...>) -> (tensor<...>) %0 = "nv_tileaa.store"(%nv_tileaa.addptr, %nv_tileaa.fp_to_fp, %arith.andi, %nv_tileaa.create_mem_token) : (tensor<...>, tensor<...>, tensor<...>, aa.ptr) -> (aa.ptr) "nv_tileaa.return"() // === Pass #2 scope=nv_tileaa.func === // === Pass #3 scope=nv_tileaa.func === // === Pass #4 scope=nv_tileaa.func === // === Pass #5 scope=nv_tileaa.func === // === Pass #6 scope=nv_tileaa.func === // === Pass #7 scope=nv_tileaa.func === // === Pass #8 scope=nv_tileaa.func === // === Pass #9 scope=nv_tileaa.func === // === Pass #10 scope=nv_tileaa.func === // === Pass #11 scope=nv_tileaa.func === // === Pass #12 scope=nv_tileaa.func === // (Lines 193-352 - final assembly with fp_to_fp conversions) // See dump for complete content including: // - 32 fp_to_fp operations for output precision conversion // - Multiple nv_tileaa.func declarations with kernel metadata // - Final memory layout preparation nv_tileas IR // nv_tileas dialect operations // Tile-level Scheduled Assembly (architecture-specific) // [within nv_tileaa.func pass] %0, %1 = "nv_tileas.load"(%nv_tileaa.addptr, %arith.cmpi, %arith.constant, %nv_tileaa.create_mem_token) : (tensor<...>, tensor<...>, tensor<...>, aa.ptr) -> (tensor<...>, aa.ptr) %0, %1 = "nv_tileas.tiled_load"(%nv_tileaa.block_tile, %nv_tileaa.extract, %nv_tileaa.create_mem_token) : (aa.mtoken, iN, aa.ptr) -> (tensor<...>, aa.ptr) %0 = "nv_tileas.view"(%nv_tileas.tiled_load) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileas.expand_dims"(%arith.floordivsi) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileas.expand_dims"(%arith.addi) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileas.convert_layout"(%nv_tileaa.broadcast) : (tensor<...>) -> (tensor<...>) %0, %1 = "nv_tileas.load"(%nv_tileaa.addptr, %arith.andi, %arith.constant, %nv_tileaa.create_mem_token) : (tensor<...>, tensor<...>, tensor<...>, aa.ptr) -> (tensor<...>, aa.ptr) %0 = "nv_tileas.convert_layout"(%nv_tileas.view) : (tensor<...>) -> (tensor<...>) %0, %1 = "nv_tileas.tiled_load"(%nv_tileaa.block_tile, %nv_tileaa.extract, %arg, %nv_tileaa.extract, %nv_tileaa.create_mem_token) : (aa.mtoken, iN, iN, iN, aa.ptr) -> (tensor<...>, aa.ptr) %0 = "nv_tileas.view"(%nv_tileas.tiled_load) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileas.convert_layout"(%nv_tileas.load) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileas.convert_layout"(%nv_tileas.view) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileas.convert_layout"(%arg) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileas.dot"(%nv_tileas.convert_layout, %nv_tileas.convert_layout, %nv_tileas.convert_layout, %arith.constant) : (tensor<...>, tensor<...>, tensor<...>, iN) -> (tensor<...>) %0 = "nv_tileas.convert_layout"(%nv_tileas.dot) : (tensor<...>) -> (tensor<...>) %0 = "nv_tileas.make_tiled_tma_desc"(%nv_tileaa.make_memref) {tmaIdx} : (aa.btile) -> (?type) // [within builtin.module pass] %0 = "nv_tileas.async.pipeline.create_pipeline"() : () -> (?type) %0 = "nv_tileas.async.pipeline.create_pipeline"() : () -> (?type) %0 = "nv_tileas.async.pipeline.create_pipeline"() : () -> (?type) %0 = "nv_tileas.async.pipeline.create_iterator"(%nv_tileas.async.pipeline.create_pipeline) : (?type) -> (?type) %0 = "nv_tileas.async.pipeline.create_iterator"(%nv_tileas.async.pipeline.create_pipeline) : (?type) -> (?type) %0 = "nv_tileas.async.pipeline.create_iterator"(%nv_tileas.async.pipeline.create_pipeline) : (?type) -> (?type) %0 = "nv_tileas.async.pipeline.create_iterator"(%nv_tileas.async.pipeline.create_pipeline) : (?type) -> (?type) %0 = "nv_tileas.async.pipeline.create_iterator"(%nv_tileas.async.pipeline.create_pipeline) : (?type) -> (?type) %0 = "nv_tileas.async.pipeline.create_iterator"(%nv_tileas.async.pipeline.create_pipeline) : (?type) -> (?type) "nv_tileas.async.pipeline.agent_switch"(%arg, ...) {4 regions} : (...) -> () // Producer-Consumer Pattern (repeated throughout) %0 = "nv_tileas.async.pipeline.producer_acquire"(%arg, %arg) : (?type, ?type) -> (?type) %0 = "nv_tileas.async.pipeline.inc_iter"(%arg) : (?type) -> (?type) %0 = "nv_tileas.async.pipeline.producer_write"(%arg, %nv_tileas.async.pipeline.producer_acquire) {1 regions} : (?type, ?type) -> (?type) "nv_tileas.async.pipeline.producer_commit"(%arg, %nv_tileas.async.pipeline.producer_write) : (?type, ?type) -> () %0 = "nv_tileas.async.pipeline.consumer_wait"(%arg, %arg) : (?type, ?type) -> (?type) %0, %1 = "nv_tileas.async.pipeline.consumer_read"(%arg, %nv_tileas.async.pipeline.consumer_wait) {consumer_idx} {1 regions} : (?type, ?type) -> (?type, tensor<...>) "nv_tileas.async.pipeline.consumer_release"(%arg, %nv_tileas.async.pipeline.consumer_read) : (?type, ?type) -> () // Dot operations (100+ for tiled matrix multiply) %0 = "nv_tileas.dot"(%nv_tileas.extract_slice, %nv_tileas.extract_slice, %arg, %arith.constant) : (tensor<...>, tensor<...>, tensor<...>, iN) -> (tensor<...>) // ... (repeated for all tile partitions) // TMA operations %0 = "nv_tileas.make_tiled_tma_desc"(%nv_tileaa.make_memref) {tmaIdx} : (aa.btile) -> (?type) %0 = "nv_tileas.async.tiled_tma_load"(%nv_tileaa.block_tile, %arg, %nv_tileas.make_tiled_tma_desc, %nv_tileaa.extract, %arg, %nv_tileaa.extract) : (...) -> (?type) // Output assembly (32 insert_slice for output tiles) %0 = "nv_tileas.insert_slice"(%nv_tileaa.fp_to_fp, %nv_tileas.alloc_tensor, %arith.constant, %arith.constant) : (tensor<...>, tensor<...>, iN, iN) -> (tensor<...>) // ... (repeated 32 times) NVVM Dialect IR // nvvm dialect operations // NVVM (NVIDIA PTX intrinsics in MLIR form) // === Barrier and Fence Operations === "nvvm.fence.mbarrier.init"() "nvvm.barrier"() "nvvm.fence.proxy"() %0 = "nvvm.read.ptx.sreg.clusterid.x"() : () -> (i32) %0 = "nvvm.read.ptx.sreg.tid.x"() : () -> (i32) // === Async Global-Shared Copies (136 instances) === "nvvm.cp.async.shared.global"(%ptr, %src, %predicate) : (ptr<3>, ptr<1>, i1) -> () // === Tensor Core Data Packing (1,088 instances) === %0 = "nvvm.cvt.packfloat.f32"(%a, %b, %mode) : (f32, f32, i32) -> (i32) // === Memory Barriers (66 instances) === "nvvm.mbarrier.init.shared"(%barrier, %count) : (ptr<3>, i32) -> () "nvvm.mbarrier.arrive.shared"(%barrier) : (ptr<3>) -> () "nvvm.mbarrier.wait.shared"(%barrier, %phase) : (ptr<3>, i32) -> () // === Matrix Load Operations (512 instances) === %0 = "nvvm.ldmatrix"(%ptr) {layout = #nvvm.mma_layout, num = 4} : (ptr<3>) -> vector<4xi32> // === Tensor Core MMA (512 instances) === %0 = "nvvm.mma.sync"(%a, %b, %c) { layoutA = #nvvm.mma_layout, layoutB = #nvvm.mma_layout, shape = #nvvm.shape } : (vector<4xi32>, vector<2xi32>, vector<4xf32>) -> vector<4xf32> // ... (2,977 lines total - tensor core operations, barriers, memory ops) LLVM IR / NVVM IR ; ModuleID = 'LLVMDialectModule' target datalayout = "e-p:64:64:64-p3:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64" target triple = "nvptx64-nvidia-cuda" ; Kernel entry point with TMA descriptors define ptx_kernel void @fused_moe_kernel( ptr addrspace(1) %A, ; Input tokens ptr addrspace(1) %B, ; Expert weights ptr addrspace(1) %C, ; Output ptr addrspace(1) %topk_weights, ptr addrspace(1) %sorted_token_ids, ptr addrspace(1) %sorted_expert_ids, i32 %num_token_replicas, i1 %mul_routed_weight, ; ... TMA descriptors appended by tileas-attach-tma-desc-args ) #0 { entry: ; Get cluster/block/thread IDs %clusterid = call i32 @llvm.nvvm.read.ptx.sreg.clusterid.x() %tid = call range(i32 0, 384) i32 @llvm.nvvm.read.ptx.sreg.tid.x() ; Initialize barriers for async pipeline call void @llvm.nvvm.mbarrier.init.shared(ptr addrspace(3) %barrier, i32 128) ; Async copy from global to shared memory call void @llvm.nvvm.cp.async.shared.global( ptr addrspace(3) %shared_dst, ptr addrspace(1) %global_src, i32 16, ; bytes i1 %pred ; predicate ) ; Tensor core matrix multiply %result = call <4 x float> @llvm.nvvm.mma.m16n8k16.row.col.f32.f16.f16.f32( <4 x i32> %a_frag, <2 x i32> %b_frag, <4 x float> %c_frag ) ; ... (full pipeline with producer/consumer synchronization) } ; NVVM intrinsic declarations declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() declare i32 @llvm.nvvm.read.ptx.sreg.clusterid.x() declare void @llvm.nvvm.mbarrier.init.shared(ptr addrspace(3), i32) declare void @llvm.nvvm.cp.async.shared.global(ptr addrspace(3), ptr addrspace(1), i32, i1) declare <4 x float> @llvm.nvvm.mma.m16n8k16.row.col.f32.f16.f16.f32(<4 x i32>, <2 x i32>, <4 x float>) PTX Assembly // // Generated by NVIDIA NVVM Compiler // Cuda compilation tools, release 13.1, V13.1.80 // Based on NVVM 21.0.0 // .version 9.1 .target sm_120a .address_size 64 .visible .entry fused_moe_kernel( .param .u64 .ptr .global .align 1 fused_moe_kernel_param_0, .param .u32 fused_moe_kernel_param_1, // ... 31 parameters total including TMA descriptors .hidden .param .align 64 .b8 fused_moe_kernel_param_31[128] ) .reqntid 384 .minnctapersm 1 { .reg .pred %p<306>; .reg .b16 %rs<500>; .reg .b32 %r<4905>; .reg .b64 %rd<348>; // 80KB shared memory for double buffering .shared .align 128 .b8 global_smem[82032]; // === Barrier Initialization === mbarrier.init.shared.b64 [global_smem+82000], %r2369; mbarrier.init.shared.b64 [global_smem+82008], %r2369; // === Matrix Load (ldmatrix for tensor cores) === ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r4645, %r4646, %r4647, %r4648}, [%r2789]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r4649, %r4650, %r4651, %r4652}, [%r2793]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r4653, %r4654, %r4655, %r4656}, [%r2797]; ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r4657, %r4658, %r4659, %r4660}, [%r2801]; // ... (512 ldmatrix instructions total) // === Tensor Core MMA (HMMA) === // Note: sm_120a uses wgmma/tcgen05 instructions in SASS // PTX shows the portable mma.sync form mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 {%f1, %f2, %f3, %f4}, {%r4645, %r4646, %r4647, %r4648}, {%r4709, %r4710}, {%f1, %f2, %f3, %f4}; // ... (512 mma.sync instructions total) // === Async Copy (cp.async for global-shared) === cp.async.cg.shared.global [%r2856], [%rd112], 16, %p116; cp.async.cg.shared.global [%r2857], [%rd113], 16, %p116; // ... (136 cp.async instructions total) // === Barrier Synchronization === mbarrier.arrive.shared.b64 _, [global_smem+82000]; mbarrier.try_wait.parity.shared.b64 %p117, [global_smem+82000], %r2371; } Citation To cite this article: @article{zhu2026tileir, title = {NVIDIA TileIR Internals: from CuTile to MLIR/LLVM to SASS}, author = {Zhu, Henry}, journal = {maknee.github.io}, year = {2026}, month = {January}, url = "https://maknee.github.io/blog/2026/NVIDIA-TileIR-Internals-from-CuTile-to-MLIR-LLVM-to-SASS/" } --------------------------------------------------------------------- Stay Updated If you enjoyed this post, subscribe to the RSS feed, follow me on X (Twitter), or get email updates. --------------------------------------------------------------------- Enjoy Reading This Article? Here are some more articles you might like to read next: Performance Hints Maybe consider putting "cutlass" in your CUDA/Triton kernels Network Storage and Scaling Characteristics of a Distributed Filesystem Network and Storage Benchmarks for LLM Training on the Cloud AI 2027 Please enable JavaScript to view the comments powered by giscus. (c) Copyright 2026 Henry Zhu. # Powered by Jekyll with al-folio theme. # Hosted by GitHub Pages. Last updated: January 31, 2026.