Rate this Page
★ ★ ★ ★ ★

Config#

The Config class represents kernel optimization parameters that control how Helion kernels are compiled and executed.

class helion.Config(*, block_sizes=None, num_threads=None, loop_orders=None, flatten_loops=None, l2_groupings=None, reduction_loops=None, range_unroll_factors=None, range_warp_specializes=None, range_num_stages=None, range_multi_buffers=None, range_flattens=None, static_ranges=None, pallas_load_buffer_count=None, load_eviction_policies=None, load_cache_modifiers=None, store_cache_modifiers=None, cute_store_policy=None, cute_exact_div_form=None, cute_async_load_stages=None, cute_async_load_lookahead=None, cute_async_load_group_rows=None, cute_async_load_cache=None, cute_async_store_policy=None, cute_bf16x2_recurrence=None, cute_signed_bitfield_bf16=None, cute_proven_bounds=None, cute_affine_scan_schedule=None, cute_rng_packet=None, cute_independent_reduction=None, cute_replicated_reduction=None, cute_vector_packet_unroll=None, cute_vloop_sink=None, cute_lane_unroll=None, cute_wide_packet_hints=None, cute_pdl=None, cute_pdl_release=None, cute_pdl_spread=None, cute_warp_aggregated_atomics=None, cute_packet_prefetch=None, cute_reduction_pipeline_depth=None, cute_host_paired_sum=None, num_warps=None, num_stages=None, pid_type=None, cross_loop_pipeline=None, num_sm_multiplier=None, maxnreg=None, host_tensor_descriptors=None, indexing=None, atomic_indexing=None, advanced_controls_file=None, epilogue_subtile=None, xcd_remap=None, reduction_fusion=None, **kwargs)[source]#

Bases: Mapping[str, object]

Parameters:
__init__(*, block_sizes=None, num_threads=None, loop_orders=None, flatten_loops=None, l2_groupings=None, reduction_loops=None, range_unroll_factors=None, range_warp_specializes=None, range_num_stages=None, range_multi_buffers=None, range_flattens=None, static_ranges=None, pallas_load_buffer_count=None, load_eviction_policies=None, load_cache_modifiers=None, store_cache_modifiers=None, cute_store_policy=None, cute_exact_div_form=None, cute_async_load_stages=None, cute_async_load_lookahead=None, cute_async_load_group_rows=None, cute_async_load_cache=None, cute_async_store_policy=None, cute_bf16x2_recurrence=None, cute_signed_bitfield_bf16=None, cute_proven_bounds=None, cute_affine_scan_schedule=None, cute_rng_packet=None, cute_independent_reduction=None, cute_replicated_reduction=None, cute_vector_packet_unroll=None, cute_vloop_sink=None, cute_lane_unroll=None, cute_wide_packet_hints=None, cute_pdl=None, cute_pdl_release=None, cute_pdl_spread=None, cute_warp_aggregated_atomics=None, cute_packet_prefetch=None, cute_reduction_pipeline_depth=None, cute_host_paired_sum=None, num_warps=None, num_stages=None, pid_type=None, cross_loop_pipeline=None, num_sm_multiplier=None, maxnreg=None, host_tensor_descriptors=None, indexing=None, atomic_indexing=None, advanced_controls_file=None, epilogue_subtile=None, xcd_remap=None, reduction_fusion=None, **kwargs)[source]#

Initialize a Config object.

