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_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]#
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) –

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_exact_div_form

cute_packet_prefetch

cute_proven_bounds

cute_reduction_pipeline_depth

cute_store_policy

cute_wide_packet_hints

epilogue_subtile

flatten_loops

host_tensor_descriptors

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_fusion

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_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#
On this page
PyTorch Libraries