Rate this Page

★ ★ ★ ★ ★

helion.Settings#

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

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) –

Methods

`__init__`(**settings)	Initialize the Settings object with the provided dictionary of settings.
`check_autotuning_disabled`()	rtype: `None`
`copy`(**overrides)	Copy values and origins, marking supplied structural overrides explicit.
`get_cute_structural_policy`()	Snapshot effective booleans without resolving or activating a policy.
`get_rebenchmark_threshold`()	Get the effective rebenchmark threshold.
`get_suspicious_rebenchmark_ratio`()	rtype: `float` \| `None`
`to_dict`()	Convert the Settings object to a dictionary.

Attributes

`backend`	Code generation backend.
`ignore_warnings`	Subtypes of exc.BaseWarning to ignore when compiling.
`index_dtype`	The dtype to use for index variables.
`dot_precision`	Precision for dot products.
`cute_rng_stream`	auto (CuTe default) uses philox4 for explicit hl.rand and preserves other RNG operations.
`fast_math`	If True, enable fast math approximations (Helion-level, including the CuTe flash-attention polynomial exp2 packets, and Inductor-level).
`cute_flash_env_overrides`	HELION_CUTE_FLASH_* variables the kernel resolves as if set in the environment.
`static_shapes`	If True, use static shapes for all tensors.
`persistent_reserved_sms`	Number of streaming multiprocessors to reserve when launching persistent kernels.
`autotune_force_persistent`	If True, restrict pid_type choices to persistent kernels only during config selection.
`distributed`	Force distributed compilation behavior (persistent-only PID types, signal-pad SM limits, and process-group resolution).
`autotune_log_level`	Log level for autotuning using Python logging levels.
`autotune_log`	Base filename for autotune logs.
`autotune_log_details`	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).
`autotune_handoff`	Hand configuration search to a native-source agent automatically.
`autotune_handoff_budget_seconds`	Wall-clock budget for the native-source agent session, excluding search and export.
`autotune_compile_timeout`	a candidate past it is killed and recorded as failed to compile, and the search goes on.
`autotune_benchmark_subprocess`	Run the autotune benchmark phase in a long-lived spawn subprocess so a hung/slow kernel can be killed without losing autotune progress.
`autotune_benchmark_timeout`	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).
`autotune_worker_start_timeout`	twice the worker's compile deadline, at least 120 seconds.
`autotune_precompile`	'fork', 'spawn', or falsy/None to disable.
`autotune_precompile_jobs`	Maximum concurrent Triton precompile processes, default to cpu count.
`autotune_random_seed`	Seed used for autotuner random number generation.
`autotune_best_of_k`	When > 1, run autotuning K times with different random seeds and keep the config with the best benchmark performance.
`autotune_accuracy_check`	If True, validate candidate configs against the baseline kernel output before accepting them during autotuning.
`autotune_rebenchmark_threshold`	If a config is within threshold*best_perf, re-benchmark it to avoid outliers.
`autotune_suspicious_rebenchmark_ratio`	When subprocess benchmarking is enabled, recheck rebenchmark timings below ratio*previous_timing before accepting them.
`autotune_search_acf`	List of PTXAS Advanced Controls Files (ACFs) to search during autotuning.
`autotune_progress_bar`	If True, show progress bar during autotuning.
`autotune_max_generations`	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`	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.
`autotune_ignore_errors`	If True, skip logging and raising autotune errors.
`autotune_adaptive_timeout`	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.
`pallas_collective_id`	Optional Pallas collective namespace override.
`pallas_interpret`	If True, run Pallas kernels in interpret mode on CPU (no TPU needed).
`pallas_topk_recall_target`	Recall target for the Pallas approximate top-k lowering.
`cute_region_fission`	If True, distribute proven row-independent materialized regions into ordered CuTe launches with independently tuned tile axes.
`cute_materialize_transformed_operands`	If True, materialize proven pure reused matrix operands before a CuTe GEMM, with independent producer and consumer tile axes.
`cute_full_slice_matmul_tiling`	If True, give eligible direct full-slice half/BF16 matmuls an explicit, tunable contraction tile loop on CuTe.
`cute_segmented_matmul_tiling`	If True, expose proven offset-delimited independent matrix rows to the CuTe grouped scheduler.
`cute_flatten_nested_reductions`	If True, flatten proven FP32 joint jagged/feature reductions and independent fresh-output stores on CuTe.
`cute_graph_cse`	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).
`cute_stage_unaligned_operands`	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.
`cute_flat_edge_tma`	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.
`cute_exact_division`	the IEEE round-to-nearest quotient bit for bit, without one div.rn expansion per element.
`cute_promote_half_atomic_outputs`	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).
`triton_do_not_specialize`	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.
`print_output_code`	If True, print the output code of the kernel to stderr.
`print_repro`	If True, print Helion kernel code, config, and caller code to stderr as a standalone repro script.
`output_origin_lines`	If True, annotate generated Triton code with source-origin comments.
`force_autotune`	If True, force autotuning even if a config is provided.
`autotune_config_overrides`	4}'.
`autotune_seed_configs`	A Config or sequence of Configs to seed the autotuner initial population without constraining the search space.
`disable_autotuner_heuristics`	If True, disable compiler seed generation and promotion.
`allow_warp_specialize`	If True, allow warp specialization for tl.range calls on CUDA devices.
`debug_dtype_asserts`	If True, emit tl.static_assert checks for dtype after each device node.
`ref_mode`	Reference mode for kernel execution.
`autotuner_fn`	Function to create an autotuner.
`autotune_effort`	Autotuning effort preset.
`autotune_baseline_fn`	Custom baseline function for computing baseline output during autotuning.
`autotune_baseline_atol`	Absolute tolerance for baseline output comparison during autotuning accuracy checks.
`autotune_baseline_rtol`	Relative tolerance for baseline output comparison during autotuning accuracy checks.
`autotune_fp8_accuracy`	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.
`autotune_baseline_accuracy_check_fn`	object, expected: object) -> None.
`autotune_cache`	The name of the autotuner cache class to use.
`autotune_benchmark_fn`	list[Callable[[], object]], *, repeat: int, desc: str \| None = None) -> list[float]. If None (default), uses the built-in benchmark function.
`autotune_best_available_max_configs`	Maximum number of cached configs to use for FROM_BEST_AVAILABLE initial population and for helion.from_cache() warm-start in FiniteSearch.
`autotune_best_available_max_cache_scan`	Maximum number of cache files to scan when searching for matching configs in FROM_BEST_AVAILABLE strategy.
`autotune_initial_population_strategy`	'from_random', 'from_best_available'.
`torch_compile_fusion`	If True, allow torch.compile to fuse this Helion kernel with surrounding Inductor ops (prologue/epilogue) when used inside torch.compile.
`autotune_config_filter`	Optional callable `(config: Config) -> Config \| None` that the autotuner calls on every candidate config before compiling or benchmarking it.
`autotune_with_torch_compile_fusion`	If True, autotuning benchmarks the fused kernel (with epilogue/prologue) to pick configs optimal for the actual fused workload.
`autotune_log_search_space`	If True, log search space analysis after autotuning including which features were enabled/disabled, total search space size, and coverage metrics.
`autotune_log_search_space_verbose`	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.
`autotune_log_search_space_path`	Optional path to save search space analysis JSON.
`cute_structural_origins`	Immutable origins captured at construction and explicit writes.

__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.