Parameters:

  • block_sizes (list[int] | None) – Controls tile sizes for hl.tile invocations.

  • num_threads (list[int] | int | None) – Target thread count per axis (backend-specific).

  • loop_orders (list[list[int]] | None) – Permutes iteration order of tiles.

  • l2_groupings (list[int] | None) – Reorders program IDs for L2 cache locality.

  • reduction_loops (list[int | None] | None) – Configures reduction loop behavior.

  • range_unroll_factors (list[int] | None) – Loop unroll factors for tl.range calls.

  • range_warp_specializes (list[bool | None] | None) – Warp specialization for tl.range calls.

  • range_num_stages (list[int] | None) – Number of stages for tl.range calls.

  • range_multi_buffers (list[bool | None] | None) – Controls disallow_acc_multi_buffer for tl.range calls.

  • range_flattens (list[bool | None] | None) – Controls flatten parameter for tl.range calls.

  • static_ranges (list[bool] | None) – Whether to use tl.static_range instead tl.range.

  • pallas_load_buffer_count (list[int] | None) – Pallas-only load buffer count (1 or 2) for each input tensor. Tensors without an existing DMA route use the ordinary path.

  • load_eviction_policies (Union[Literal['', 'first', 'last'], list[Literal['', 'first', 'last']], None]) – Eviction policies for load operations. A single value applies to every load; a list specifies one value per load. Valid values are “”, “first”, and “last”. CuTe also supports “nc_l2_first” for proven read-only 16-byte or 32-byte packets.

  • load_cache_modifiers (list[Literal['', '.cg']] | None) – Cache modifiers for load operations (“”, “.cg”).

  • store_cache_modifiers (list[Literal['', '.cs', '.wt']] | None) – Cache modifiers for store operations (“”, “.cs”, “.wt”).

  • cute_store_policy (Optional[Literal['default', 'streaming', 'l2_evict_last']]) – Cache policy for proven full aligned 16-byte CuTe global store packets (“default”, “streaming” = evict-first st.global.cs for outputs written once, or “l2_evict_last”).

  • cute_async_load_stages (int | None) – Shared-memory ring stages for eligible CuTe in-place 16-byte state loads. Zero disables the transformation.

  • cute_async_load_lookahead (int | None) – Async-copy groups kept ahead of compute.

  • cute_async_load_group_rows (int | None) – Per-thread row iterations in each group.

  • cute_exact_div_form (Optional[Literal['redo', 'plain']]) – How CuTe lowers an fp32 division by a loop-invariant divisor (a row sum, a scalar argument); both forms give the IEEE quotient bit for bit. "redo" (default) speculates the Markstein sequence over the loop and re-runs the loop with div.rn on a hard case (a loop the redo cannot take keeps its plain division), "plain" keeps one div.rn per element.

  • cute_async_load_cache (Optional[Literal['cg', 'ca']]) – PTX cp.async cache policy (“cg” or “ca”).

  • cute_async_store_policy (Optional[Literal['default', 'l2_evict_last']]) – Cache policy for the matching in-place 16-byte state store (“default” or “l2_evict_last”).

  • cute_bf16x2_recurrence (bool | None) – Pack a structurally proven BF16 rank-one recurrence into native BF16x2 operations.

  • cute_signed_bitfield_bf16 (bool | None) – Convert proved signed byte fields to BF16 using the existing vector load/store packets. Disabled by default.

  • cute_proven_bounds (bool | None) – Remove CuTe index guards only when exact launch dimensions and cache-specialized tensor sizes prove them true.

  • cute_affine_scan_schedule (Optional[Literal['ordinary', 'direct_m16n8_v1', 'direct_m16n16_v1']]) – Physical schedule for a compatible affine scan. "ordinary" disables the direct lowering; "direct_m16n8_v1" and "direct_m16n16_v1" select the measured direct schedule profiles.

  • cute_independent_reduction (bool | None) – Sum each vector lane independently in FP32 before combining lanes, preserving the element expression. This changes summation order. Disabled by default.

  • cute_replicated_reduction (bool | None) – Replicate single-use whole-CTA sums in each warp using one shared-memory barrier. Disabled by default.

  • cute_vector_packet_unroll (bool | None) – Fully unroll small straight-line vector packet loops. Disabled by default to retain compact code.

  • cute_vloop_sink (bool | None) – Sink a grid constexpr vector loop into the serial reduction nest it wraps (column sums): one vector load per row, per-lane register accumulators, one grouped cross-thread combine per row tile. Disabled by default. Kernels where nothing can be sunk generate exactly the knob-off code.

  • cute_lane_unroll (int | None) – Unroll factor (1, 2, 4, 8, 16) for the row lane loop of a sunk vector nest; every copy’s vector loads issue before any copy’s accumulates. Only applies with cute_vloop_sink.

  • cute_wide_packet_hints (bool | None) – Carry the l2_last / l1_l2_first / l1_l2_last load eviction policies onto 32-byte CuTe packets (cute_vector_widths 8 for fp32, 16 for bf16/fp16) through one 256-bit cache-hinted load; SM100+ with CUDA 13 or newer, elsewhere inert. Off by default: a hinted 32-byte packet then loads without the hint, as it always has.

  • cute_pdl (bool | None) – Launch the kernel as a programmatic dependent of the previous kernel in the stream: its launch and prologue overlap that kernel’s tail and it waits (griddepcontrol.wait) ahead of its first global memory access; when its own host zero fills launch as dependent fill kernels, ahead of its first access to a filled tensor instead (the fill kernel releases only after its own wait, so every earlier kernel has completed), which lets a split-K kernel’s operand loads overlap the fill of its atomic output. Off by default. Kernels that already take part in dependent launch (native plans launched with it, bodies with their own wait or release, the materialized-fission producer) and pre-sm_90 targets generate exactly the knob-off code.

  • cute_pdl_release (Optional[Literal['auto', 'entry', 'none', 'exit', 'last_wave']]) – Where a cute_pdl kernel releases its own dependents (griddepcontrol.launch_dependents), which is what lets the next kernel of a sequence launch early too. entry releases right after the wait, so the dependents’ CTAs become resident while this kernel runs; exit releases after the body’s last statement; none leaves the implicit trigger at grid completion, after which a dependent starts later than a plain stream edge would have launched it; last_wave releases at entry in the CTAs of the launch’s last wave only (the CTAs whose dispatch-order index lies within the resident capacity of the grid’s end; the earlier CTAs count as released when they exit), so the dependents launch as the last wave starts and fill the slots its tail frees, which is plain entry for a grid that fits in one wave and keeps the middle waves of a larger grid undisturbed. auto (the default) picks entry while the launch grid is known to fit in one wave of resident CTAs (the SM count times the CTAs an SM holds, from the launch block or cute_min_blocks_per_mp) and last_wave past it or when the grid is sized at runtime; the search leaves the placement at auto. Every placement reads the same data: a dependent’s own wait blocks until this grid has completed and flushed. On the tcgen05 GEMM kernels, whose lowering waits in its warp roles and releases after the last MMA, entry and last_wave add the release right after each of those waits and exit / none keep the late release alone. Inert when cute_pdl is off.

  • cute_pdl_spread (bool | None) – Pad a cute_pdl kernel’s shared memory so that at most ceil(CTAs / SMs) of its CTAs fit on one SM, its density on an idle device. A programmatic dependent launches while its primary runs and the hardware dispatches its CTAs onto whatever SMs have room; when the primary’s CTAs leave none beside them, a small dependent’s CTAs all land on the few SMs the primary left idle and run packed there once it completes. With the pad the CTAs beyond the idle SMs wait for the primary’s to retire and take their SMs one each. Off by default; inert when cute_pdl is off, when the kernel’s own allocation already holds its CTAs to that density, when the grid is sized at runtime, and on the families that own their dependent launch.

  • cute_warp_aggregated_atomics (bool | None) – Combine relaxed Int32 zero/one ticket reservations for equal addresses within each active warp. Each lane retains its own valid fetch-add return value. Off by default.

  • cute_packet_prefetch (int | None) – Prefetch 2, 4, or 8 independent vector packets in a proved complete tile; 0 (the default) keeps the original order. Requires cute_proven_bounds and storage-disjointness proof.

  • cute_reduction_pipeline_depth (int | None) – Shared-memory ring depth (2 or 4) for the proved CuTe pipelined resident-reduction schedule. Defaults to 2. Depth 4 prefetches three row groups ahead and requires enough shared memory for all four slots and reduction scratch.

  • cute_host_paired_sum (Optional[Literal['off', 'mapped', 'narrow']]) – Fuse a proved host FP32 sum(0)/cast pair or terminal sum/cast. "off" (the default) retains Torch calls; "mapped" and "narrow" change independent-column ownership while preserving the audited FP32 sum tree. Unknown Torch implementations or unsupported bindings use the original operations. Available only for typed, independent host pairs or a terminal sum/cast with an effect-free return prefix.

  • num_warps (int | None) – Number of warps per block.

  • num_stages (int | None) – Number of stages for software pipelining.

  • pid_type (Optional[Literal['flat', 'xyz', 'persistent_blocked', 'persistent_interleaved']]) – Program ID type strategy (“flat”, “xyz”, “persistent_blocked”, “persistent_interleaved”).

  • cross_loop_pipeline (Optional[Literal['barrier', 'static', 'dynamic']]) – Execution strategy for kernels with compiler-inferred cross-loop dependencies. "barrier" uses grid synchronization. "static" and "dynamic" (Triton only) execute the same compiler-derived dependency schedule with fixed worker ownership or one-shot packet dispatch, respectively. Unsupported kernels reject this field.

  • num_sm_multiplier (int | None) – Positive integer multiplier for the number of SMs in persistent kernels. The autotuner searches powers of two, but explicit configurations may use intermediate values. Controls multi-occupancy by launching N * num_sms thread blocks instead of just num_sms.

  • maxnreg (int | None) – None or a positive integer maximum number of registers per thread (1–256). The autotuner searches None, 32, 64, 128, and 256, but explicit configurations may use intermediate values. Lower values allow higher occupancy but may hurt performance. Used with persistent kernels to ensure multi-occupancy can be achieved.

  • host_tensor_descriptors (bool | None) – Construct tensor descriptors once in the host launcher instead of independently in every device program. This only affects operations using tensor-descriptor indexing.

  • indexing (Union[Literal['pointer', 'tensor_descriptor', 'block_ptr'], list[Literal['pointer', 'tensor_descriptor', 'block_ptr']], None]) –

    Indexing strategy for load and store operations. Can be: - A single strategy string (all loads/stores use this strategy):

    indexing=”block_ptr” # backward compatible

    • A list of strategies (one per load/store operation, must specify all): indexing=[“pointer”, “block_ptr”, “tensor_descriptor”]

    • Empty/omitted (all loads/stores default to “pointer”)

    Valid strategies: “pointer”, “tensor_descriptor”, “block_ptr” (“block_ptr” needs Triton < 3.9; newer Triton lowers it as “pointer”)

  • atomic_indexing (Union[Literal['pointer', 'tensor_descriptor', 'block_ptr'], list[Literal['pointer', 'tensor_descriptor', 'block_ptr']], None]) – Indexing strategy for atomic operations (e.g., hl.atomic_add). Same format as indexing (a single string or a list per atomic op). Defaults to “pointer” when omitted.

  • advanced_controls_file (str | None) – Path to a PTXAS control file applied during compilation, or empty string for none.

  • epilogue_subtile (int | None) – Split factor for the epilogue (post-matmul pointwise + store) along the N dimension. None = disabled (default), valid values are 2 or 4.

  • xcd_remap (bool | None) – AMD CDNA only. Remap program IDs into contiguous per-XCD regions to improve L2 locality on multi-XCD GPUs (MI300/MI350). Supported for pid_type “flat”, “persistent_blocked”, and “persistent_interleaved”; composes with l2_groupings.

  • reduction_fusion (bool | None) – When multiple independent reductions in a looped (rolled) reduction read the same operand, fold them into one loop with a single shared load instead of re-reading the operand per reduction. Trades global-memory bandwidth for register pressure, so it is autotuned. Defaults to False.

  • **kwargs (object) – Additional user-defined configuration parameters.

  • flatten_loops (list[bool] | None) –

  • cute_rng_packet (bool | None) –

