Skip to content

Commit ebb93b2

Browse files
Connor-XYchtruong814claude
authored
Make Triton autotune configuration selection deterministic (#6076)
Signed-off-by: Yan Xu <yxu1@nvidia.com> Co-authored-by: Charlie Truong <chtruong@nvidia.com> Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
1 parent 9445e48 commit ebb93b2

39 files changed

Lines changed: 4425 additions & 185 deletions

‎MANIFEST.in‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,6 @@
11
include megatron/core/requirements.txt
22
include megatron/core/README.md
33
include megatron/core/package_info.py
4+
recursive-include megatron/core/tuning/tables *.json
45
global-exclude LICENSE
56
recursive-include requirements *

‎docs/developer/determinism/op-catalog.md‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,7 @@ Selected by `torch.are_deterministic_algorithms_enabled()` or
3838
| Vocab-parallel embedding | `megatron/core/tensor_parallel/layers.py` | direct indexing `weight[idx]` (deterministic backward) | `F.embedding` (non-deterministic atomic backward) |
3939
| Gated-delta-net kernel | `megatron/core/ssm/gated_delta_net.py` | torch `chunk_gated_delta_rule` | FLA fused kernel |
4040
| Gated-delta-net causal conv1d | `megatron/core/ssm/gated_delta_net/` | `F.conv1d` (plus transposes) | FLA `causal_conv1d` |
41-
| Mamba/SSM Triton ops | `megatron/core/ssm/ops/common/determinism.py` | one fixed autotune config plus a zero-initialized tiled workspace reduced with an ordered `sum` | timing-based autotune, uninitialized workspace |
41+
| Mamba/SSM Triton ops | `megatron/core/tuning` (config selection), `megatron/core/ssm/ops/common/determinism.py` (workspaces) | one fixed autotune config plus a zero-initialized tiled workspace reduced with an ordered `sum` | timing-based autotune, uninitialized workspace |
4242
| Mamba/GDP causal conv1d | `megatron/core/ssm/causal_conv1d.py` | causal_conv1d >= 1.6.0 — per-block workspace for the weight and bias gradients, reduced with an ordered `sum` | `atomicAdd` accumulation (order varies per launch) |
4343
| Transformer Engine attention | `megatron/core/extensions/transformer_engine.py` | requires `NVTE_ALLOW_NONDETERMINISTIC_ALGO=0`, under which TE selects only backends that support deterministic execution (including deterministic FlashAttention backward) | TE picks freely, including atomic-accumulation attention backward |
4444
| Inference DP scheduling and RL rollout order | `megatron/core/inference/engines/dynamic_engine.py`, `megatron/rl/rl_utils.py` | sort by stable key | completion order |

‎docs/user-guide/deterministic-training.md‎

Lines changed: 9 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -23,17 +23,17 @@ When enabled, Megatron applies the env vars and config overrides below via `mega
2323

2424
## Environment variables
2525

26-
Each variable may be set by the launcher or left unset. If set, the value must be one that has been validated as deterministic — anything else fails hard with an assertion. If unset, `apply_determinism_env` fills the canonical default (except `MAMBA_DETERMINISTIC` and `CAUSAL_CONV1D_DETERMINISTIC`, which their kernels auto-detect from `torch.are_deterministic_algorithms_enabled()`). Must be set before the first cuBLAS / Transformer Engine call — `apply_determinism_to_args` runs early in `validate_args` to guarantee this.
26+
Each variable may be set by the launcher or left unset. If set, the value must be one that has been validated as deterministic — anything else fails hard with an assertion. If unset, `apply_determinism_env` fills the canonical default (except `MAMBA_DETERMINISTIC` and `CAUSAL_CONV1D_DETERMINISTIC`, which the external `mamba_ssm` and `causal_conv1d` packages auto-detect from `torch.are_deterministic_algorithms_enabled()`). Must be set before the first cuBLAS / Transformer Engine call — `apply_determinism_to_args` runs early in `validate_args` to guarantee this.
2727

2828
| Variable | Accepted values (or unset) | Default filled if unset | Reason |
2929
|---|---|---|---|
3030
| `NCCL_ALGO` | subset of `{Ring, CollnetDirect, CollnetChain, ^NVLS}` | `Ring` | Conservative default — `Ring`'s reduction order is fixed by topology, so it is bit-exact across runs on every supported NCCL version |
3131
| `NVTE_ALLOW_NONDETERMINISTIC_ALGO` | `0` | `0` | Forces Transformer Engine to use deterministic algorithms |
3232
| `CUBLAS_WORKSPACE_CONFIG` | `:4096:8` or `:16:8` | `:4096:8` | Deterministic cuBLAS workspace (both sizes are reproducible per NVIDIA docs; `:4096:8` is faster, `:16:8` uses less memory) |
33-
| `TRITON_CACHE_AUTOTUNING` | `0` or `1` | *(none — opt-in)* | Persists each Triton autotune winner so every rank reuses one choice instead of re-timing it. Left unset, deterministic mode instead pins the cheapest config, which needs no cache — see [Triton autotuning](#triton-autotuning) |
33+
| `TRITON_CACHE_AUTOTUNING` | `0` or `1` | *(none — opt-in)* | Persists Triton autotune winners for autotuners **outside** the pinned scope, so ranks reuse one choice instead of re-timing it. Kernels inside the scope are pinned without timing either way — see [Triton autotuning](#triton-autotuning) |
3434
| `TRITON_CACHE_DIR` | any shared-filesystem path | *(none — required only with `TRITON_CACHE_AUTOTUNING=1`)* | No safe default exists: unset, Triton uses a node-local directory and each node autotunes on its own. Required rather than filled in |
35-
| `TRITON_PRINT_AUTOTUNING` | `1` | *(none — recommended, not set)* | Logs the config each rank selected. Changes no numerics, so it is recommended rather than forced; when `TRITON_CACHE_AUTOTUNING=1` and this is unset, a startup line reminds you. See [Verifying kernel-config agreement](#verifying-kernel-config-agreement) |
36-
| `MAMBA_DETERMINISTIC` | any string starting with `'1'` | *(none — SSM auto-detects)* | Mamba SSM auto-follows `torch.are_deterministic_algorithms_enabled()` when unset; only an explicit non-deterministic override is rejected |
35+
| `TRITON_PRINT_AUTOTUNING` | `1` | *(none — not set)* | Logs the config Triton times, which only happens for autotuners outside the pinned scope. For pinned kernels use `--triton-autotune-enumerate-autotuners` and `--triton-autotune-verify-every` — see [Verifying kernel-config agreement](#verifying-kernel-config-agreement) |
36+
| `MAMBA_DETERMINISTIC` | any string starting with `'1'` | *(none — SSM auto-detects)* | Controls the external `mamba_ssm` package, which auto-follows `torch.are_deterministic_algorithms_enabled()` when unset; only an explicit non-deterministic override is rejected. Megatron's in-tree SSM kernels (`megatron.core.ssm.ops`) do not read it and follow `--deterministic-mode` instead |
3737
| `CAUSAL_CONV1D_DETERMINISTIC` | any string starting with `'1'` | *(none — the kernel auto-detects)* | causal_conv1d ≥ 1.6.0 auto-follows `torch.are_deterministic_algorithms_enabled()` when unset, reducing the conv weight/bias gradients through a workspace instead of `atomicAdd`; the Mamba and GDP mixers reject a deterministic run without it |
3838

3939
If you override `NCCL_ALGO`, the value must be a subset of `{Ring, CollnetDirect, CollnetChain, ^NVLS}`. `Tree` is intentionally excluded: its intra-node chain reduction order is not user-controllable, and the inter-node tree topology can vary across runs without a pinned topology file, so it cannot be vouched for as bit-exact across stacks. `^NVLS` is accepted (banning NVLS is a legitimate user choice on hardware that exposes it); the user is responsible for ensuring whatever NCCL falls back to is deterministic on their environment.
@@ -69,20 +69,17 @@ torch.utils.deterministic.fill_uninitialized_memory = True
6969

7070
## Triton autotuning
7171

72-
Triton picks a kernel config by timing its candidates, so the winner depends on the machine at that instant and ranks can disagree. Deterministic mode offers two ways to remove that variance:
72+
Triton picks a kernel config by timing its candidates, so the winner depends on the machine at that instant and ranks can disagree. `--deterministic-mode` pins it: the `megatron.core.tuning` adapter selects one candidate per kernel and shape without timing anything, from a tuned table when one matches (tables ship for `sm100` and `sm103`, covering the `mamba_ssm` SSD kernels and the in-tree MLA RoPE and mHC kernels, or can be recorded with `--triton-autotune-mode record`), otherwise from the cheapest candidate by a pure function of the candidate list. Pinned kernels never consult Triton's autotune cache, so every rank computes the same answer by construction. The cheapest candidate is not necessarily the fastest one; a recorded table recovers the throughput.
7373

74-
| Strategy | How to select it | Determinism rests on |
75-
|---|---|---|
76-
| **Pinned config** (default) | leave `TRITON_CACHE_AUTOTUNING` unset | Nothing external. `autotune_configs` picks the cheapest config by a pure function of the candidate list, so every rank computes the same answer without timing anything. Slower, since the pinned config is not necessarily the fastest one. |
77-
| **Cached autotuning** | `TRITON_CACHE_AUTOTUNING=1` **and** `TRITON_CACHE_DIR=<shared path>` | Every rank reading one warm cache. Autotuning still runs and still picks fast configs, but a rank that misses the cache re-times the selection on its own and can pick differently. |
74+
The default scope is `mamba_ssm`, `transformer_engine` and `megatron.core` (`--triton-autotune-modules`). Kernels whose outputs do not depend on the config — by default only pure data movement, Transformer Engine's MoE permutation kernels, listed in `AutotunePolicy.config_invariant` (`--triton-autotune-config-invariant`) — keep Triton's timed choice, since pinning them would only cost throughput. See `megatron/core/tuning/README.md` for recording tables, precedence rules and the full option list.
7875

79-
Cached autotuning is opt-in because its determinism is conditional: the pinned default holds by construction, the cached path holds only while the shared cache does. Setting `TRITON_CACHE_AUTOTUNING=1` without `TRITON_CACHE_DIR` is rejected — unset, Triton falls back to a node-local directory, which is exactly the case the cache is meant to prevent.
76+
Autotuners outside the scope still time their candidates. `TRITON_CACHE_AUTOTUNING=1` with a shared `TRITON_CACHE_DIR` makes those ranks reuse one cached winner, but a rank that misses the cache re-times the selection on its own and can pick differently, so prefer adding their modules to `--triton-autotune-modules`. Setting `TRITON_CACHE_AUTOTUNING=1` without `TRITON_CACHE_DIR` is rejected — unset, Triton falls back to a node-local directory, which is exactly the case the cache is meant to prevent.
8077

8178
## Verifying kernel-config agreement
8279

83-
Applies to cached autotuning; the pinned default has nothing to compare. `TRITON_PRINT_AUTOTUNING=1` makes each rank log the config it selects per kernel; group those lines by kernel and key across the per-rank logs, and every group should hold exactly one distinct config.
80+
`--triton-autotune-verify-every N` compares, every N training steps, the configs that ranks chose for the same kernel and shape, and logs any disagreement (`--triton-autotune-verify-strict` raises instead). Each check exchanges only the choices made since the previous one. `--triton-autotune-enumerate-autotuners` logs, per rank, every multi-config autotuner the run reaches and whether it is pinned, timed as config-invariant, or outside the scope.
8481

85-
Note the limit: a rank only logs when it *tunes*, so a run where some ranks hit the cache and others miss cannot be compared this way — the hitting ranks log nothing.
82+
For autotuners outside the scope, `TRITON_PRINT_AUTOTUNING=1` makes each rank log the config Triton selects; a rank only logs when it tunes, so a run where some ranks hit the autotune cache and others miss cannot be compared this way.
8683

8784
## Verifying determinism
8885

‎megatron/core/ssm/ops/common/causal_conv1d_varlen.py‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -10,7 +10,7 @@
1010
import triton
1111
import triton.language as tl
1212

13-
from megatron.core.ssm.ops.common.determinism import autotune_configs
13+
from megatron.core.tuning import autotune_configs
1414

1515

1616
# Two block-dim regimes:

‎megatron/core/ssm/ops/common/determinism.py‎

Lines changed: 20 additions & 92 deletions
Original file line numberDiff line numberDiff line change
@@ -5,102 +5,21 @@
55
# This source code is licensed under the Apache license found in the
66
# LICENSE file in the root directory of this source tree.
77

8-
import os
9-
import warnings
10-
11-
import torch
12-
from packaging import version
13-
14-
try:
15-
import triton
16-
17-
TRITON_VERSION = version.parse(triton.__version__)
18-
except ImportError:
19-
TRITON_VERSION = version.parse("0.0.0")
20-
21-
TRITON_HAS_CACHE_RESULTS = TRITON_VERSION >= version.parse("3.4.0")
22-
_autotune_warning_issued = False
23-
24-
_deterministic_override = None
25-
26-
27-
def use_deterministic_mode():
28-
"""Use torch deterministic mode."""
29-
if _deterministic_override is not None:
30-
return _deterministic_override
31-
env = os.environ.get('MAMBA_DETERMINISTIC')
32-
if env:
33-
return env[0] == '1'
34-
return torch.are_deterministic_algorithms_enabled()
35-
36-
37-
def set_deterministic_mode(value):
38-
"""Set torch deterministic mode."""
39-
global _deterministic_override
40-
_deterministic_override = value
8+
"""Deterministic-mode helpers for the SSM Triton kernels.
419
10+
Kernel-config selection used to live here, but it is not specific to SSM: the
11+
same wall-clock autotuning affects Transformer Engine's MoE permutation kernels,
12+
and a model without Mamba needs it just as much. It moved to
13+
:mod:`megatron.core.tuning`, which is also where the policy is installed from.
4214
43-
def _estimate_config_cost(cfg):
44-
"""Estimate shared memory cost of a config. Lower is cheaper.
15+
What remains here is the tile workspaces that let the SSD kernels accumulate in
16+
a fixed order, plus re-exports of the names this module defined before the move
17+
so out-of-tree callers keep working.
18+
"""
4519

46-
Returns a tuple (block_cost, num_warps) so that ties in block cost
47-
are broken deterministically by warp count (fewer warps = cheaper).
48-
"""
49-
block_product = 1
50-
for key, val in cfg.kwargs.items():
51-
if key.startswith('BLOCK') and isinstance(val, int):
52-
block_product *= val
53-
stages = getattr(cfg, 'num_stages', 1) or 1
54-
warps = getattr(cfg, 'num_warps', 1) or 1
55-
return (block_product * stages, warps)
56-
57-
58-
def _filter_configs_by_block_sizes(configs):
59-
"""Filter configs by TRITON_AUTOTUNE_BLOCK_* env vars.
60-
61-
Scans environment for any variable matching TRITON_AUTOTUNE_BLOCK_*
62-
(e.g. TRITON_AUTOTUNE_BLOCK_SIZE_M, TRITON_AUTOTUNE_BLOCK_SIZE_H,
63-
TRITON_AUTOTUNE_BLOCK_T, TRITON_AUTOTUNE_BLOCK_C, TRITON_AUTOTUNE_BLOCK_SIZE)
64-
and maps them to the corresponding kernel kwarg (BLOCK_SIZE_M, BLOCK_SIZE_H,
65-
BLOCK_T, BLOCK_C, BLOCK_SIZE).
66-
"""
67-
prefix = "TRITON_AUTOTUNE_"
68-
env_filters = {}
69-
for env_key, env_val in os.environ.items():
70-
if env_key.startswith(prefix + "BLOCK") and env_val:
71-
kwarg_name = env_key[len(prefix) :]
72-
env_filters[kwarg_name] = int(env_val)
73-
if not env_filters:
74-
return None
75-
matching = configs
76-
for key, target in sorted(env_filters.items()):
77-
matching = [c for c in matching if c.kwargs.get(key) == target]
78-
return matching[:1] if matching else None
79-
80-
81-
def autotune_configs(configs):
82-
"""Select autotune configs for deterministic mode.
20+
import torch
8321

84-
Uses cached autotuning (TRITON_CACHE_AUTOTUNING=1) if Triton >= 3.4.0,
85-
otherwise auto-selects the cheapest config by block size * stages.
86-
"""
87-
if not configs or not use_deterministic_mode():
88-
return configs
89-
if TRITON_HAS_CACHE_RESULTS and os.environ.get("TRITON_CACHE_AUTOTUNING") == "1":
90-
return configs
91-
global _autotune_warning_issued
92-
if not _autotune_warning_issued:
93-
_autotune_warning_issued = True
94-
msg = (
95-
"Deterministic mode: set TRITON_CACHE_AUTOTUNING=1 for cached autotuning."
96-
if TRITON_HAS_CACHE_RESULTS
97-
else "Deterministic mode: upgrade to Triton >= 3.4.0 for cached autotuning."
98-
)
99-
warnings.warn(msg)
100-
filtered = _filter_configs_by_block_sizes(configs)
101-
if filtered:
102-
return filtered
103-
return [min(configs, key=_estimate_config_cost)]
22+
from megatron.core.tuning import autotune_configs, set_deterministic_mode, use_deterministic_mode
10423

10524

10625
def alloc_tile_workspace(base_shape, tile_dim, dtype, device, deterministic, *, zero_init=True):
@@ -121,3 +40,12 @@ def finalize_tile_workspace(tensor, deterministic):
12140
if deterministic:
12241
tensor = tensor.sum(dim=-1)
12342
return tensor
43+
44+
45+
__all__ = [
46+
"alloc_tile_workspace",
47+
"autotune_configs",
48+
"finalize_tile_workspace",
49+
"set_deterministic_mode",
50+
"use_deterministic_mode",
51+
]

‎megatron/core/ssm/ops/gdp/chunk_h.py‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@
2020

2121
import torch
2222

23-
from megatron.core.ssm.ops.common.determinism import autotune_configs
23+
from megatron.core.tuning import autotune_configs
2424

2525
from .common import HAVE_TRITON, exp2, prepare_chunk_indices, prepare_chunk_offsets, tl, triton
2626

‎megatron/core/ssm/ops/gdp/chunk_o.py‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@
1717

1818
import torch
1919

20-
from megatron.core.ssm.ops.common.determinism import autotune_configs
20+
from megatron.core.tuning import autotune_configs
2121

2222
from .common import (
2323
HAVE_TRITON,

‎megatron/core/ssm/ops/gdp/common.py‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,7 @@
2727

2828
import torch
2929

30-
from megatron.core.ssm.ops.common.determinism import autotune_configs
30+
from megatron.core.tuning import autotune_configs
3131

3232
try:
3333
import triton

‎megatron/core/ssm/ops/gdp/cumsum.py‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,7 @@
1414

1515
import torch
1616

17-
from megatron.core.ssm.ops.common.determinism import autotune_configs
17+
from megatron.core.tuning import autotune_configs
1818

1919
from .common import HAVE_TRITON, prepare_chunk_indices, tl, triton
2020

‎megatron/core/ssm/ops/gdp/scaled_dot_kkt.py‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,7 @@
1414

1515
import torch
1616

17-
from megatron.core.ssm.ops.common.determinism import autotune_configs
17+
from megatron.core.tuning import autotune_configs
1818

1919
from .common import HAVE_TRITON, exp2, prepare_chunk_indices, tl, triton
2020

0 commit comments

Comments
 (0)