Contents Menu Expand Light mode Dark mode Auto light/dark, in light mode Auto light/dark, in dark mode Skip to content
TileLang
0.1.13 documentation
Light Logo Dark Logo
TileLang
0.1.13 documentation

GET STARTED

  • Installation Guide
  • The Tile Language: A Brief Introduction
  • Understanding Targets

TUTORIALS

  • Debugging Tile Language Programs
  • Auto-Tuning Techniques for Performance Optimization
  • Logging in Tilelang/TVM

TOOLS

  • Tools
  • Compile-Only Tool
  • Performance Analyzer
  • Layout Visualization
  • AutoDD
  • IR Lower Trace
  • Pass Diff
  • IKET Profiling

PROGRAMMING GUIDES

  • Programming Guides Overview
  • Language Basics
  • Instructions
  • Control Flow
  • Software Pipeline Annotations
  • Python Compatibility
  • Autotuning
  • Type System

DEEP LEARNING OPERATORS

  • ElementWise Operators
  • General Matrix-Vector Multiplication (GEMV)
  • General Matrix-Matrix Multiplication with Tile Library
  • Sparse Matrix-Matrix Multiplication with Tile Library
  • 🚀 Write High Performance FlashMLA with TileLang on Hopper

COMPILER INTERNALS

  • LetStmt Inlining in TileLang
  • InjectFenceProxy Pass
  • Tensor Checks (Host-Side Auto-Validation)
  • TileLang Metal Backend Internals

DEVELOPER GUIDE

  • C++ Style Guide