config: dict[str, object]#
to_json()[source]#

Convert the config to a JSON string.

Return type:

str

classmethod from_dict(config_dict)[source]#

Create a Config from a plain dictionary.

Parameters:

config_dict (Mapping[str, object]) –

Return type:

Config

classmethod from_json(json_str)[source]#

Create a Config object from a JSON string.

Parameters:

json_str (str) –

Return type:

Config

minimize(config_spec)[source]#

Return a new Config with values matching effective defaults removed.

This produces a minimal config representation by removing any values that match what the config_spec would use as defaults.

Parameters:

config_spec (object) – The ConfigSpec that defines the defaults for this kernel.

Return type:

Config

Returns:

A new Config with default values removed.

save(path)[source]#

Save the config to a JSON file.

Parameters:

path (str | Path) –

Return type:

None

classmethod load(path)[source]#

Load a config from a JSON file.

Parameters:

path (str | Path) –

Return type:

Config

property block_sizes: list[int]#
property loop_orders: list[list[int]]#
property num_threads: list[int]#
property flatten_loops: list[bool]#
property reduction_loops: list[int | None]#
property num_warps: int#
property num_stages: int#
property l2_groupings: list[int]#
property pid_type: Literal['flat', 'xyz', 'persistent_blocked', 'persistent_interleaved']#
property cross_loop_pipeline: Literal['barrier', 'static', 'dynamic']#
property xcd_remap: bool#
property num_sm_multiplier: int#
property maxnreg: int | None#
property host_tensor_descriptors: bool#
property reduction_fusion: bool#
property range_unroll_factors: list[int]#
property advanced_controls_file: str#
property range_warp_specializes: list[bool | None]#
property range_num_stages: list[int]#
property range_multi_buffers: list[bool | None]#
property range_flattens: list[bool | None]#
property static_ranges: list[bool]#
property pallas_load_buffer_count: list[int]#
property load_eviction_policies: Literal['', 'first', 'last'] | list[Literal['', 'first', 'last']]#
property load_cache_modifiers: list[Literal['', '.cg']]#
property store_cache_modifiers: list[Literal['', '.cs', '.wt']]#
property cute_async_load_stages: int#
property cute_async_load_lookahead: int#
property cute_async_load_group_rows: int#
property cute_async_load_cache: Literal['cg', 'ca']#
property cute_async_store_policy: Literal['default', 'l2_evict_last']#
property cute_store_policy: Literal['default', 'streaming', 'l2_evict_last']#
property cute_wide_packet_hints: bool#
property cute_exact_div_form: Literal['redo', 'plain']#
property cute_bf16x2_recurrence: bool#
property cute_proven_bounds: bool#
property cute_affine_scan_schedule: Literal['ordinary', 'direct_m16n8_v1', 'direct_m16n16_v1']#
property cute_packet_prefetch: int#
property cute_reduction_pipeline_depth: int#
property indexing: Literal['pointer', 'tensor_descriptor', 'block_ptr'] | list[Literal['pointer', 'tensor_descriptor', 'block_ptr']]#
property atomic_indexing: Literal['pointer', 'tensor_descriptor', 'block_ptr'] | list[Literal['pointer', 'tensor_descriptor', 'block_ptr']]#
property epilogue_subtile: int | None#

