Rate this Page

★ ★ ★ ★ ★

Settings#

The Settings class controls compilation behavior and debugging options for Helion kernels.

class helion.Settings(**settings)[source]#

Bases: _Settings

Settings can be passed to hl.kernel as kwargs and control the behavior of the compilation process. Unlike a Config, settings are not auto-tuned and set by the user.

Parameters:: settings (object) –

__init__(**settings)[source]#

Initialize the Settings object with the provided dictionary of settings.

Parameters:: settings (object) –

copy(**overrides)[source]#

Copy values and origins, marking supplied structural overrides explicit.

Like dataclasses.replace, ordinary mutable field values are shared. In contrast, dataclasses.replace and Settings(**settings.to_dict()) supply every effective field to the constructor and therefore record every structural value as explicit, even when its value is unchanged.

Parameters:: overrides (object) –
Return type:: Settings

property cute_structural_origins: CuteStructuralOrigins#: Immutable origins captured at construction and explicit writes.

get_cute_structural_policy()[source]#

Snapshot effective booleans without resolving or activating a policy.

Return type:: CuteStructuralPolicy

to_dict()[source]#

Convert the Settings object to a dictionary.

Returns:: A dictionary representation of the Settings object.
Return type:: dict[str, object]

check_autotuning_disabled()[source]#

Return type:: None

get_rebenchmark_threshold()[source]#

Get the effective rebenchmark threshold. Uses the explicit setting if provided, otherwise falls back to the effort profile default.

Returns:: The rebenchmark threshold value.
Return type:: float

get_suspicious_rebenchmark_ratio()[source]#

Return type:: float | None

backend: str#: Code generation backend. One of ‘triton’ (default), ‘pallas’ (JAX/Pallas), ‘cute’ (CUTLASS CuTe DSL), or ‘metal’ (Apple Metal MSL). Set HELION_BACKEND=<backend> to override.

ignore_warnings: list[type[BaseWarning]]#: Subtypes of exc.BaseWarning to ignore when compiling. Set HELION_IGNORE_WARNINGS=WarningA,WarningB (names from helion.exc) to configure via env.

index_dtype: dtype | None#: The dtype to use for index variables. Default auto-selects torch.int32 or torch.int64 based on input sizes. Override with HELION_INDEX_DTYPE=<dtype> (or set to ‘auto’).

dot_precision: Literal['tf32', 'tf32x3', 'ieee', 'default', 'high', 'highest']#: Precision for dot products. For Triton backend, see triton.language.dot (can be ‘tf32’, ‘tf32x3’, ‘ieee’). For JAX/Pallas backend, accepted values emit Pallas default precision on TPU. Unified mappings exist so that any value can be used on any backend.

cute_rng_stream: Literal['auto', 'word0', 'philox4']#

