tilelang.transform.pass_config¶

Classes¶

PassConfigKey

Pass configuration keys for TileLang compiler.

Module Contents¶

class tilelang.transform.pass_config.PassConfigKey¶

Bases: str, enum.Enum

Pass configuration keys for TileLang compiler.

TL_SIMPLIFY = 'tl.Simplify'¶

True

Type:: Enable/disable TileLang simplification passes. Default

TL_DYNAMIC_ALIGNMENT = 'tl.dynamic_alignment'¶

Type:: Memory alignment requirement for dynamic shapes. Default

TL_DISABLE_DYNAMIC_TAIL_SPLIT = 'tl.disable_dynamic_tail_split'¶

False

Type:: Disable dynamic tail splitting optimization. Default

TL_DISABLE_WARP_SPECIALIZED = 'tl.disable_warp_specialized'¶

False

Type:: Disable warp specialization optimization. Default

TL_DISABLE_FAST_MATH = 'tl.disable_fast_math'¶

False

Type:: Disable fast math optimization. Default

TL_ENABLE_PTXAS_VERBOSE_OUTPUT = 'tl.enable_ptxas_verbose_output'¶

False

Type:: Enable ptxas verbose output. Default

TL_CONFIG_INDEX_BITWIDTH = 'tl.config_index_bitwidth'¶

Type:: Bitwidth for configuration indices. Default

TL_DISABLE_TMA_LOWER = 'tl.disable_tma_lower'¶

False

Type:: Disable TMA (Tensor Memory Access) lowering. Default

TL_DISABLE_SAFE_MEMORY_ACCESS = 'tl.disable_safe_memory_legalize'¶

False

Type:: Disable safe memory access optimization. Default

TL_DEBUG_MERGE_SHARED_MEMORY_ALLOCATIONS = 'tl.debug_merge_shared_memory_allocations'¶

False

Type:: Enable debug information for merge shared memory allocations. Default

TL_ENABLE_AGGRESSIVE_SHARED_MEMORY_MERGE = 'tl.enable_aggressive_shared_memory_merge'¶

False

Type:: Enable aggressive merge of shared memory allocations. Default

TIR_ENABLE_EQUIV_TERMS_IN_CSE = 'tir.enable_equiv_terms_in_cse_tir'¶

True

Type:: Enable equivalent terms in TIR Common Subexpression Elimination. Default

TIR_DISABLE_CSE = 'tir.disable_cse_tir'¶

False

Type:: Disable TIR Common Subexpression Elimination. Default

TIR_SIMPLIFY = 'tir.Simplify'¶

True

Type:: Enable/disable TIR simplification passes. Default

TIR_DISABLE_STORAGE_REWRITE = 'tir.disable_storage_rewrite'¶

False

Type:: Disable storage rewrite optimization. Default

TIR_DISABLE_VECTORIZE = 'tir.disable_vectorize'¶

False

Type:: Disable vectorization optimization. Default

TIR_USE_ASYNC_COPY = 'tir.use_async_copy'¶

True

Type:: Enable asynchronous memory copy operations. Default

TIR_ENABLE_DEBUG = 'tir.enable_debug'¶

False

Type:: Enable debug information in generated code. Default

TIR_MERGE_STATIC_SMEM = 'tir.merge_static_smem'¶

True

Type:: Merge static shared memory allocations. Default

TIR_ADD_LOWER_PASS = 'tir.add_lower_pass'¶

None

Type:: Additional lowering passes to be applied. Default

TIR_NOALIAS = 'tir.noalias'¶

True

Type:: Enable pointer non-aliasing assumptions. Default

CUDA_KERNELS_OUTPUT_DIR = 'cuda.kernels_output_dir'¶

empty string

Type:: Output directory for generated CUDA kernels. Default