Overview#

Config objects specify optimization parameters that control how Helion kernels run on the hardware.

Key Characteristics#

  • Performance-focused: Control GPU resource allocation, memory access patterns, and execution strategies

  • Autotuned: The autotuner searches through different Config combinations to find optimal performance

  • Kernel-specific: Each kernel can have different optimal Config parameters based on its computation pattern

  • Hardware-dependent: Optimal configs vary based on GPU architecture and problem size

Config vs Settings#

Aspect

Config

Settings

Purpose

Control execution performance

Control compilation behavior

Autotuning

✅ Automatically optimized

❌ Never autotuned

Examples

block_sizes, num_warps, indexing

print_output_code, print_repro, autotune_effort

When to use

Performance optimization

Development, debugging, environment setup

Configs are typically discovered automatically through autotuning, but can also be manually specified for more control.

Configuration Parameters#

Block Sizes and Resources#

Config.block_sizes#

List of tile sizes for hl.tile() loops. Each value controls the number of elements processed per GPU thread block for the corresponding tile dimension.

Config.reduction_loops#

Configuration for reduction operations within loops.

Config.num_warps#

Number of warps (groups of 32 threads) per thread block. Higher values increase parallelism but may reduce occupancy.

Config.num_stages#

Number of pipeline stages for software pipelining. Higher values can improve memory bandwidth utilization.

