Rate this Page
★ ★ ★ ★ ★

helion.Config#

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_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_proven_bounds=None, cute_affine_scan_schedule=None, num_warps=None, num_stages=None, pid_type=None, cross_loop_pipeline=None, num_sm_multiplier=None, maxnreg=None, indexing=None, atomic_indexing=None, advanced_controls_file=None, epilogue_subtile=None, xcd_remap=None, **kwargs)[source]#
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_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_proven_bounds=None, cute_affine_scan_schedule=None, num_warps=None, num_stages=None, pid_type=None, cross_loop_pipeline=None, num_sm_multiplier=None, maxnreg=None, indexing=None, atomic_indexing=None, advanced_controls_file=None, epilogue_subtile=None, xcd_remap=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”.

  • 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_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_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_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.

  • 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" 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 (Optional[Literal[1, 2, 4, 8]]) – Multiplier for the number of SMs in persistent kernels (1, 2, 4, 8). Controls multi-occupancy by launching N * num_sms thread blocks instead of just num_sms.

  • maxnreg (Optional[Literal[32, 64, 128, 256]]) – Maximum number of registers per thread (None, 32, 64, 128, 256). Lower values allow higher occupancy but may hurt performance. Used with persistent kernels to ensure multi-occupancy can be achieved.

  • 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”

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

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

  • flatten_loops (list[bool] | None) –

Methods

__init__(*[, block_sizes, num_threads, ...])

Initialize a Config object.

from_dict(config_dict)

Create a Config from a plain dictionary.

from_json(json_str)

Create a Config object from a JSON string.

get(k[,d])

items()

keys()

load(path)

Load a config from a JSON file.

minimize(config_spec)

Return a new Config with values matching effective defaults removed.

save(path)

Save the config to a JSON file.

to_json()

Convert the config to a JSON string.

values()

Attributes

advanced_controls_file

atomic_indexing

block_sizes

cross_loop_pipeline

cute_affine_scan_schedule

cute_async_load_cache

cute_async_load_group_rows

cute_async_load_lookahead

cute_async_load_stages

cute_async_store_policy

cute_bf16x2_recurrence

cute_proven_bounds

epilogue_subtile

flatten_loops

indexing

l2_groupings

load_cache_modifiers

load_eviction_policies

loop_orders

maxnreg

num_sm_multiplier

num_stages

num_threads

num_warps

pallas_load_buffer_count

pid_type

range_flattens

range_multi_buffers

range_num_stages

range_unroll_factors

range_warp_specializes

reduction_loops

static_ranges

store_cache_modifiers

xcd_remap

config

__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_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_proven_bounds=None, cute_affine_scan_schedule=None, num_warps=None, num_stages=None, pid_type=None, cross_loop_pipeline=None, num_sm_multiplier=None, maxnreg=None, indexing=None, atomic_indexing=None, advanced_controls_file=None, epilogue_subtile=None, xcd_remap=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”.

  • 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_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_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_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.

  • 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" 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 (Optional[Literal[1, 2, 4, 8]]) – Multiplier for the number of SMs in persistent kernels (1, 2, 4, 8). Controls multi-occupancy by launching N * num_sms thread blocks instead of just num_sms.

  • maxnreg (Optional[Literal[32, 64, 128, 256]]) – Maximum number of registers per thread (None, 32, 64, 128, 256). Lower values allow higher occupancy but may hurt performance. Used with persistent kernels to ensure multi-occupancy can be achieved.

  • 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”

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

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

  • flatten_loops (list[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 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_bf16x2_recurrence: bool#
property cute_proven_bounds: bool#
property cute_affine_scan_schedule: Literal['ordinary', 'direct_m16n8_v1', 'direct_m16n16_v1']#
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#
On this page
PyTorch Libraries