Skip to main content
Ctrl+K
Modern GPU Programming For MLSys - Home Modern GPU Programming For MLSys - Home

Part I, Understanding the GPU

  • GPU Execution Model
  • What Makes a Kernel Fast
  • Data Layout and Its Notation
  • The Evolution of Tensor Core Data Layouts
  • Async Data Movement: TMA
  • Tensor Cores: tcgen05
  • Special Memory: TMEM
  • Async Coordination: mbarriers
  • Advanced: Cluster Launch Control

Part II, TIRx Overview

  • Introduction to TIRx
  • TIRx Layout API

Part III, GEMM: Tiled to SOTA

  • Building a Tiled GEMM
  • Pipelining GEMM with TMA
  • Scaling GEMM with Warp Specialization and Clusters

Part IV, Flash Attention 4

  • Flash Attention 4

Reference

  • Reference
  • Debugging Warp-Specialized Kernels
  • Compiler Internals
    • TIRx lowering pipeline
  • TIRx Language Reference
    • Parser utilities
    • Data types and expressions
    • Buffers and memory
    • Control flow
    • CUDA C++/PTX intrinsics
  • Repository
  • Show source
  • Suggest edit
  • Open issue

Compiler Internals

Compiler Internals#

Internals of the TIRx compiler, for contributors.

  • TIRx lowering pipeline

previous

Debugging Warp-Specialized Kernels

next

TIRx lowering pipeline

By MLC Community

© Copyright 2026, MLC Community.