API Reference

  • tilelang
    • tilelang.analysis
      • tilelang.analysis.ast_printer
      • tilelang.analysis.fragment_loop_checker
      • tilelang.analysis.layout_visual
      • tilelang.analysis.nested_loop_checker
      • tilelang.analysis.parallel_local_index_checker
    • tilelang.autodd
    • tilelang.autotuner
      • tilelang.autotuner.capture
      • tilelang.autotuner.grouped_compile
      • tilelang.autotuner.param
      • tilelang.autotuner.tuner
    • tilelang.backend
      • tilelang.backend.device_codegen
      • tilelang.backend.execution_backend
      • tilelang.backend.host_codegen
      • tilelang.backend.module
      • tilelang.backend.pass_pipeline
        • tilelang.backend.pass_pipeline.pipeline
        • tilelang.backend.pass_pipeline.pipeline_utils
      • tilelang.backend.target
    • tilelang.cache
      • tilelang.cache.cuda_binary_cache
      • tilelang.cache.kernel_cache
    • tilelang.carver
      • tilelang.carver.analysis
      • tilelang.carver.arch
        • tilelang.carver.arch.arch_base
        • tilelang.carver.arch.cdna
        • tilelang.carver.arch.cpu
        • tilelang.carver.arch.cuda
        • tilelang.carver.arch.driver
          • tilelang.carver.arch.driver.cuda_driver
        • tilelang.carver.arch.metal
        • tilelang.carver.arch.rdna
      • tilelang.carver.common_schedules
      • tilelang.carver.matmul_analysis
      • tilelang.carver.roller
        • tilelang.carver.roller.bestfit
        • tilelang.carver.roller.hint
        • tilelang.carver.roller.node
        • tilelang.carver.roller.policy
          • tilelang.carver.roller.policy.common
          • tilelang.carver.roller.policy.default
          • tilelang.carver.roller.policy.tensorcore
        • tilelang.carver.roller.rasterization
        • tilelang.carver.roller.shape_inference
          • tilelang.carver.roller.shape_inference.tir
      • tilelang.carver.template
        • tilelang.carver.template.base
        • tilelang.carver.template.conv
        • tilelang.carver.template.elementwise
        • tilelang.carver.template.flashattention
        • tilelang.carver.template.gemv
        • tilelang.carver.template.general_reduce
        • tilelang.carver.template.matmul
      • tilelang.carver.utils
    • tilelang.contrib
      • tilelang.contrib.cc
      • tilelang.contrib.cutedsl
        • tilelang.contrib.cutedsl.atomic
        • tilelang.contrib.cutedsl.cluster
        • tilelang.contrib.cutedsl.cpasync
        • tilelang.contrib.cutedsl.gemm_tcgen05
        • tilelang.contrib.cutedsl.gemm_v1
        • tilelang.contrib.cutedsl.gemm_v2
        • tilelang.contrib.cutedsl.grid_sync
        • tilelang.contrib.cutedsl.ieee_math
        • tilelang.contrib.cutedsl.ldsm
        • tilelang.contrib.cutedsl.math
        • tilelang.contrib.cutedsl.ptx_mma
        • tilelang.contrib.cutedsl.quantize
        • tilelang.contrib.cutedsl.reduce
        • tilelang.contrib.cutedsl.threadblock_swizzle
        • tilelang.contrib.cutedsl.utils
        • tilelang.contrib.cutedsl.warp
      • tilelang.contrib.hip_resource_info
      • tilelang.contrib.hipcc
      • tilelang.contrib.msvc
      • tilelang.contrib.nvcc
      • tilelang.contrib.nvrtc
      • tilelang.contrib.rocm
    • tilelang.cpu
      • tilelang.cpu.backend
      • tilelang.cpu.codegen
      • tilelang.cpu.execution_backend
      • tilelang.cpu.language
      • tilelang.cpu.op
        • tilelang.cpu.op.gemm
          • tilelang.cpu.op.gemm.gemm_scalar
      • tilelang.cpu.pipeline
      • tilelang.cpu.transform
    • tilelang.cuda
      • tilelang.cuda.backend
      • tilelang.cuda.codegen
      • tilelang.cuda.cutedsl_backend
      • tilelang.cuda.debug
      • tilelang.cuda.execution_backend
      • tilelang.cuda.intrinsics
        • tilelang.cuda.intrinsics.layout
          • tilelang.cuda.intrinsics.layout.mma_layout
          • tilelang.cuda.intrinsics.layout.mma_sm100_layout
          • tilelang.cuda.intrinsics.layout.mma_sm70_layout
          • tilelang.cuda.intrinsics.layout.mma_sp_layout
          • tilelang.cuda.intrinsics.layout.utils
        • tilelang.cuda.intrinsics.macro
          • tilelang.cuda.intrinsics.macro.mma_macro_generator
          • tilelang.cuda.intrinsics.macro.mma_sm120_macro_generator
          • tilelang.cuda.intrinsics.macro.mma_sm70_macro_generator
          • tilelang.cuda.intrinsics.macro.mma_sm75_macro_generator
          • tilelang.cuda.intrinsics.macro.mma_sp_macro_generator
          • tilelang.cuda.intrinsics.macro.tcgen05_macro_generator
          • tilelang.cuda.intrinsics.macro.wgmma_macro_generator
          • tilelang.cuda.intrinsics.macro.wgmma_sp_macro_generator
        • tilelang.cuda.intrinsics.sparse_layout
      • tilelang.cuda.language
        • tilelang.cuda.language.cluster
        • tilelang.cuda.language.intrinsics
        • tilelang.cuda.language.pdl
        • tilelang.cuda.language.print
        • tilelang.cuda.language.random
        • tilelang.cuda.language.tir
        • tilelang.cuda.language.warpgroup
      • tilelang.cuda.op
        • tilelang.cuda.op.gemm
          • tilelang.cuda.op.gemm.gemm_fma
          • tilelang.cuda.op.gemm.gemm_mma
          • tilelang.cuda.op.gemm.gemm_mma_sm120
          • tilelang.cuda.op.gemm.gemm_mma_sm70
          • tilelang.cuda.op.gemm.gemm_mma_sm75
          • tilelang.cuda.op.gemm.gemm_tcgen05
          • tilelang.cuda.op.gemm.gemm_wgmma
        • tilelang.cuda.op.gemm_sp
          • tilelang.cuda.op.gemm_sp.gemm_sp_mma
          • tilelang.cuda.op.gemm_sp.gemm_sp_wgmma
      • tilelang.cuda.pipeline
      • tilelang.cuda.target
      • tilelang.cuda.transform
    • tilelang.dtypes
    • tilelang.engine
      • tilelang.engine.callback
      • tilelang.engine.lower
      • tilelang.engine.param
      • tilelang.engine.semantic_check
    • tilelang.env
    • tilelang.errors
    • tilelang.instrumentation
      • tilelang.instrumentation.events
      • tilelang.instrumentation.session
    • tilelang.ir
    • tilelang.jit
      • tilelang.jit.abi
      • tilelang.jit.adapter
        • tilelang.jit.adapter.base
        • tilelang.jit.adapter.cutedsl
          • tilelang.jit.adapter.cutedsl.adapter
          • tilelang.jit.adapter.cutedsl.checks
          • tilelang.jit.adapter.cutedsl.kernel_cache
          • tilelang.jit.adapter.cutedsl.libgen
          • tilelang.jit.adapter.cutedsl.wrapper
        • tilelang.jit.adapter.cython
          • tilelang.jit.adapter.cython.adapter
          • tilelang.jit.adapter.cython.kernel_cache
        • tilelang.jit.adapter.kernel_cache
        • tilelang.jit.adapter.libgen
        • tilelang.jit.adapter.nvrtc
          • tilelang.jit.adapter.nvrtc.adapter
          • tilelang.jit.adapter.nvrtc.include_paths
          • tilelang.jit.adapter.nvrtc.kernel_cache
          • tilelang.jit.adapter.nvrtc.libgen
          • tilelang.jit.adapter.nvrtc.wrapper
        • tilelang.jit.adapter.torch
          • tilelang.jit.adapter.torch.kernel_cache
          • tilelang.jit.adapter.torch.metal
        • tilelang.jit.adapter.tvm_ffi
        • tilelang.jit.adapter.utils
        • tilelang.jit.adapter.wrapper
      • tilelang.jit.diagnostics
      • tilelang.jit.exceptions
      • tilelang.jit.kernel
      • tilelang.jit.param
    • tilelang.language
      • tilelang.language.allocate
      • tilelang.language.annotations
      • tilelang.language.atomic
      • tilelang.language.builtin
      • tilelang.language.common
      • tilelang.language.copy_op
      • tilelang.language.customize
      • tilelang.language.dtypes
      • tilelang.language.eager
        • tilelang.language.eager.ast
        • tilelang.language.eager.builder
        • tilelang.language.eager.utils
      • tilelang.language.experimental
        • tilelang.language.experimental.gemm_sp_op
      • tilelang.language.fill_op
      • tilelang.language.fp8
      • tilelang.language.frame
      • tilelang.language.gemm_op
      • tilelang.language.kernel
      • tilelang.language.logical
      • tilelang.language.loop
      • tilelang.language.math_intrinsics
      • tilelang.language.meta
      • tilelang.language.overrides
        • tilelang.language.overrides.buffer
      • tilelang.language.proxy
      • tilelang.language.reduce_op
      • tilelang.language.scan_op
      • tilelang.language.symbolics
      • tilelang.language.tile_schedule
      • tilelang.language.tir
        • tilelang.language.tir.common
        • tilelang.language.tir.entry
        • tilelang.language.tir.exports
        • tilelang.language.tir.ir
        • tilelang.language.tir.op
      • tilelang.language.utils
      • tilelang.language.ws_schedule
    • tilelang.layout
      • tilelang.layout.cute
      • tilelang.layout.fragment
      • tilelang.layout.gemm_sp
      • tilelang.layout.layout
      • tilelang.layout.swizzle
      • tilelang.layout.swizzle_mode
    • tilelang.math
    • tilelang.metal
      • tilelang.metal.backend
      • tilelang.metal.codegen
      • tilelang.metal.execution_backend
      • tilelang.metal.intrinsics
        • tilelang.metal.intrinsics.metal_macro_generator
      • tilelang.metal.language
        • tilelang.metal.language.tir
      • tilelang.metal.op
        • tilelang.metal.op.gemm
          • tilelang.metal.op.gemm.gemm_metal
      • tilelang.metal.pipeline
      • tilelang.metal.target
      • tilelang.metal.transform
        • tilelang.metal.transform.mark_host_metal_context
        • tilelang.metal.transform.metal_fragment_to_simdgroup
      • tilelang.metal.utils
    • tilelang.profiler
      • tilelang.profiler.bench
    • tilelang.rocm
      • tilelang.rocm.backend
      • tilelang.rocm.codegen
      • tilelang.rocm.execution_backend
      • tilelang.rocm.intrinsics
        • tilelang.rocm.intrinsics.mfma_layout
        • tilelang.rocm.intrinsics.mfma_macro_generator
        • tilelang.rocm.intrinsics.utils
        • tilelang.rocm.intrinsics.wmma_layout
        • tilelang.rocm.intrinsics.wmma_macro_generator
      • tilelang.rocm.language
        • tilelang.rocm.language.intrinsics
        • tilelang.rocm.language.tir
      • tilelang.rocm.op
        • tilelang.rocm.op.gemm
          • tilelang.rocm.op.gemm.gemm_mfma
          • tilelang.rocm.op.gemm.gemm_wmma
      • tilelang.rocm.pipeline
      • tilelang.rocm.target
    • tilelang.testing
      • tilelang.testing.perf_regression
    • tilelang.tileop
      • tilelang.tileop.base
      • tilelang.tileop.gemm
        • tilelang.tileop.gemm.gemm_base
        • tilelang.tileop.gemm.registry
      • tilelang.tileop.gemm_sp
        • tilelang.tileop.gemm_sp.gemm_sp_base
        • tilelang.tileop.gemm_sp.registry
    • tilelang.tools
      • tilelang.tools.Analyzer
      • tilelang.tools.compile_only
      • tilelang.tools.cuda
        • tilelang.tools.cuda.iket
          • tilelang.tools.cuda.iket.cli
          • tilelang.tools.cuda.iket.codegen
          • tilelang.tools.cuda.iket.frontend
          • tilelang.tools.cuda.iket.metadata
          • tilelang.tools.cuda.iket.session
      • tilelang.tools.lower_trace
        • tilelang.tools.lower_trace.core
        • tilelang.tools.lower_trace.diff
        • tilelang.tools.lower_trace.html
      • tilelang.tools.pass_timing
      • tilelang.tools.pass_visualizer
        • tilelang.tools.pass_visualizer.core
        • tilelang.tools.pass_visualizer.viewer
      • tilelang.tools.plot_layout
    • tilelang.transform
      • tilelang.transform.add_bufstore_wrapper
      • tilelang.transform.decouple_type_cast
      • tilelang.transform.hoist_broadcast_values
      • tilelang.transform.pass_config
      • tilelang.transform.simplify
    • tilelang.utils
      • tilelang.utils.deprecated
      • tilelang.utils.device
      • tilelang.utils.language
      • tilelang.utils.pass_diff
      • tilelang.utils.pass_diff_hook
      • tilelang.utils.tensor
    • tilelang.webgpu
      • tilelang.webgpu.backend
      • tilelang.webgpu.codegen
      • tilelang.webgpu.language
      • tilelang.webgpu.pipeline

Privacy

  • Privacy
Back to top
View this page

tilelang.metal¶

Submodules¶

  • tilelang.metal.backend
  • tilelang.metal.codegen
  • tilelang.metal.execution_backend
  • tilelang.metal.intrinsics
  • tilelang.metal.language
  • tilelang.metal.op
  • tilelang.metal.pipeline
  • tilelang.metal.target
  • tilelang.metal.transform
  • tilelang.metal.utils
Next
tilelang.metal.backend
Previous
tilelang.math
Copyright © 2025-2025, Tile Lang Contributors
Made with Sphinx and @pradyunsg's Furo
On this page
  • tilelang.metal
    • Submodules