Loop Optimizations#

Config.loop_orders#

Permutation of loop iteration order for each hl.tile() loop. Used to optimize memory access patterns.

Config.flatten_loops#

Whether to flatten nested loops for each hl.tile() invocation.

Config.range_unroll_factors#

Unroll factors for tl.range loops in generated Triton code.

Config.range_warp_specializes#

Whether to enable warp specialization for tl.range loops.

Config.range_num_stages#

Number of pipeline stages for tl.range loops.

Config.range_multi_buffers#

Controls disallow_acc_multi_buffer parameter for tl.range loops.

Config.range_flattens#

Controls flatten parameter for tl.range loops.

Config.static_ranges#

Whether to use tl.static_range instead of tl.range.

Config.reduction_fusion#

Fold multiple independent looped reductions that read the same operand into one loop with a single shared load. Trades global-memory bandwidth for register pressure, so it is autotuned.

Execution and Indexing#

Config.pid_type#

Program ID layout strategy:

  • "flat": Standard linear program ID assignment

  • "xyz": 3D program ID layout

  • "persistent_blocked": Persistent kernels with blocked work distribution

  • "persistent_interleaved": Persistent kernels with interleaved distribution

Config.l2_groupings#

Controls reordering of program IDs to improve L2 cache locality.

Config.xcd_remap#

