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.12 documentation
Light Logo Dark Logo
TileLang
0.1.12 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

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)

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.autodd
    • tilelang.autotuner
      • tilelang.autotuner.capture
      • tilelang.autotuner.grouped_compile
      • tilelang.autotuner.param
      • tilelang.autotuner.tuner
    • tilelang.backend
      • tilelang.backend.common
      • tilelang.backend.device_codegen
      • tilelang.backend.execution_backend
      • tilelang.backend.host_codegen
      • 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.common
          • 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.common
      • tilelang.common.transform_kind
    • tilelang.contrib
      • tilelang.contrib.cc
      • tilelang.contrib.cutedsl
        • tilelang.contrib.cutedsl.atomic
        • 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.dlpack
      • tilelang.contrib.hip_resource_info
      • tilelang.contrib.hipcc
      • tilelang.contrib.msvc
      • tilelang.contrib.nvcc
      • tilelang.contrib.nvrtc
      • tilelang.contrib.rocm
    • tilelang.cpu
      • tilelang.cpu.codegen
      • tilelang.cpu.execution_backend
      • tilelang.cpu.op
        • tilelang.cpu.op.gemm
          • tilelang.cpu.op.gemm.gemm_scalar
      • tilelang.cpu.pipeline
    • tilelang.cuda
      • tilelang.cuda.codegen
      • 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_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_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.op
        • tilelang.cuda.op.gemm
          • tilelang.cuda.op.gemm.gemm_mma
          • 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.intrinsics
    • tilelang.ir
    • tilelang.jit
      • 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.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.env
      • 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.cluster
      • 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.fastmath
      • 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.pdl
      • tilelang.language.print_op
      • tilelang.language.proxy
      • tilelang.language.random
      • tilelang.language.reduce_op
      • tilelang.language.scan_op
      • tilelang.language.symbolics
      • tilelang.language.tile_schedule
      • tilelang.language.tir
        • tilelang.language.tir.entry
        • tilelang.language.tir.ir
        • tilelang.language.tir.op
      • tilelang.language.utils
      • tilelang.language.warpgroup
    • 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.codegen
      • tilelang.metal.execution_backend
      • tilelang.metal.intrinsics
        • tilelang.metal.intrinsics.metal_macro_generator
      • 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.profiler
      • tilelang.profiler.bench
    • tilelang.quantize
      • tilelang.quantize.lop3
      • tilelang.quantize.mxfp
      • tilelang.quantize.quantization
      • tilelang.quantize.utils
    • tilelang.rocm
      • 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.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.gemm_sp_wgmma
        • tilelang.tileop.gemm_sp.registry
    • tilelang.tools
      • tilelang.tools.Analyzer
      • 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.sparse
      • tilelang.utils.tensor
    • tilelang.webgpu
      • tilelang.webgpu.codegen
      • tilelang.webgpu.pipeline

Privacy

  • Privacy
Back to top
View this page

tilelang.metal.transform¶

Metal-specific transformation frontends.

Submodules¶

  • tilelang.metal.transform.mark_host_metal_context
  • tilelang.metal.transform.metal_fragment_to_simdgroup
Next
tilelang.metal.transform.mark_host_metal_context
Previous
tilelang.metal.target
Copyright © 2025-2025, Tile Lang Contributors
Made with Sphinx and @pradyunsg's Furo
On this page
  • tilelang.metal.transform
    • Submodules