auto (CuTe default) uses philox4 for explicit hl.rand and preserves other RNG operations. philox4 maps logical offset i to Philox(seed, i // 4)[i % 4], independently of configuration. Set word0 for the previous reproducible stream; other backends default to word0. Explicit philox4 rejects unsupported RNG operations. Override with HELION_CUTE_RNG_STREAM.

Type:: CuTe RNG policy

fast_math: bool#: If True, enable fast math approximations (Helion-level, including the CuTe flash-attention polynomial exp2 packets, and Inductor-level). May reduce numerical precision and change NaN/Inf behavior. Set HELION_FAST_MATH=1 to enable.

cute_flash_env_overrides: dict[str, str]#

HELION_CUTE_FLASH_* variables the kernel resolves as if set in the environment. The channel is captured once when the kernel binds (the process environment overlaid with this mapping); later environment changes do not reach the bound kernel.

Type:: CuTe flash-attention knob channel

static_shapes: bool#: If True, use static shapes for all tensors. This is a performance optimization. Set HELION_STATIC_SHAPES=0 to disable.

persistent_reserved_sms: int#: Number of streaming multiprocessors to reserve when launching persistent kernels. Set HELION_PERSISTENT_RESERVED_SMS=N (default 0) or pass persistent_reserved_sms=N to helion.kernel.

autotune_force_persistent: bool#: If True, restrict pid_type choices to persistent kernels only during config selection. Set HELION_AUTOTUNE_FORCE_PERSISTENT=1 to force persistent kernel autotuning globally.

distributed: bool#: Force distributed compilation behavior (persistent-only PID types, signal-pad SM limits, and process-group resolution). Normally auto-detected from symmetric-memory tensor arguments; set this for distributed kernels whose symmetric memory is not passed as a detectable tensor. Set HELION_DISTRIBUTED=1 to force globally.

autotune_log_level: int#: Log level for autotuning using Python logging levels. Default is logging.INFO. Use HELION_AUTOTUNE_LOG_LEVEL to override or set 0 to disable output.

autotune_log: str | None#: Base filename for autotune logs. Set HELION_AUTOTUNE_LOG=/tmp/run to write /tmp/run.csv and /tmp/run.log with per-config metrics and debug logs.

autotune_log_details: bool#: Opt-in (HELION_AUTOTUNE_LOG_DETAILS=1) to write /tmp/run.trace.jsonl with config evaluations, rebenchmarks, selection, exact LLM prompts/responses, and timestamps, plus the cost-model dataset sidecar /tmp/run.meta.jsonl (kernel identity and configs keyed by config_id). Off by default; needs autotune_log.

autotune_handoff: bool#: Hand configuration search to a native-source agent automatically. Off by default; enable with HELION_AUTOTUNE_HANDOFF=1. Uses the existing search budget or search completion, then saves validated standalone source beside the autotune logs.

autotune_handoff_budget_seconds: int#: Wall-clock budget for the native-source agent session, excluding search and export. Default 1500 seconds; override with HELION_AUTOTUNE_HANDOFF_BUDGET_SECONDS.

autotune_compile_timeout: int#

a candidate past it is killed and recorded as failed to compile, and the search goes on. Default is 60 seconds. In-process compiles cannot be interrupted by it; the autotune budget bounds them instead.

Type:: Timeout in seconds for one autotune precompile subprocess (autotune_precompile); a CuTe render worker past it is given up and its config rendered in-process instead (never failed for the timeout alone). Also the benchmark worker’s deadline for a candidate’s compile and first call (the backend JIT; never below autotune_benchmark_timeout, armed once the worker is ready so a fresh worker’s start is not the candidate’s; lowered along with the setting by autotune_adaptive_timeout)

autotune_benchmark_subprocess: bool#: Run the autotune benchmark phase in a long-lived spawn subprocess so a hung/slow kernel can be killed without losing autotune progress. Enabled by default. Set HELION_AUTOTUNE_BENCHMARK_SUBPROCESS=0 to disable.

autotune_benchmark_timeout: int#: Per-config wall-clock timeout in seconds for the subprocess benchmark phase, counted once the candidate’s compile and first call are done (those have autotune_compile_timeout). Only applies when autotune_benchmark_subprocess is enabled. Default 30 seconds, raised automatically on environments with a heavier subprocess cold-start.

autotune_worker_start_timeout: int | None#

twice the worker’s compile deadline, at least 120 seconds. Set HELION_AUTOTUNE_WORKER_START_TIMEOUT to override.

Type:: Seconds a fresh benchmark worker may take to start (the spawned interpreter, its imports and the device initialization) before it is killed and the candidate recorded as a worker failure. The start is never charged to the candidate’s compile deadline (autotune_compile_timeout), which is armed once the worker reports ready. Default None

autotune_precompile: Optional[Literal['spawn', 'fork']]#

‘fork’, ‘spawn’, or falsy/None to disable. Defaults to ‘fork’ on non-Windows platforms. With a pool a batch of configs is generated, precompiled in parallel, then timed; without one (disabled, or a backend such as CuTe that has no pool) each config is compiled and timed in turn, so results arrive as the batch runs and the autotune budget can stop it between configs. CuTe generates each candidate’s code in forked workers ahead of its turn in ‘fork’ mode (its only mode; ‘spawn’ disables the workers), so the serial part per config is the module load, the DSL JIT at first call and the timing.

Type:: Autotuner precompile mode

autotune_precompile_jobs: int | None#: Maximum concurrent Triton precompile processes, default to cpu count. Also the number of CuTe render workers alive at once, which default to the CPUs this process may run on, divided by the autotuning processes of the host (LOCAL_WORLD_SIZE or the process group size), capped at 32. Does not affect results or the autotune cache key.

autotune_random_seed: int#: Seed used for autotuner random number generation. Defaults to HELION_AUTOTUNE_RANDOM_SEED or a time-based seed.

autotune_best_of_k: int#: When > 1, run autotuning K times with different random seeds and keep the config with the best benchmark performance. Each trial uses seed = autotune_random_seed + i for i in range(K), so two runs with the same base seed produce the same K configs. Cost is K× the autotune wall time. Cache entries are keyed on K, so a K≥2 result will not be overwritten by a later K=1 run and vice versa. Default 1 (current single-trial behavior). Set HELION_AUTOTUNE_BEST_OF_K=N to override.

autotune_accuracy_check: bool#: If True, validate candidate configs against the baseline kernel output before accepting them during autotuning. Unless autotune_baseline_fn is set, the baseline is the output of the config the kernel runs by default (a compiler-promoted default when there is one, else the conservative autotuning reference, which is also the fallback when the promoted default fails to compile or run).

autotune_rebenchmark_threshold: float | None#: If a config is within threshold*best_perf, re-benchmark it to avoid outliers. Defaults to effort profile value. Set HELION_REBENCHMARK_THRESHOLD to override.

autotune_suspicious_rebenchmark_ratio: float | None#: When subprocess benchmarking is enabled, recheck rebenchmark timings below ratio*previous_timing before accepting them. Defaults to 0.9. Set HELION_AUTOTUNE_SUSPICIOUS_REBENCHMARK_RATIO=0 to disable.

autotune_search_acf: list[str]#: List of PTXAS Advanced Controls Files (ACFs) to search during autotuning. ACFs are highly specialized configurations for specific hardware and use cases; when autotuning with ACFs, default -O3 is always considered. Empty list disables.

autotune_progress_bar: bool#: If True, show progress bar during autotuning. Default is True. Set HELION_AUTOTUNE_PROGRESS_BAR=0 to disable.

autotune_max_generations: int | None#: Override the maximum number of generations for Pattern Search and Differential Evolution Search autotuning algorithms with HELION_AUTOTUNE_MAX_GENERATIONS=N or @helion.kernel(autotune_max_generations=N).

autotune_budget_seconds: int | None#

once a few configs have been timed and one has succeeded, the measured cost per config decides how many more are attempted and the rest are skipped, never below a CuTe flash surface’s structural coverage rows, so slow-compiling kernels keep time for the search generations. A population that fits its share is evaluated in full. Set with HELION_AUTOTUNE_BUDGET_SECONDS=N or @helion.kernel(autotune_budget_seconds=N). Default None (no budget).

Type:: Wall-clock budget in seconds for the entire autotune. When the budget is exceeded the search returns the best config found so far. The budget is checked between generations and before each config’s compile (with a precompile pool also before each config’s timing; without one a config is timed right after its compile). Once it is out nothing further is compiled or timed, except that while no config has been measured successfully the search still attempts up to 3 more candidates, so a budget that runs out early returns a config when one of those works; the overrun is bounded by those attempts (and, with a pool, by the precompile wait of the batch in flight). Without a precompile pool the initial population may take at most half of the search phase (the budget minus the time kept for the final selection), counted from the start of each of its batches

autotune_ignore_errors: bool#: If True, skip logging and raising autotune errors. Set HELION_AUTOTUNE_IGNORE_ERRORS=1 to enable globally. The one-line summary of the initial population (configs measured, failed, skipped for the budget, and the wall time per config) and the end-of-run failure counts are logged either way.

autotune_adaptive_timeout: bool#

CuTe render workers keep the configured autotune_compile_timeout (a worker past it is rendered in-process), and a timeout fitted on one shape’s render times would be shared by every shape of the kernel. The benchmark worker’s compile deadline follows the lowered setting.

Type:: If True, set the compile timeout threshold to be smaller for Triton compilation,based on a quantile of initial compile times (with a lower bound). Lower bound and quantile are set by the effort profile. Set HELION_AUTOTUNE_ADAPTIVE_TIMEOUT=0 to disable. Applies only to a Triton-style precompile pool

pallas_collective_id: int | None#: Optional Pallas collective namespace override. Distributed kernels derive a stable ID automatically; set this only when otherwise identical kernels use incompatible runtime communication groups.

pallas_interpret: bool#: If True, run Pallas kernels in interpret mode on CPU (no TPU needed). Defaults to HELION_PALLAS_INTERPRET env var.

pallas_topk_recall_target: float#: Recall target for the Pallas approximate top-k lowering. Must be in (0, 1]; use 1.0 when exact top-k results are required. Default 0.99.

cute_region_fission: bool#: If True, distribute proven row-independent materialized regions into ordered CuTe launches with independently tuned tile axes. Defaults to HELION_CUTE_REGION_FISSION (False).

cute_materialize_transformed_operands: bool#: If True, materialize proven pure reused matrix operands before a CuTe GEMM, with independent producer and consumer tile axes. Input dimensions used by the proof are specialized with normal runtime guards. Defaults to HELION_CUTE_MATERIALIZE_TRANSFORMED_OPERANDS (False).

cute_full_slice_matmul_tiling: bool#: If True, give eligible direct full-slice half/BF16 matmuls an explicit, tunable contraction tile loop on CuTe. Defaults to HELION_CUTE_FULL_SLICE_MATMUL_TILING (False).

cute_segmented_matmul_tiling: bool#: If True, expose proven offset-delimited independent matrix rows to the CuTe grouped scheduler. Defaults to HELION_CUTE_SEGMENTED_MATMUL_TILING (False).

cute_flatten_nested_reductions: bool#: If True, flatten proven FP32 joint jagged/feature reductions and independent fresh-output stores on CuTe. Defaults to HELION_CUTE_FLATTEN_NESTED_REDUCTIONS (False).

cute_graph_cse: bool#: If True (the default), merge duplicate reductions and the pure chains between them in the traced device IR before CuTe code generation (one amax/exp/sum per operand for a softmax and a log_softmax of it). Results are bitwise unchanged; set False (HELION_CUTE_GRAPH_CSE=0) to bisect a structural-matcher regression against the traced graph.

cute_stage_unaligned_operands: bool#: If True (the default), the CuTe backend copies a static-shape matmul operand whose outer strides are not 16-byte multiples (e.g. odd-sized bf16 rows; a misaligned base offset is not visible to the pass) into a padded row-major buffer on the host before the launch, so the tcgen05 pipeline keeps its TMA producers instead of the scalar shared-memory producers. Set False (HELION_CUTE_STAGE_UNALIGNED_OPERANDS=0) to run such operands in place.

cute_flat_edge_tma: bool#: If True (the default), flat one-CTA tcgen05 matmul kernels load partial output tiles and the K tail through TMA (zero-filled past the extents) and store them with predicated vector copies. Set False (HELION_CUTE_FLAT_EDGE_TMA=0) for the scalar shared-memory edge staging those tiles used before.

cute_exact_division: bool#

the IEEE round-to-nearest quotient bit for bit, without one div.rn expansion per element. Set False (HELION_CUTE_EXACT_DIVISION=0) to keep the plain per-element division; the config knob cute_exact_div_form picks the form (redo / plain) per kernel. Independent of fast_math, whose fp32 division is approximate.

Type:: If True (the default), an fp32 division by a loop-invariant scalar (a row sum, a scalar kernel argument) on CuTe hoists the divisor’s correctly rounded reciprocal and admission window out of the element loop and divides with the Markstein sequence, re-running the loop with div.rn when a dividend left the window

cute_promote_half_atomic_outputs: bool#: If True, a fresh float16 allocation that the CuTe kernel only ever writes through hl.atomic_add is accumulated in a float32 twin and cast to float16 once on return (one extra elementwise kernel per call). False (the default, HELION_CUTE_PROMOTE_HALF_ATOMIC_OUTPUTS=0) adds into the float16 storage itself, rounding to float16 on every add as the program and the Triton backend do; the vector collective epilogues use native packed half atomics, the scalar half atomic is a compare-and-swap loop on sm_100.

triton_do_not_specialize: bool#: If True, pass do_not_specialize for every dynamic size/stride/symbol arg so a single Triton binary handles any value, avoiding Triton-level recompiles when shapes cross 0/1 or alignment boundaries. Disables Triton’s value/alignment specializations and can significantly slow memory-bound kernels (vectorized loads rely on inner stride being specialized to constexpr 1). Only takes effect when static_shapes=False. Set HELION_TRITON_DO_NOT_SPECIALIZE=1 to enable globally.

print_output_code: bool#: If True, print the output code of the kernel to stderr.

print_repro: bool#: If True, print Helion kernel code, config, and caller code to stderr as a standalone repro script.

output_origin_lines: bool#: If True, annotate generated Triton code with source-origin comments. Set HELION_OUTPUT_ORIGIN_LINES=0 to disable.

force_autotune: bool#: If True, force autotuning even if a config is provided. The result is still written to the cache so subsequent runs can reuse it. Set HELION_SKIP_CACHE=1 instead to skip both reading and writing the cache.

autotune_config_overrides: dict[str, object]#

4}’.

Type:: Dictionary of config key/value pairs forced during autotuning. Accepts HELION_AUTOTUNE_CONFIG_OVERRIDES=’{“num_warps”

autotune_seed_configs: Union[Config, dict[str, object], CuteStructuralConfig, Sequence[Config | dict[str, object] | CuteStructuralConfig], None]#: A Config or sequence of Configs to seed the autotuner initial population without constraining the search space.

disable_autotuner_heuristics: bool#: If True, disable compiler seed generation and promotion. Correctness fact hooks still run so explicit configs normalize safely. User-provided autotune_seed_configs are unaffected. Set HELION_DISABLE_AUTOTUNER_HEURISTICS=1 to disable globally.

allow_warp_specialize: bool#: If True, allow warp specialization for tl.range calls on CUDA devices.

debug_dtype_asserts: bool#: If True, emit tl.static_assert checks for dtype after each device node.

ref_mode: RefMode#: Reference mode for kernel execution. Can be RefMode.OFF or RefMode.EAGER.

autotuner_fn: AutotunerFunction#: Function to create an autotuner. Override by passing a callable to @helion.kernel(…, autotuner_fn=…).

autotune_effort: Literal['none', 'quick', 'full']#: Autotuning effort preset. One of ‘none’, ‘quick’, ‘full’.

autotune_baseline_fn: Optional[Callable[..., object]]#: Custom baseline function for computing baseline output during autotuning. If provided, this function will be called instead of running the default config. Should have the same signature as the kernel function. Pass as @helion.kernel(…, autotune_baseline_fn=my_baseline_fn).

autotune_baseline_atol: float | None#: Absolute tolerance for baseline output comparison during autotuning accuracy checks. When unset, defaults to 1e-2 scaled per output tensor by max(1, rms(baseline)) so large-magnitude reductions tolerate accumulation-order noise, or 0.0 for fp8 dtypes under autotune_fp8_accuracy=’bitwise’. Setting an explicit value disables the RMS scaling and is applied exactly; an explicit 0 together with autotune_baseline_rtol=0 compares fp8 outputs bitwise. Pass as @helion.kernel(…, autotune_baseline_atol=1e-3).

autotune_baseline_rtol: float | None#: Relative tolerance for baseline output comparison during autotuning accuracy checks. Defaults to 1e-2, or 0.0 for fp8 dtypes under autotune_fp8_accuracy=’bitwise’. Pass as @helion.kernel(…, autotune_baseline_rtol=1e-3). See autotune_baseline_atol for how the absolute-tolerance floor scales when unset.

autotune_fp8_accuracy: Literal['ulp', 'bitwise']#

every element must lie within one fp8 ulp of max(|actual|, |baseline|) plus the atol/rtol tolerances (atol scaled by the baseline RMS as usual), and the one-ulp differences must not be systematically one-directional. This admits the one-bucket flips a kernel that quantizes an intermediate to fp8 (e.g. an online-softmax P) produces when its reduction blocking changes, while still rejecting multi-ulp and biased differences. ‘bitwise’ requires identical bytes whenever every output is fp8 and no explicit tolerance is set (the previous behavior). Explicit zero tolerances (autotune_baseline_atol=0 with autotune_baseline_rtol=0) ask for exact outputs and select ‘bitwise’ whatever this setting says. Set HELION_AUTOTUNE_FP8_ACCURACY=bitwise to override.

Type:: How autotuning accuracy checks compare fp8 output tensors against the baseline. ‘ulp’ (default) compares dequantized values

autotune_baseline_accuracy_check_fn: Optional[Callable[[object, object], None]]#

object, expected: object) -> None. Should raise AssertionError on mismatch. When set, replaces the default torch.testing.assert_close-based check (atol/rtol settings are ignored). Useful for scenarios where a small fraction of elements may have large relative differences, e.g. checking that mismatch percentage < X AND max relative diff < Y. A built-in utility helion._testing.assert_close_with_mismatch_tolerance is provided for this common pattern; use functools.partial(assert_close_with_mismatch_tolerance, ...) to customize thresholds. Pass as @helion.kernel(…, autotune_baseline_accuracy_check_fn=my_check_fn).