AMD CDNA only. Remaps program IDs into contiguous per-XCD regions to improve L2 locality on multi-XCD GPUs (MI300/MI350). Supported for pid_type "flat" (remap the program id), "persistent_blocked" (remap the worker’s block assignment), and "persistent_interleaved" (remap each virtual pid); composes with l2_groupings. Best for tiled-reuse kernels (matmul, MoE, attention); not for simple rowwise streaming kernels.

Config.indexing#

Memory indexing strategy for load and store operations. Can be specified as:

Single strategy (applies to all loads and stores - backward compatible):

indexing="block_ptr"  # All loads and stores use block pointers

Per-operation strategies (list, one per load/store in execution order):

# 2 loads + 1 store = 3 indexing strategies
indexing=["pointer", "pointer", "block_ptr"]  # loads use pointer, store uses block_ptr

Empty/omitted (defaults to "pointer" for all operations):

# indexing not specified - all loads and stores use pointer indexing

Valid strategies:

  • "pointer": Pointer-based indexing (default)

  • "tensor_descriptor": Tensor descriptor indexing (requires Hopper+ GPU)

  • "block_ptr": Block pointer indexing (Triton < 3.9 only; newer Triton removed block pointers, so this value is lowered as "pointer" and the autotuner no longer searches it)

Note

When using a list, provide one strategy for each load and store operation in the order they appear in the kernel. The indexing list is ordered as: [load1, load2, ..., loadN, store1, store2, ..., storeM]

Epilogue Optimization#

Config.epilogue_subtile#

Split factor for the epilogue (pointwise ops + store) along a tile dimension. Splits the store from [BLOCK_M, BLOCK_N] into SUBTILE_FACTOR × [BLOCK_M, BLOCK_N / SUBTILE_FACTOR], reducing the accumulator shared-memory footprint and enabling extra pipeline stages.

Valid values:

  • None: Disabled (default)

  • 2: Split epilogue into 2 sub-tiles

  • 4: Split epilogue into 4 sub-tiles (when K ≥ 16384)

Requirements:

  • Blackwell (sm_100+) GPU with tensor descriptor support

  • Automatically discovered by the autotuner when the K dimension is ≥ 1024

Interactions:

  • Incompatible with flatten_loops=True

  • Forces store indexing to "tensor_descriptor"

See the epilogue subtiling example for usage patterns.

Memory and Caching#

Config.load_eviction_policies#

Eviction policies for load operations issued from device loops. Provide one policy per hl.load site discovered in the kernel. Allowed values:

  • "": No eviction policy (omitted)

  • "first": Maps to Triton eviction_policy='evict_first'

  • "last": Maps to Triton eviction_policy='evict_last'

Notes:

  • The number of entries must match the number of load sites considered tunable by the kernel.

  • An explicit eviction_policy=... argument passed to hl.load overrides this config.

Usage Examples#

Manual Config Creation#

import torch
import helion
import helion.language as hl

# Create a specific configuration
config = helion.Config(
    block_sizes=[64, 32],      # 64 elements per tile in dim 0, 32 in dim 1
    num_warps=8,               # Use 8 warps (256 threads) per block
    num_stages=4,              # 4-stage pipeline
    pid_type="xyz"             # Use 3D program ID layout
)

# Use with kernel
@helion.kernel(config=config)
def my_kernel(x: torch.Tensor) -> torch.Tensor:
    result = torch.zeros_like(x)
    for i, j in hl.tile(x.shape):
        result[i, j] = x[i, j] * 2
    return result

Eviction Policy Example#

import torch
import helion
import helion.language as hl

@helion.kernel(
    config={
        "block_size": 16,
        "load_eviction_policies": ["", "last"],  # second load uses evict_last
    }
)
def kernel_with_eviction(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
    out = torch.empty_like(x)
    for tile in hl.tile(x.size(0)):
        a = hl.load(x, [tile])                 # No eviction policy
        b = hl.load(y, [tile])                 # Will use evict_last from config
        out[tile] = a + b
    return out

# Explicit policy on hl.load overrides config:
# hl.load(x, [tile], eviction_policy="evict_first")

Per-Load Indexing Example#

import torch
import helion
import helion.language as hl

# Single indexing strategy for all loads and stores (backward compatible)
@helion.kernel(config={"indexing": "block_ptr"})
def kernel_uniform_indexing(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
    out = torch.empty_like(x)
    for tile in hl.tile(x.size(0)):
        a = hl.load(x, [tile])  # Load: uses block_ptr
        b = hl.load(y, [tile])  # Load: uses block_ptr
        out[tile] = a + b       # Store: uses block_ptr
    return out

# Per-operation indexing strategies for fine-grained control
# Indexing list is ordered: [load1, load2, ..., store1, store2, ...]
@helion.kernel(
    config={
        "block_size": 16,
        "indexing": ["pointer", "pointer", "block_ptr"],  # 2 loads + 1 store
    }
)
def kernel_mixed_indexing(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
    out = torch.empty_like(x)
    for tile in hl.tile(x.size(0)):
        a = hl.load(x, [tile])  # First load: pointer indexing
        b = hl.load(y, [tile])  # Second load: pointer indexing
        out[tile] = a + b       # Store: block_ptr indexing
    return out

Config Serialization#

# Save config to file
config.save("my_config.json")

# Load config from file
loaded_config = helion.Config.load("my_config.json")

# JSON serialization
config_dict = config.to_json()
restored_config = helion.Config.from_json(config_dict)

Autotuning with Restricted Configs#

# Restrict autotuning to specific configurations
configs = [
    helion.Config(block_sizes=[32, 32], num_warps=4),
    helion.Config(block_sizes=[64, 16], num_warps=8),
    helion.Config(block_sizes=[16, 64], num_warps=4),
]

@helion.kernel(configs=configs)
def matrix_multiply(a: torch.Tensor, b: torch.Tensor) -> torch.Tensor:
    m, k = a.size()
    k2, n = b.size()
    assert k == k2, f"size mismatch {k} != {k2}"
    out = torch.empty([m, n], dtype=a.dtype, device=a.device)

    for tile_m, tile_n in hl.tile([m, n]):
        acc = hl.zeros([tile_m, tile_n], dtype=torch.float32)
        for tile_k in hl.tile(k):
            acc = torch.addmm(acc, a[tile_m, tile_k], b[tile_k, tile_n])
        out[tile_m, tile_n] = acc

    return out

See Also#

On this page
Edit on GitHub
Show Source
PyTorch Libraries