Type:: Custom accuracy check function for comparing autotuning candidate outputs against the baseline. Signature
Type:: (actual

autotune_cache: str#: The name of the autotuner cache class to use. Set HELION_AUTOTUNE_CACHE=StrictLocalAutotuneCache to enable strict caching. Set HELION_AUTOTUNE_CACHE=RemoteAutotuneCache (or StrictRemoteAutotuneCache) and HELION_REMOTE_CACHE_BACKEND=mypackage.module.MyBackend to enable remote caching via a user-provided RemoteCacheBackend subclass. Defaults to ‘LocalAutotuneCache’.

autotune_benchmark_fn: Optional[Callable[..., list[float]]]#

list[Callable[[], object]], *, repeat: int, desc: str | None = None) -> list[float]. If None (default), uses the built-in benchmark function.

Type:: Custom benchmark function for rebenchmarking during autotuning. Should have the following signature
Type:: (fns

autotune_best_available_max_configs: int#: Maximum number of cached configs to use for FROM_BEST_AVAILABLE initial population and for helion.from_cache() warm-start in FiniteSearch. Set HELION_BEST_AVAILABLE_MAX_CONFIGS=N to override. Default is 20.

autotune_best_available_max_cache_scan: int#: Maximum number of cache files to scan when searching for matching configs in FROM_BEST_AVAILABLE strategy. Set HELION_BEST_AVAILABLE_MAX_CACHE_SCAN=N to override. Default is 500.

autotune_initial_population_strategy: Optional[Literal['from_random', 'from_best_available']]#

‘from_random’, ‘from_best_available’. When set, takes precedence over the HELION_AUTOTUNER_INITIAL_POPULATION env var and the effort profile default.

Type:: Override the initial population strategy for autotuning. Valid values

torch_compile_fusion: bool#: If True, allow torch.compile to fuse this Helion kernel with surrounding Inductor ops (prologue/epilogue) when used inside torch.compile. Default False. Set HELION_TORCH_COMPILE_FUSION=1 to enable globally.

autotune_config_filter: Optional[Callable[[Config], Config | None]]#: Optional callable (config: Config) -> Config | None that the autotuner calls on every candidate config before compiling or benchmarking it. If the callable returns None, the config is skipped entirely (no compilation, no benchmarking). If it returns a Config (which may be a modified copy of the original), that config is used for benchmarking. Also filters the explicit configs=[...] list when one is provided. Pass as @helion.kernel(…, autotune_config_filter=my_filter_fn).

autotune_with_torch_compile_fusion: bool#: If True, autotuning benchmarks the fused kernel (with epilogue/prologue) to pick configs optimal for the actual fused workload. Default False. Has no effect unless torch_compile_fusion is also True. Set HELION_AUTOTUNE_WITH_TORCH_COMPILE_FUSION=1 to enable globally.

autotune_log_search_space: bool#: If True, log search space analysis after autotuning including which features were enabled/disabled, total search space size, and coverage metrics. Off by default; set HELION_AUTOTUNE_LOG_SEARCH_SPACE=1 to enable.

autotune_log_search_space_verbose: bool#: If True, additionally log each search-space restriction (disabled pid_types, tcgen05 narrowing, etc.) live at INFO the moment it is applied during compilation, on top of the end-of-run summary. Implies autotune_log_search_space. Off by default; set HELION_AUTOTUNE_LOG_SEARCH_SPACE_VERBOSE=1 to enable.

autotune_log_search_space_path: str | None#: Optional path to save search space analysis JSON. Set HELION_AUTOTUNE_LOG_SEARCH_SPACE_PATH=/path/to/analysis.json to save.

Overview#

Settings control the compilation process and development environment for Helion kernels.

Key Characteristics#

Not autotuned: Settings remain constant across all kernel configurations
Meta-compilation: Control the compilation process itself, debugging output, and development features
Environment-driven: Often configured via environment variables
Development-focused: Primarily used for debugging, logging, and development workflow optimization

Settings vs Config#

Aspect	Settings	Config
Purpose	Control compilation behavior	Control execution performance
Autotuning	❌ Never autotuned	✅ Automatically optimized
Examples	`print_output_code`, `autotune_effort`	`block_sizes`, `num_warps`
When to use	Development, debugging, environment setup	Performance optimization

Settings can be configured via:

Environment variables
Keyword arguments to @helion.kernel

If both are provided, decorator arguments take precedence.

Note

Helion reads the environment variables for Settings when the @helion.kernel decorator defines the function (typically at import time). One can modify Kernel.settings to change settings for an already defined kernel.

Configuration Examples#

Using Environment Variables#

env HELION_PRINT_OUTPUT_CODE=1  HELION_AUTOTUNE_EFFORT=none my_kernel.py

Using Decorator Arguments#

import logging
import helion
import helion.language as hl

@helion.kernel(
    autotune_effort="none",           # Skip autotuning
    print_output_code=True,            # Debug: show generated Triton code
    print_repro=True,                  # Debug: show Helion kernel code, config, and caller code as a standalone repro script
)
def my_kernel(x: torch.Tensor) -> torch.Tensor:
    result = torch.zeros_like(x)
    for i in hl.grid(x.size(0)):
        result[i] = x[i] * 2
    return result

Settings Reference#

Core Compilation Settings#

Settings.index_dtype: dtype | None#

The dtype to use for index variables. Default auto-selects torch.int32 or torch.int64 based on input sizes. Override with HELION_INDEX_DTYPE=<dtype> (or set to ‘auto’).

The data type used for index variables in generated code. By default Helion auto-selects between torch.int32 and torch.int64 based on whether any input tensor exceeds torch.iinfo(torch.int32).max elements. Override via HELION_INDEX_DTYPE=<dtype> (or set it to auto to keep the automatic behavior).

Settings.dot_precision: Literal['tf32', 'tf32x3', 'ieee', 'default', 'high', 'highest']#

Precision for dot products. For Triton backend, see triton.language.dot (can be ‘tf32’, ‘tf32x3’, ‘ieee’). For JAX/Pallas backend, accepted values emit Pallas default precision on TPU. Unified mappings exist so that any value can be used on any backend.

Precision mode for dot product operations. - For Triton backend, this is initialized from the TRITON_F32_DEFAULT environment variable (defaulting to "tf32"). - For Pallas backend, this is initialized from the JAX_DEFAULT_MATMUL_PRECISION environment variable (defaulting to "default").

Supported values depend on the backend: - Triton (GPU): "tf32", "tf32x3", "ieee" - Pallas (TPU): "default", "high", "highest" and portable Triton aliases are accepted.

However, unified mappings exist so you can use any value on any backend: - On Triton: "default" maps to "tf32", "high" maps to "tf32x3", and "highest" maps to "ieee". - On Pallas/TPU: all values currently emit JAX default precision. JAX "high"/"highest" fp32 dot precision is not used because it is less compatible with PyTorch eager references on the supported TPU stack.

Settings.cute_rng_stream: Literal['auto', 'word0', 'philox4']#

auto (CuTe default) uses philox4 for explicit hl.rand and preserves other RNG operations. philox4 maps logical offset i to Philox(seed, i // 4)[i % 4], independently of configuration. Set word0 for the previous reproducible stream; other backends default to word0. Explicit philox4 rejects unsupported RNG operations. Override with HELION_CUTE_RNG_STREAM.

Type:: CuTe RNG policy

CuTe defaults to "auto". Explicit hl.rand uses all four Philox outputs: logical offset i selects word i % 4 from counter i // 4. The stream depends on the seed and logical offset, independently of tile size or packet vectorization. This changes the sequence from the earlier CuTe default.

Set cute_rng_stream="word0" or HELION_CUTE_RNG_STREAM=word0 to retain the previous sequence. Other backends still default to "word0". Under "auto", implicit Torch random operations, hl.randint and hl.rand4x retain their existing behavior. Explicit "philox4" selects the same uniform stream but rejects these other RNG operations. The policy is serialized with settings and emitted reproduction decorators.

Settings.cute_graph_cse: bool#

If True (the default), merge duplicate reductions and the pure chains between them in the traced device IR before CuTe code generation (one amax/exp/sum per operand for a softmax and a log_softmax of it). Results are bitwise unchanged; set False (HELION_CUTE_GRAPH_CSE=0) to bisect a structural-matcher regression against the traced graph.

CuTe only. When True (the default) duplicate reductions of the traced device IR, and the pure chains between them, are merged before code generation: a softmax and a log_softmax of one operand share their amax, exp and sum sweeps. Two identical pure operators on the same inputs are bitwise identical, so results never change. Set cute_graph_cse=False or HELION_CUTE_GRAPH_CSE=0 to keep the graph as traced, which bisects a structural-matcher regression against the pass. The two graphs roll different reduction sweeps, so their configs differ in shape; printed repros record the value in their @helion.kernel decorator and replay the graph they were emitted from whatever the replay environment.

Settings.cute_promote_half_atomic_outputs: bool#

If True, a fresh float16 allocation that the CuTe kernel only ever writes through hl.atomic_add is accumulated in a float32 twin and cast to float16 once on return (one extra elementwise kernel per call). False (the default, HELION_CUTE_PROMOTE_HALF_ATOMIC_OUTPUTS=0) adds into the float16 storage itself, rounding to float16 on every add as the program and the Triton backend do; the vector collective epilogues use native packed half atomics, the scalar half atomic is a compare-and-swap loop on sm_100.

CuTe only. hl.atomic_add into a float16 tensor rounds to float16 on every add (the reference semantics, what the Triton backend emits). When True the CuTe backend instead accumulates a fresh float16 allocation that the kernel only ever writes through hl.atomic_add in a float32 twin and casts it to float16 once on return: more accurate (one rounding instead of one per add), at the price of one extra elementwise kernel per call (about 1.5 us in a CUDA graph) and results that differ from the other backends. False (the default, HELION_CUTE_PROMOTE_HALF_ATOMIC_OUTPUTS=0) keeps the requested float16 storage; the collective vector epilogues then use native packed f16x2 / bf16x2 reductions (the scalar half atomic is a compare-and-swap loop on sm_100, so a lane-loop fallback is slow).

With the promotion off, a split-K kernel’s rounding depends on the config the autotuner picks: the atomic routes (the collective vector epilogue, and the split-K workspace schedule’s reducer-free epilogue, cute_split_k_atomic_epilogue) round each K partition’s fp32 partial to float16 as it is added, while the workspace schedule’s reducer form sums the fp32 partials and rounds once; two configs of the same kernel can thus differ by a few float16 ulps on large-K outputs. True restores config-independent fp32 accumulation (every route accumulates the fp32 twin and casts once) at the cost of the extra cast kernel.

Settings.cute_exact_division: bool#

the IEEE round-to-nearest quotient bit for bit, without one div.rn expansion per element. Set False (HELION_CUTE_EXACT_DIVISION=0) to keep the plain per-element division; the config knob cute_exact_div_form picks the form (redo / plain) per kernel. Independent of fast_math, whose fp32 division is approximate.

Type:: If True (the default), an fp32 division by a loop-invariant scalar (a row sum, a scalar kernel argument) on CuTe hoists the divisor’s correctly rounded reciprocal and admission window out of the element loop and divides with the Markstein sequence, re-running the loop with div.rn when a dividend left the window

CuTe only. When True (the default) an fp32 division by a scalar that is invariant across an element loop (a softmax row sum, a kernel-wide temperature) hoists the divisor’s correctly rounded reciprocal and its admission window out of the loop and divides with the Markstein sequence (q = RN(x*y), r = RN(x - d*q), RN(q + r*y)), speculated over the loop: when some dividend left the window the loop runs again as written, with div.rn.f32 itself. The result is the IEEE round-to-nearest quotient bit for bit (helion/_compiler/cute/exact_division.py states the argument); set cute_exact_division=False or HELION_CUTE_EXACT_DIVISION=0 to keep one div.rn expansion per element; the config knob cute_exact_div_form ("redo", "plain") picks the form per kernel, so the autotuner can keep the plain division where a kernel sits at the register cap (the knob is exposed for every kernel whose traced graph divides fp32 by a scalar or a broadcast row value; a loop the redo form cannot re-execute keeps its plain division, and where the pass proves no divisor loop-invariant both forms render identically). fast_math is unaffected: its fp32 division is the approximate cute.math.div.

Settings.static_shapes: bool#

If True, use static shapes for all tensors. This is a performance optimization. Set HELION_STATIC_SHAPES=0 to disable.

When enabled, tensor shapes are treated as compile-time constants for optimization. Default is True. Set HELION_STATIC_SHAPES=0 the default if you need a compiled kernel instance to serve many shape variants.

Settings.fast_math: bool#

If True, enable fast math approximations (Helion-level, including the CuTe flash-attention polynomial exp2 packets, and Inductor-level). May reduce numerical precision and change NaN/Inf behavior. Set HELION_FAST_MATH=1 to enable.

If True, enable fast math approximations. This activates both Helion-level optimizations (e.g. fast sigmoid, the CuTe flash-attention polynomial exp2 packets) and Inductor-level fast math (flush-to-zero exp, fast online softmax, etc.). May reduce numerical precision and change NaN/Inf behavior. Default is False. Controlled by HELION_FAST_MATH=1.

Settings.cute_flash_env_overrides: dict[str, str]#

HELION_CUTE_FLASH_* variables the kernel resolves as if set in the environment. The channel is captured once when the kernel binds (the process environment overlaid with this mapping); later environment changes do not reach the bound kernel.

Type:: CuTe flash-attention knob channel

CuTe flash-attention knob channel: a mapping of HELION_CUTE_FLASH_* names to values the kernel resolves as if they were set in the environment. The channel is captured once when the kernel binds (the process environment overlaid with this mapping); environment changes made after the bind do not reach the bound kernel. Default is {}.

Settings.persistent_reserved_sms: int#

Number of streaming multiprocessors to reserve when launching persistent kernels. Set HELION_PERSISTENT_RESERVED_SMS=N (default 0) or pass persistent_reserved_sms=N to helion.kernel.

Reserve this many streaming multiprocessors when launching persistent kernels. Default is 0 (use all SMs). Configure globally with HELION_PERSISTENT_RESERVED_SMS or per-kernel via @helion.kernel(..., persistent_reserved_sms=N).

Autotuning Settings#

Settings.force_autotune: bool#

If True, force autotuning even if a config is provided. The result is still written to the cache so subsequent runs can reuse it. Set HELION_SKIP_CACHE=1 instead to skip both reading and writing the cache.

Force autotuning even when explicit configs are provided. Default is False. Controlled by HELION_FORCE_AUTOTUNE=1. The result is still saved to the cache so subsequent runs can reuse it. Use HELION_SKIP_CACHE=1 instead to skip both reading and writing the cache.

Settings.autotune_force_persistent: bool#

If True, restrict pid_type choices to persistent kernels only during config selection. Set HELION_AUTOTUNE_FORCE_PERSISTENT=1 to force persistent kernel autotuning globally.

Restrict pid_type choices to the persistent strategies ("persistent_blocked" or "persistent_interleaved"). Default is False. Enable globally with HELION_AUTOTUNE_FORCE_PERSISTENT=1 or per kernel via @helion.kernel(..., autotune_force_persistent=True).

Settings.autotune_log_level: int#

Log level for autotuning using Python logging levels. Default is logging.INFO. Use HELION_AUTOTUNE_LOG_LEVEL to override or set 0 to disable output.

Controls verbosity of autotuning output using Python logging levels:

logging.CRITICAL: No autotuning output
logging.WARNING: Only warnings and errors
logging.INFO: Standard progress messages (default)
logging.DEBUG: Verbose debugging output

You can also use 0 to completely disable all autotuning output. Controlled by HELION_AUTOTUNE_LOG_LEVEL.

Settings.autotune_log: str | None#

Base filename for autotune logs. Set HELION_AUTOTUNE_LOG=/tmp/run to write /tmp/run.csv and /tmp/run.log with per-config metrics and debug logs.

When set, Helion writes per-config telemetry (run id, timestamp, config id, generation, status, perf, compile time, config) to <value>.csv and mirrors the autotune log to <value>.log (for population-based autotuners: PatternSearch, DifferentialEvolution). Both append, so runs sharing one base path accumulate. CSV rows join to .meta.jsonl via two hashes: run_id for the invocation — hash(kernel source, shapes, dtypes, hardware, codegen settings) — and config_id for the config. The codegen settings hashed into run_id are: backend, dot_precision, fast_math, static_shapes, index_dtype, allow_warp_specialize, triton_do_not_specialize, pallas_interpret, debug_dtype_asserts, persistent_reserved_sms (the full settings are recorded in .meta.jsonl). Same config_id across started/ok/error and re-benchmarks. The full config is written inline in the trailing config column of the CSV; it is also stored in the .meta.jsonl configs map, which is written only when autotune_log_details is enabled (see below). Controlled by HELION_AUTOTUNE_LOG.

Settings.autotune_log_details: bool#

Opt-in (HELION_AUTOTUNE_LOG_DETAILS=1) to write /tmp/run.trace.jsonl with config evaluations, rebenchmarks, selection, exact LLM prompts/responses, and timestamps, plus the cost-model dataset sidecar /tmp/run.meta.jsonl (kernel identity and configs keyed by config_id). Off by default; needs autotune_log.

Opt-in to the cost-model dataset sidecar. When enabled (HELION_AUTOTUNE_LOG_DETAILS=1) and autotune_log is set, Helion appends one JSON record per run to <autotune_log>.meta.jsonl: the kernel identity (run_id, name, source, shapes, dtypes, hardware), the full helion.settings (JSON-safe via json.dumps(default=str), so torch.dtype/callables become strings), an ir_graph (see below), and a configs map from config_id to the config tested. ir_graph is a config-independent, node-link dump of the kernel’s device IR (the per-tile computation lowered to the Triton kernel body), captured once per run. Load it with networkx.node_link_graph(record["ir_graph"], edges="edges") (requires networkx>=3.4). It is null when the device IR is unavailable or extraction fails. Block dimensions stay symbolic in the dump; recover concrete tile sizes by joining with each config’s block_sizes. Because the device IR is a function of run_id’s inputs, every record sharing a run_id carries an identical ir_graph. Recover a measured (config, perf) sample by joining a CSV row to its record: meta[run_id]["configs"][row["config_id"]]. run_id may recur (re-runs, processes, autotune_best_of_k), but the configs maps are union-safe (same config_id implies the same config), so de-duplicating on run_id is lossless. Searches restricted to user-pinned configs (without force_autotune) are excluded as a biased slice (.csv/.log still written); setting this without autotune_log collects nothing and warns once. Controlled by HELION_AUTOTUNE_LOG_DETAILS.

This flag also appends <autotune_log>.trace.jsonl with timestamped search stages, config evaluations, rebenchmarks, selection, and exact LLM prompts and responses. Recording is disabled by default; set HELION_AUTOTUNE_LOG=/tmp/run and HELION_AUTOTUNE_LOG_DETAILS=1 before importing the kernel. Each trace uses a per-search run_id, distinct from the identity used by the CSV and metadata.

Settings.autotune_log_search_space: bool#

If True, log search space analysis after autotuning including which features were enabled/disabled, total search space size, and coverage metrics. Off by default; set HELION_AUTOTUNE_LOG_SEARCH_SPACE=1 to enable.

Enable search space analysis logging after autotuning. When enabled, Helion logs:

Search space summary: Total search space size, enabled/disabled features with reasons, shape constraints
Coverage metrics: Configs tested vs. total space, coverage percentage
Feature exploration report: Per-feature statistics showing how many options were tested (e.g., “pid_type: 2/8 options tested (25.0%)”)
Exploration quality: Average and minimum feature coverage, highlighting poorly explored features (<50%)

The total search space size is the exact combinatorial product of every dimension’s cardinality (an arbitrary-precision integer, so for large kernels it can be very large); it is reported as unknown only when some dimension’s cardinality cannot be determined.

Default is False. Controlled by HELION_AUTOTUNE_LOG_SEARCH_SPACE (set to 1 to enable).

See Autotuner Module for example output and the helion.autotuner.search_space_logger module for implementation details.

Settings.autotune_log_search_space_verbose: bool#

If True, additionally log each search-space restriction (disabled pid_types, tcgen05 narrowing, etc.) live at INFO the moment it is applied during compilation, on top of the end-of-run summary. Implies autotune_log_search_space. Off by default; set HELION_AUTOTUNE_LOG_SEARCH_SPACE_VERBOSE=1 to enable.

Additionally log each individual search-space restriction live (at INFO) the moment it is applied during kernel compilation, on top of the end-of-run summary. This surfaces why the search space shrank as it happens, for example:

Autotuner feature restriction: pid_type='xyz' disabled (data-dependent loop bounds require a persistent kernel ...)
Autotuner feature restriction: tcgen05 search narrowed to validated configs (matmul kernel with CuTe tcgen05 backend)

Enabling this implies autotune_log_search_space (the summary is produced as well). Default is False. Controlled by HELION_AUTOTUNE_LOG_SEARCH_SPACE_VERBOSE (set to 1 to enable).

Settings.autotune_log_search_space_path: str | None#

Optional path to save search space analysis JSON. Set HELION_AUTOTUNE_LOG_SEARCH_SPACE_PATH=/path/to/analysis.json to save.

Optional path to save search space analysis as JSON. When set, Helion writes one JSON file per autotuned kernel/shape. The kernel name and the autotuner’s stable cache hash are injected into the filename stem (matching the .best_config cache key) so distinct kernels/shapes do not overwrite each other:

<stem>.<kernel>.<hash>.json: a single report combining the search-space description (dimensions, possible values, constraints, total size) with the exploration outcome (per-dimension observed values and coverage, configs tested, valid/invalid attempts).

For example, autotune_log_search_space_path="/tmp/analysis.json" yields a file like /tmp/analysis.my_kernel.3f9a1c2e.json. A directory (existing path or trailing separator) is also accepted, in which case the default autotune_search_space.json filename is used inside it. Requires autotune_log_search_space to be enabled. Default is None (only logs to console). Controlled by HELION_AUTOTUNE_LOG_SEARCH_SPACE_PATH.

Settings.autotune_compile_timeout: int#

a candidate past it is killed and recorded as failed to compile, and the search goes on. Default is 60 seconds. In-process compiles cannot be interrupted by it; the autotune budget bounds them instead.

Type:: Timeout in seconds for one autotune precompile subprocess (autotune_precompile); a CuTe render worker past it is given up and its config rendered in-process instead (never failed for the timeout alone). Also the benchmark worker’s deadline for a candidate’s compile and first call (the backend JIT; never below autotune_benchmark_timeout, armed once the worker is ready so a fresh worker’s start is not the candidate’s; lowered along with the setting by autotune_adaptive_timeout)

Timeout in seconds for one precompile subprocess during autotuning, and the benchmark worker’s deadline for a candidate’s compile and first call (never below autotune_benchmark_timeout; a candidate past it is recorded as failed to compile and the search goes on). The worker’s deadline is armed once the worker is ready, so a fresh worker’s start is not the candidate’s, and autotune_adaptive_timeout does not lower it. Default is 60. Controlled by HELION_AUTOTUNE_COMPILE_TIMEOUT.

Settings.autotune_worker_start_timeout: int | None#

twice the worker’s compile deadline, at least 120 seconds. Set HELION_AUTOTUNE_WORKER_START_TIMEOUT to override.

Type:: Seconds a fresh benchmark worker may take to start (the spawned interpreter, its imports and the device initialization) before it is killed and the candidate recorded as a worker failure. The start is never charged to the candidate’s compile deadline (autotune_compile_timeout), which is armed once the worker reports ready. Default None

Seconds a fresh benchmark worker may take to start (the spawned interpreter, its imports and the device initialization) before it is killed and the candidate recorded as a worker failure. Default is None: twice the worker’s compile deadline, at least 120. Controlled by HELION_AUTOTUNE_WORKER_START_TIMEOUT.

Settings.autotune_precompile: Optional[Literal['spawn', 'fork']]#

‘fork’, ‘spawn’, or falsy/None to disable. Defaults to ‘fork’ on non-Windows platforms. With a pool a batch of configs is generated, precompiled in parallel, then timed; without one (disabled, or a backend such as CuTe that has no pool) each config is compiled and timed in turn, so results arrive as the batch runs and the autotune budget can stop it between configs. CuTe generates each candidate’s code in forked workers ahead of its turn in ‘fork’ mode (its only mode; ‘spawn’ disables the workers), so the serial part per config is the module load, the DSL JIT at first call and the timing.

Type:: Autotuner precompile mode

Select the autotuner precompile mode, which adds parallelism and checks for errors/timeouts. "fork" (default) is faster but does not include the error check run, "spawn" runs kernel warm-up in a fresh process including running to check for errors, or None to disables precompile checks altogether. Controlled by HELION_AUTOTUNE_PRECOMPILE.

Settings.autotune_random_seed: int#

Seed used for autotuner random number generation. Defaults to HELION_AUTOTUNE_RANDOM_SEED or a time-based seed.

Seed used for autotuner random number generation. Defaults to HELION_AUTOTUNE_RANDOM_SEED if set, otherwise a time-based value.

Settings.autotune_precompile_jobs: int | None#

Maximum concurrent Triton precompile processes, default to cpu count. Also the number of CuTe render workers alive at once, which default to the CPUs this process may run on, divided by the autotuning processes of the host (LOCAL_WORLD_SIZE or the process group size), capped at 32. Does not affect results or the autotune cache key.

Cap the number of concurrent Triton precompile subprocesses. None (default) uses the machine CPU count. Controlled by HELION_AUTOTUNE_PRECOMPILE_JOBS. When using "spawn" precompile mode, Helion may automatically lower this cap if free GPU memory is limited.

Settings.autotune_max_generations: int | None#

Override the maximum number of generations for Pattern Search and Differential Evolution Search autotuning algorithms with HELION_AUTOTUNE_MAX_GENERATIONS=N or @helion.kernel(autotune_max_generations=N).

Override the default number of generations set for Pattern Search and Differential Evolution Search autotuning algorithms with HELION_AUTOTUNE_MAX_GENERATIONS=N or @helion.kernel(autotune_max_generations=N).

Lower values result in faster autotuning but may find less optimal configurations.

Settings.autotune_budget_seconds: int | None#

once a few configs have been timed and one has succeeded, the measured cost per config decides how many more are attempted and the rest are skipped, never below a CuTe flash surface’s structural coverage rows, so slow-compiling kernels keep time for the search generations. A population that fits its share is evaluated in full. Set with HELION_AUTOTUNE_BUDGET_SECONDS=N or @helion.kernel(autotune_budget_seconds=N). Default None (no budget).

Type:: Wall-clock budget in seconds for the entire autotune. When the budget is exceeded the search returns the best config found so far. The budget is checked between generations and before each config’s compile (with a precompile pool also before each config’s timing; without one a config is timed right after its compile). Once it is out nothing further is compiled or timed, except that while no config has been measured successfully the search still attempts up to 3 more candidates, so a budget that runs out early returns a config when one of those works; the overrun is bounded by those attempts (and, with a pool, by the precompile wait of the batch in flight). Without a precompile pool the initial population may take at most half of the search phase (the budget minus the time kept for the final selection), counted from the start of each of its batches

Wall-clock budget in seconds for autotuning. When the budget is exceeded, Helion returns the best configuration found so far. Controlled by HELION_AUTOTUNE_BUDGET_SECONDS.

Settings.autotune_handoff: bool#

Hand configuration search to a native-source agent automatically. Off by default; enable with HELION_AUTOTUNE_HANDOFF=1. Uses the existing search budget or search completion, then saves validated standalone source beside the autotune logs.

Enable automatic native-source optimization after configuration search with HELION_AUTOTUNE_HANDOFF=1. Off by default. Uses the existing search budget or search completion, then runs one continuous Codex session and saves standalone kernels beside the autotune logs. Each candidate submission records a round.

Settings.autotune_handoff_budget_seconds: int#

Wall-clock budget for the native-source agent session, excluding search and export. Default 1500 seconds; override with HELION_AUTOTUNE_HANDOFF_BUDGET_SECONDS.

Time budget for native-source rounds, excluding search and export. Defaults to 1500 seconds. Controlled by HELION_AUTOTUNE_HANDOFF_BUDGET_SECONDS.

Settings.autotune_ignore_errors: bool#

If True, skip logging and raising autotune errors. Set HELION_AUTOTUNE_IGNORE_ERRORS=1 to enable globally. The one-line summary of the initial population (configs measured, failed, skipped for the budget, and the wall time per config) and the end-of-run failure counts are logged either way.

Continue autotuning even when candidate configurations raise recoverable runtime errors (for example, GPU out-of-memory). Default is False. Controlled by HELION_AUTOTUNE_IGNORE_ERRORS.

Settings.autotune_accuracy_check: bool#

If True, validate candidate configs against the baseline kernel output before accepting them during autotuning. Unless autotune_baseline_fn is set, the baseline is the output of the config the kernel runs by default (a compiler-promoted default when there is one, else the conservative autotuning reference, which is also the fallback when the promoted default fails to compile or run).

Validate each candidate configuration against a baseline output before accepting it. Default is True. Controlled by HELION_AUTOTUNE_ACCURACY_CHECK.

Settings.autotune_baseline_atol: float | None#

Absolute tolerance for baseline output comparison during autotuning accuracy checks. When unset, defaults to 1e-2 scaled per output tensor by max(1, rms(baseline)) so large-magnitude reductions tolerate accumulation-order noise, or 0.0 for fp8 dtypes under autotune_fp8_accuracy=’bitwise’. Setting an explicit value disables the RMS scaling and is applied exactly; an explicit 0 together with autotune_baseline_rtol=0 compares fp8 outputs bitwise. Pass as @helion.kernel(…, autotune_baseline_atol=1e-3).

Absolute tolerance for baseline output comparison during autotune accuracy checks. Default is 1e-2.

Settings.autotune_baseline_rtol: float | None#

Relative tolerance for baseline output comparison during autotuning accuracy checks. Defaults to 1e-2, or 0.0 for fp8 dtypes under autotune_fp8_accuracy=’bitwise’. Pass as @helion.kernel(…, autotune_baseline_rtol=1e-3). See autotune_baseline_atol for how the absolute-tolerance floor scales when unset.

Relative tolerance for baseline output comparison during autotune accuracy checks. Default is 1e-2.

Settings.autotune_fp8_accuracy: Literal['ulp', 'bitwise']#

every element must lie within one fp8 ulp of max(|actual|, |baseline|) plus the atol/rtol tolerances (atol scaled by the baseline RMS as usual), and the one-ulp differences must not be systematically one-directional. This admits the one-bucket flips a kernel that quantizes an intermediate to fp8 (e.g. an online-softmax P) produces when its reduction blocking changes, while still rejecting multi-ulp and biased differences. ‘bitwise’ requires identical bytes whenever every output is fp8 and no explicit tolerance is set (the previous behavior). Explicit zero tolerances (autotune_baseline_atol=0 with autotune_baseline_rtol=0) ask for exact outputs and select ‘bitwise’ whatever this setting says. Set HELION_AUTOTUNE_FP8_ACCURACY=bitwise to override.

Type:: How autotuning accuracy checks compare fp8 output tensors against the baseline. ‘ulp’ (default) compares dequantized values

How autotune accuracy checks compare fp8 output tensors against the baseline. "ulp" (default) compares dequantized values: each element must lie within one fp8 ulp plus the tolerances, and the one-ulp differences must not be systematically one-directional. "bitwise" requires identical bytes whenever every output is fp8 and no explicit tolerance is set. Explicit zero tolerances (autotune_baseline_atol=0 with autotune_baseline_rtol=0) select "bitwise" whatever this setting says. Controlled by HELION_AUTOTUNE_FP8_ACCURACY.

Settings.autotune_search_acf: list[str]#

List of PTXAS Advanced Controls Files (ACFs) to search during autotuning. ACFs are highly specialized configurations for specific hardware and use cases; when autotuning with ACFs, default -O3 is always considered. Empty list disables.

List of PTXAS config file paths to search during autotuning. Empty list (default) disables the feature. An empty string entry represents not passing a config. Controlled by HELION_AUTOTUNE_SEARCH_ACF (comma-separated paths).

Settings.autotune_rebenchmark_threshold: float | None#

If a config is within threshold*best_perf, re-benchmark it to avoid outliers. Defaults to effort profile value. Set HELION_REBENCHMARK_THRESHOLD to override.

Controls how aggressively Helion re-runs promising configs to avoid outliers. Default is 1.5 (re-benchmark anything within 1.5x of the best).

Settings.autotune_progress_bar: bool#

If True, show progress bar during autotuning. Default is True. Set HELION_AUTOTUNE_PROGRESS_BAR=0 to disable.

Toggle the interactive progress bar during autotuning. Default is True. Controlled by HELION_AUTOTUNE_PROGRESS_BAR.

Settings.autotune_config_overrides: dict[str, object]#

4}’.

Type:: Dictionary of config key/value pairs forced during autotuning. Accepts HELION_AUTOTUNE_CONFIG_OVERRIDES=’{“num_warps”

Dict of config key/value pairs to force during autotuning. Useful for disabling problematic candidates or pinning experimental options. Provide JSON via HELION_AUTOTUNE_CONFIG_OVERRIDES='{"num_warps": 4}' for global overrides.

Settings.autotune_effort: Literal['none', 'quick', 'full']#

Autotuning effort preset. One of ‘none’, ‘quick’, ‘full’.

Select the autotuning effort preset. Available values:

"none" – skip autotuning and run the default configuration.
"quick" – limited search for faster runs with decent performance. Uses from_best_available initial population strategy (no random padding).
"full" – exhaustive autotuning (current default behavior). Uses from_random initial population strategy.

Each preset also sets a default initial population strategy (see Deployment and Autotuning for details). Users can still override individual autotune_* settings; explicit values win over the preset. Controlled by HELION_AUTOTUNE_EFFORT.

Settings.autotune_best_available_max_configs: int#

Maximum number of cached configs to use for FROM_BEST_AVAILABLE initial population and for helion.from_cache() warm-start in FiniteSearch. Set HELION_BEST_AVAILABLE_MAX_CONFIGS=N to override. Default is 20.

Maximum number of cached configs to use when seeding the initial population with the from_best_available strategy. Default is 20. Controlled by HELION_BEST_AVAILABLE_MAX_CONFIGS.

Settings.autotune_best_available_max_cache_scan: int#

Maximum number of cache files to scan when searching for matching configs in FROM_BEST_AVAILABLE strategy. Set HELION_BEST_AVAILABLE_MAX_CACHE_SCAN=N to override. Default is 500.

Maximum number of cache files to scan when searching for matching configs in the from_best_available strategy. Default is 500. Controlled by HELION_BEST_AVAILABLE_MAX_CACHE_SCAN.

Autotuning Cache#

Helion stores the best-performing configs discovered during autotuning in an on-disk cache so subsequent runs can skip the search.

HELION_CACHE_DIR: Override the directory used to store cache entries. Defaults to PyTorch’s torch._inductor cache path (typically /tmp/torchinductor_$USER/helion).
HELION_SKIP_CACHE: Set to 1 to skip both reading and writing the autotuning cache. Useful for one-off experiments that should not affect cached state. To re-tune and save the result, use HELION_FORCE_AUTOTUNE=1 instead.
TRITON_STORE_BINARY_ONLY: During autotuning, Helion sets this Triton environment variable to 1 by default, skipping storage of intermediate representations (.ttir, .ttgir, .llir, etc.) and keeping only compiled binaries and metadata. This reduces Triton cache disk usage by approximately 40%. To retain IRs for debugging, set TRITON_STORE_BINARY_ONLY=0 before running.
HELION_KEEP_CACHE: Set to 1 to keep the backend compile-cache entries for all candidate configs evaluated during autotuning. By default, Helion uses an ephemeral cache directory during autotuning (TRITON_CACHE_DIR for the Triton backend, CUTE_DSL_CACHE_DIR for the CuTe backend) and only preserves the winning config’s cache entry, avoiding significant disk bloat. Enable this if you need to inspect the compiled artifacts of non-winning configs for debugging. (HELION_KEEP_TRITON_CACHE is a deprecated alias that still works for the Triton backend.)

See :class:helion.autotuner.LocalAutotuneCache for details on cache keys and behavior.

Debugging and Development#

Settings.print_output_code: bool#

If True, print the output code of the kernel to stderr.

Print generated Triton code to stderr. Default is False. Controlled by HELION_PRINT_OUTPUT_CODE=1.

Settings.print_repro: bool#

If True, print Helion kernel code, config, and caller code to stderr as a standalone repro script.

Print Helion kernel code, config, and caller code to stderr as a standalone repro script. Default is False. Controlled by HELION_PRINT_REPRO=1.

Settings.output_origin_lines: bool#

If True, annotate generated Triton code with source-origin comments. Set HELION_OUTPUT_ORIGIN_LINES=0 to disable.

Annotate generated Triton code with # src[<file>:<line>] comments indicating the originating Helion statements. Default is True. Controlled by HELION_OUTPUT_ORIGIN_LINES (set to 0 to disable).

Settings.ignore_warnings: list[type[BaseWarning]]#

Subtypes of exc.BaseWarning to ignore when compiling. Set HELION_IGNORE_WARNINGS=WarningA,WarningB (names from helion.exc) to configure via env.

List of warning types to suppress during compilation. Default is an empty list. Accepts comma-separated warning class names from helion.exc via HELION_IGNORE_WARNINGS (for example, HELION_IGNORE_WARNINGS=TensorOperationInWrapper).

Settings.debug_dtype_asserts: bool#

If True, emit tl.static_assert checks for dtype after each device node.

Emit tl.static_assert dtype checks after each lowering step. Default is False. Controlled by HELION_DEBUG_DTYPE_ASSERTS.

Device Execution Modes#

Settings.allow_warp_specialize: bool#

If True, allow warp specialization for tl.range calls on CUDA devices.

Allow warp specialization for tl.range calls. Default is True. Controlled by HELION_ALLOW_WARP_SPECIALIZE.

Settings.ref_mode: RefMode#

Reference mode for kernel execution. Can be RefMode.OFF or RefMode.EAGER.

Select the reference execution strategy. RefMode.OFF runs compiled kernels (default); RefMode.EAGER runs the interpreter for debugging. Controlled by HELION_INTERPRET.

Autotuner Hooks#

Settings.autotuner_fn: AutotunerFunction#

Function to create an autotuner. Override by passing a callable to @helion.kernel(…, autotuner_fn=…).

Override the callable that constructs autotuner instances. Accepts the same signature as helion.runtime.settings.default_autotuner_fn(). Pass a replacement callable via @helion.kernel(..., autotuner_fn=...) or helion.kernel(autotuner_fn=...) at definition time.

Settings.autotune_benchmark_fn: Optional[Callable[..., list[float]]]#

list[Callable[[], object]], *, repeat: int, desc: str | None = None) -> list[float]. If None (default), uses the built-in benchmark function.

Type:: Custom benchmark function for rebenchmarking during autotuning. Should have the following signature
Type:: (fns

Custom benchmark function for rebenchmarking during autotuning. Should have the signature (fns: list[Callable[[], object]], *, repeat: int, desc: str | None = None) -> list[float]. If None (default), uses the built-in benchmark function. Pass a replacement callable via @helion.kernel(..., autotune_benchmark_fn=...) at definition time.

Built-in values for HELION_AUTOTUNER include "LFBOTreeSearch" (default), "LFBOPatternSearch", "DESurrogateHybrid", "PatternSearch", "DifferentialEvolutionSearch", "FiniteSearch", and "RandomSearch".

Environment Variable Reference#

Environment Variable	Maps To	Description
`TRITON_F32_DEFAULT`	`dot_precision`	Sets default floating-point precision for Triton dot products (`"tf32"`, `"tf32x3"`, `"ieee"`). This variable does not apply to the Pallas backend.
`JAX_DEFAULT_MATMUL_PRECISION`	`dot_precision`	Accepts JAX matmul precision values for Pallas dot products (`"default"`, `"high"`, `"highest"`, etc., mapped to Helion `DotPrecision`). On TPU these values emit Pallas default precision. This variable does not apply to the Triton backend.
`HELION_INDEX_DTYPE`	`index_dtype`	Choose the index dtype (accepts any `torch.<dtype>` name, e.g. `int64`), or set to `auto`/unset to allow Helion to pick `int32` vs `int64` based on input sizes.
`HELION_STATIC_SHAPES`	`static_shapes`	Set to `0`/`false` to disable global static shape specialization.
`HELION_FAST_MATH`	`fast_math`	Set to `1` to enable fast math approximations (Helion-level and Inductor-level). May reduce numerical precision and change NaN/Inf behavior.
`HELION_PERSISTENT_RESERVED_SMS`	`persistent_reserved_sms`	Reserve this many streaming multiprocessors when launching persistent kernels (`0` uses all available SMs).
`HELION_FORCE_AUTOTUNE`	`force_autotune`	Force the autotuner to run even when explicit configs are provided. The result is saved to the cache.
`HELION_AUTOTUNE_FORCE_PERSISTENT`	`autotune_force_persistent`	Restrict `pid_type` to persistent kernel strategies during config search.
`HELION_DISALLOW_AUTOTUNING`	`check_autotuning_disabled`	Hard-disable autotuning; kernels must supply explicit configs when this is `1`.
`HELION_AUTOTUNE_COMPILE_TIMEOUT`	`autotune_compile_timeout`	Maximum seconds for one precompile subprocess, and for a candidate’s compile and first call in the benchmark worker (a candidate past it is recorded as failed to compile).
`HELION_AUTOTUNE_WORKER_START_TIMEOUT`	`autotune_worker_start_timeout`	Maximum seconds a fresh benchmark worker may take to start before it is killed; the start is never charged to a candidate’s compile deadline. Default: twice the compile deadline, at least `120`.
`HELION_AUTOTUNE_LOG_LEVEL`	`autotune_log_level`	Adjust logging verbosity; accepts names like `INFO` or numeric levels.
`HELION_AUTOTUNE_LOG`	`autotune_log`	Base filename for per-config CSV telemetry and mirrored autotune logs.
`HELION_AUTOTUNE_LOG_DETAILS`	`autotune_log_details`	Enable cost-model dataset sidecar (`.meta.jsonl`) with full configs map.
`HELION_AUTOTUNE_LOG_SEARCH_SPACE`	`autotune_log_search_space`	Enable search space analysis logging (default `0`). Set to `1` to enable.
`HELION_AUTOTUNE_LOG_SEARCH_SPACE_VERBOSE`	`autotune_log_search_space_verbose`	Also log each search-space restriction live at `INFO` as it is applied (default `0`). Implies `HELION_AUTOTUNE_LOG_SEARCH_SPACE`.
`HELION_AUTOTUNE_LOG_SEARCH_SPACE_PATH`	`autotune_log_search_space_path`	Optional path to save search space analysis JSON files.
`HELION_AUTOTUNE_PRECOMPILE`	`autotune_precompile`	Select the autotuner precompile mode (`"fork"` (default), `"spawn"`, or disable when empty).
`HELION_AUTOTUNE_PRECOMPILE_JOBS`	`autotune_precompile_jobs`	Cap the number of concurrent Triton precompile subprocesses.
`HELION_CUTE_RNG_STREAM`	`cute_rng_stream`	CuTe RNG policy: `auto`, `word0` or `philox4`.
`HELION_CUTE_GRAPH_CSE`	`cute_graph_cse`	Set to `0` to keep duplicate reductions (and the pure chains between them) of the traced device IR as traced on CuTe, for bisecting structural-matcher regressions. Default `1`; results are bitwise identical either way.
`HELION_CUTE_EXACT_DIVISION`	`cute_exact_division`	Set to `0` to keep one `div.rn.f32` per element for fp32 divisions by a loop-invariant scalar on CuTe instead of the hoisted bit-exact Markstein sequence. Default `1`; results are bitwise identical either way.
`HELION_CUTE_PROMOTE_HALF_ATOMIC_OUTPUTS`	`cute_promote_half_atomic_outputs`	Set to `1` to accumulate a fresh float16 `hl.atomic_add` output in float32 and cast once on return (more accurate, one extra kernel per call, differs from the Triton backend). Default `0`: add into the float16 storage with native packed half reductions.
`HELION_AUTOTUNE_RANDOM_SEED`	`autotune_random_seed`	Seed used for randomized autotuning searches.
`HELION_AUTOTUNE_MAX_GENERATIONS`	`autotune_max_generations`	Upper bound on generations for Pattern Search and Differential Evolution.
`HELION_AUTOTUNE_BUDGET_SECONDS`	`autotune_budget_seconds`	Wall-clock budget for an autotune run.
`HELION_AUTOTUNE_HANDOFF`	`autotune_handoff`	Automatically hand configuration search to a native-source agent (default off).
`HELION_AUTOTUNE_HANDOFF_BUDGET_SECONDS`	`autotune_handoff_budget_seconds`	Native-source session budget in seconds (default 1500).
`HELION_HANDOFF_AGENT`	(source agent)	Agent CLI: `codex` (default) or `claude`.
`HELION_HANDOFF_MODEL`	(source agent)	Model override; defaults to `gpt-6-astra` for Codex, `opus` for Claude.
`HELION_HANDOFF_EFFORT`	(source agent)	Reasoning effort supported by the selected CLI; defaults to `ultra` for Codex, `max` for Claude.
`HELION_AUTOTUNE_ACCURACY_CHECK`	`autotune_accuracy_check`	Toggle baseline validation for candidate configs.
`HELION_AUTOTUNE_EFFORT`	`autotune_effort`	Select autotuning preset (`"none"`, `"quick"`, `"full"`).
`HELION_AUTOTUNE_SEARCH_ACF`	`autotune_search_acf`	Comma-separated list of PTXAS config file paths to search during autotuning.
`HELION_AUTOTUNER_INITIAL_POPULATION`	(effort profile)	Override the initial population strategy (`"from_random"`, `"from_best_available"`).
`HELION_BEST_AVAILABLE_MAX_CONFIGS`	`autotune_best_available_max_configs`	Maximum cached configs to seed when using `from_best_available` strategy.
`HELION_BEST_AVAILABLE_MAX_CACHE_SCAN`	`autotune_best_available_max_cache_scan`	Maximum cache files to scan when using `from_best_available` strategy.
`HELION_REBENCHMARK_THRESHOLD`	`autotune_rebenchmark_threshold`	Re-run configs whose performance is within a multiplier of the current best.
`HELION_AUTOTUNE_FINAL_SELECTION_ROUNDS`	(effort profile `final_selection`)	Rounds of the full-effort final selection, which picks the returned config with harness-grade windows after the search (and after the finishing phase and any backend refinement): a 5 s thermal warmup, one 500+500 ms screen of the retained shortlist when it is longer than the finalists, then the best 8 survivors plus the best seed and the search’s incumbent run 4 rounds of 500+500 ms steady windows in alternating order; the median across rounds decides and a finalist that misses a round cannot win. Where the kernel’s surface names family knobs (knobs whose value picks a distinct kernel body, such as the CuTe flash pipeline family), the best shortlisted config of up to 2 families the ranking would drop, each within 1.25x of the best reading, joins the shortlist and the survivors (`F` rows of the final-selection table): the search’s short windows run at boost clocks while the steady windows run at the sustained, power-capped clock, which depends on the body’s energy per call, so a body the search ranks 1-6% behind can win the windows. This adds about 60-150 s to every full-effort autotune on every backend (the search’s own 50-200 ms windows mis-rank finalists by 3-6% on power-limited GPUs, and a window needs about 500 ms of warmup after the idle gap between windows before the B200 power limiter settles). `0` restores the single-pass final verification and its `HELION_AUTOTUNE_FINAL_REBENCHMARK_*` knobs, which the selection otherwise ignores (warned once). Under `HELION_AUTOTUNE_BUDGET_SECONDS` the search phase keeps the phase’s estimated time back plus the wall time of one search generation (the generation in flight finishes its benchmark pass after the budget check); together they take at most a quarter of the budget, so a long generation shortens the selection rather than the search, except that the selection’s smallest plan (up to four finalists: two survivors plus the seed and the incumbent, two rounds, no screen or thermal warmup, about 11-17 s) stays reserved past that generation whenever a quarter of the budget covers it (budgets under about 70 s do not, and skip the selection). The phase is sized to whatever is left when it starts: the screen goes first, then the second regime’s column, then the family quota, then rounds (down to 2), then survivors (down to 2), then the thermal warmup; it is skipped when even that does not fit.
`HELION_AUTOTUNE_TARGET_REGIME`	(effort profile `final_selection`)	The L2 regime the full-effort final selection selects by. The search times candidates in the regime of `HELION_AUTOTUNE_L2_FLUSH` (`call`: flush before every timed call, the default; `replay`: one flush, then a burst of back-to-back calls; `0`: never flushed) and the final selection in `HELION_AUTOTUNE_FINAL_L2_FLUSH` (default: the search’s regime, `replay` for a resident search). A warm regime dilutes a kernel’s cold start to a fraction of a reading, so finalists that differ only in how the first call treats its operands (an L2 eviction hint on a state load, a prefetch, a pipeline depth) tie there while they differ by 20-40% per call flushed. Unset or `auto`: the finalists tied within the policy’s `tie_tolerance` (1%) of the best median in a `replay` or `0` selection are timed again for `tiebreak_rounds` (2) rounds of the same windows in the `call` regime and the cold medians break the tie (the warm pick stands unless a tied rival is faster cold by more than the tolerance); a `call` selection takes no second column. `call` / `replay` / `0`: every finalist is timed in that regime as well (when it is not the selection’s regime already) and that column selects outright, the finalists’ recorded time being its median; set it to the regime the kernel is scored in. The final-selection table reports both columns; the second column costs `tiebreak_rounds` windows per finalist timed in it, is the second item shed under a budget, and is reserved for only where the regimes call for one. This is not `HELION_AUTOTUNE_FINAL_L2_FLUSH`, which moves the whole phase (shortlist screen, survivors and rounds) into one regime: under `TARGET_REGIME=call` on a `replay` selection the shortlist is still screened warm, so a shortlist member that is good only cold can be screened out before the `call` column sees it; `FINAL_L2_FLUSH=call` is the cheaper choice when only the per-call flushed pick matters, the target regime the one for seeing both columns.
`HELION_AUTOTUNE_PROGRESS_BAR`	`autotune_progress_bar`	Enable or disable the progress bar UI during autotuning.
`HELION_AUTOTUNE_IGNORE_ERRORS`	`autotune_ignore_errors`	Continue autotuning even when recoverable runtime errors occur.
`HELION_AUTOTUNE_CONFIG_OVERRIDES`	`autotune_config_overrides`	Supply JSON forcing particular autotuner config key/value pairs.
`TRITON_STORE_BINARY_ONLY`	Triton (autotuning)	Set to `1` during autotuning to skip Triton intermediate IRs, reducing cache size ~40%. Set to `0` to retain IRs for debugging.
`HELION_CACHE_DIR`	`LocalAutotuneCache`	Override the on-disk directory used for cached autotuning artifacts.
`HELION_SKIP_CACHE`	`LocalAutotuneCache`	When set to `1`, skip both reading and writing the autotuning cache entirely.
`HELION_ASSERT_CACHE_HIT`	`AutotuneCacheBase`	When set to `1`, require a cache hit; raises `CacheAssertionError` on cache miss with detailed diagnostics.
`HELION_AUTOTUNE_CACHE`	`autotune_cache`	Cache class to use (`"LocalAutotuneCache"` (default), `"StrictLocalAutotuneCache"`, `"RemoteAutotuneCache"`, `"StrictRemoteAutotuneCache"`, `"AOTAutotuneCache"`).
`HELION_REMOTE_CACHE_BACKEND`	(used by `RemoteAutotuneCache` and warm-start)	Fully-qualified class path to a `RemoteCacheBackend` subclass (e.g. `mypackage.cache.RedisBackend`); enables remote read-through/write-through caching and, when the backend overrides `list()`, remote warm-start lookups for `from_best_available` / `helion.from_cache`.
`HELION_BENCHMARK_CUDAGRAPH`	(benchmarking)	Wrap `run_example`’s timing loop in CUDA-graph capture/replay so launch overhead is amortized for both helion and the torch baseline. Defaults on inside `run_example`, off elsewhere. Set `0` to opt out; `1` to enable globally. Capture silently falls back to non-CG if unavailable (HIP, no CUDA, nested capture). Note: `=1` globally also applies during autotune, which is counterproductive — each of thousands of configs triggers a fresh capture (warmup + sync per config, no reuse across configs since launch parameters differ), captured graphs accumulate memory pool allocations, and many configs raise on capture due to shape/stride differences across the search space. Prefer leaving unset (default off in autotune, on in `run_example`). When it is on, kernels whose single replay takes under 64 us are timed as one captured graph holding several flushed calls minus a graph of the same flushes, because a lone replay carries the graph launch latency (4-6 us on B200, in ~2 us steps) and would read the same value for every short kernel.
`HELION_BENCHMARK_GRAPH_WARMUP`	(benchmarking)	Before the first CUDA graph a benchmark captures, drive the process into its steady graph-dispatch state (16 eager launches on side streams, kept alive; `steady_graph_dispatch`). A fresh process replays graphs with shorter inter-kernel gaps until it has launched about ten kernels on non-default streams, then every graph of the process replays slower for its lifetime, so without the warm-up the first candidates are ranked in one state and the rest in another. Runs once per device per process; a warm-up that cannot run is skipped (and retried at the next capture), never fatal. Defaults on; set `0` to measure the fast state itself (the example catalogue drivers pass this variable through to their workers). Any other value logs a warning and keeps the default.
`HELION_PRINT_OUTPUT_CODE`	`print_output_code`	Print generated Triton code to stderr for inspection.
`HELION_PRINT_REPRO`	`print_repro`	Print Helion kernel code, config, and caller code to stderr as a standalone repro script.
`HELION_OUTPUT_ORIGIN_LINES`	`output_origin_lines`	Include `# src[...]` comments in generated Triton code; set to `0` to disable.
`HELION_IGNORE_WARNINGS`	`ignore_warnings`	Comma-separated warning names defined in `helion.exc` to suppress.
`HELION_ALLOW_WARP_SPECIALIZE`	`allow_warp_specialize`	Permit warp-specialized code generation for `tl.range`.
`HELION_DEBUG_DTYPE_ASSERTS`	`debug_dtype_asserts`	Inject dtype assertions after each lowering step.
`HELION_INTERPRET`	`ref_mode`	Run kernels through the reference interpreter when set to `1` (maps to `RefMode.EAGER`).
`HELION_AUTOTUNER`	`default_autotuner_fn`	Select which autotuner implementation to instantiate. Default is `"LFBOTreeSearch"`. Other options: `"LFBOPatternSearch"`, `"DESurrogateHybrid"`, `"PatternSearch"`, `"DifferentialEvolutionSearch"`, `"FiniteSearch"`, `"RandomSearch"`, `"LLMGuidedSearch"`, `"LLMSeededSearch"`, `"LLMSeededLFBOTreeSearch"`. See :doc:`autotuner` for LLM configuration.
`HELION_BACKEND`	`backend`	Code generation backend (`"triton"` (default), `"pallas"`, `"cute"`, `"tileir"`). `"cute"` requires PyTorch built against CUDA 13 or later.

Settings#

Overview#

Key Characteristics#

Settings vs Config#

Configuration Examples#

Using Environment Variables#

Using Decorator Arguments#

Settings Reference#

Core Compilation Settings#

Autotuning Settings#

Autotuning Cache#

Debugging and Development#

Device Execution Modes#

Autotuner Hooks#

Environment Variable Reference#

See Also#