Skip to navigation

nemo_automodel.components.models.common

View as Markdown

Common/shared components for model implementations.

Subpackages

Submodules

Package Contents

Classes

NameDescription
BackendConfigBackend configuration for model components.
CudaGraphConfigConfiguration for scoped partial CUDA graphs.
HFCheckpointingMixinMixin providing HF-compatible API using NeMo’s checkpointing infrastructure.
MoKBackendConfigConfiguration for the optional Mixture-of-Kittens MoE backend.

Functions

NameDescription
compute_lm_head_logitsProject hidden states through lm_head and wrap the result.
get_rope_configExtract rope configuration from config.rope_parameters.
initialize_linear_moduleInitialize Linear module with the specified backend.
initialize_rms_norm_moduleInitialize RMSNorm module with the specified backend.

Data

AttentionBackend

HAVE_DEEP_EP

HAVE_TE

API

class nemo_automodel.components.models.common.utils.BackendConfig(
sparse_attn: typing.Literal['generic', 'msa'] = 'generic',
linear: typing.Literal['torch', 'te', 'quack'] = 'te' if HAVE_TE and torch.c...,
rms_norm: typing.Literal['torch', 'torch_fp32', 'te', 'quack'] = 'torch_fp32',
rope: typing.Literal['torch', 'quack', 'triton'] = 'torch',
rope_fusion: bool = HAVE_TE and torch.cuda.is_a...,
experts: typing.Literal['torch', 'te', 'gmm', 'torch_mm', 'torch_mm_mxfp8'] = 'torch_mm' if torch.cuda.is...,
dispatcher: typing.Literal['torch', 'deepep', 'hybridep', 'uccl_ep', 'mok'] = 'deepep' if HAVE_DEEP_EP an...,
te_grouped_tensor: bool = False,
dispatcher_num_sms: int = 20,
dispatcher_share_token_dispatcher: bool = True,
dispatcher_async_dispatch: bool = False,
dispatcher_capacity_factor: float | None = None,
dispatcher_equal_token_counts: bool = False,
dispatcher_recompute_cache: str = 'none',
dispatcher_hybridep_permute_fusion: bool = False,
dispatcher_hybridep_compact_routing: bool = False,
dispatcher_hybridep_num_sms_preprocessing: int | None = None,
dispatcher_hybridep_num_blocks_permute: int | None = None,
dispatcher_hybridep_num_blocks_unpermute: int | None = None,
enable_deepep: bool | None = None,
fake_balanced_gate: bool = False,
fake_gate_noise: float = 0.0,
enable_hf_state_dict_adapter: bool = True,
enable_fsdp_optimizations: bool = False,
gate_precision: str | torch.dtype | None = None,
compile_attn: bool = False,
compile_situ: bool = False,
compile_norm: bool = False,
compile_hc: bool = False,
shared_expert_overlap: bool = False,
compile_router_weight: bool = False,
benchmark_static_routing: bool = False,
)
Dataclass

Backend configuration for model components.

attn
AttentionBackend

Attention backend (“torch”, “te”, “sdpa”, “flex”, “eager”, “tilelang”, “cudnn”, “fa4”, or “magi”). “fa4” selects native FlashAttention-4 (CuTe) and requires an INSTALL_FA4=true image. HF-dispatched models use “flash_attention_4”. A rank-2 binary padding mask is converted to varlen metadata, while dense/block masks and indexed document masks without explicit packed metadata are rejected. Inputs must use BSHD layout; pre-packed THD is unsupported. Packed Qwen3Next and DeepSeek V3.2 sparse attention do not support FA4. For DeepSeek V4, “tilelang” enables the TileLang sparse attention, indexer, and Sinkhorn kernels together; “cudnn” keeps the TileLang indexer/Sinkhorn and selects FlashMLA forward with cuDNN backward. For GLM DSA, “tilelang” and “cudnn” select their respective packed sparse-attention kernels. For Qwen3.8-Flash-Next, “flex” selects FlexAttention sparse GQA on CUDA BF16; CPU execution retains the PyTorch numerical oracle.

benchmark_static_routing
bool = False

Benchmark-only. Requires fake_balanced_gate=True with fake_gate_noise=0.0, where routing metadata (tokens per expert, permuted token counts) is identical for every microbatch. Skips the per-microbatch device-to-host reads of that metadata (.tolist() / count_nonzero / dispatcher size checks) by caching the first microbatch’s values, removing recurring host-sync stalls on the hot path. Never enable with a learned gate: cached metadata would go stale and silently corrupt expert dispatch.

compile_attn
bool = False

torch.compile(fullgraph) the attention module’s forward — both the DeepSeek-V3 MLA and standard GQA attention (e.g. Qwen3-MoE) honor it. Requires attn=“sdpa”, linear=“torch”, rms_norm=“torch”, rope_fusion=False.

compile_hc
bool = False

torch.compile the mHC (hyper-connection) cores of models that opt in (currently DeepSeek V4 / V4.1): coefficient projection, stream collapse and stream expand, one fused kernel each, static shapes (a new sequence length compiles them again; past dynamo’s recompile limit a core runs eager). Same lazy once-per-process pattern as compile_situ; numerics are allclose to eager but not bitwise-identical.

compile_norm
bool = False

torch.compile the fp32 RMSNorm chain of models that opt in (currently Kimi K3 and DeepSeek V4.1), fusing cast/pow/mean/rsqrt/mul into one kernel. Same lazy once-per-process pattern as compile_situ; numerics are allclose to eager but not bitwise-identical.

compile_router_weight
bool = False

torch.compile the fp32 router-weight multiply applied to expert outputs (GroupedExperts / GroupedExpertsDeepEP / GroupedExpertsTE, all MoE models), so each pass is one fused kernel instead of the eager per-chunk cast/multiply/ cast/slice-assign loop. Same lazy once-per-process pattern as compile_situ; numerics are allclose to eager but not bitwise-identical.

compile_situ
bool = False

torch.compile the SiTU expert activation of models that opt in (currently Kimi K3): the weighted-SiTU forward and backward run as one fused whole-tensor kernel each (low-precision in/out, fp32 math inside, no chunk loop), the dense / shared-expert SituAndMul uses a compiled core, and the attn-res mixing chain is compiled as well. Compiled numerics are allclose to eager but not bitwise-identical.

cuda_graph
CudaGraphConfig = field(default_factory=CudaGraphConfig)

Scoped partial CUDA-graph configuration.

dispatcher
Literal['torch', 'deepep', 'hybridep', 'uccl_ep', 'mok']

MoE token dispatcher. “torch” uses DTensor all-gather/reduce-scatter, “deepep” uses DeepEP, “hybridep” uses HybridEP, “uccl_ep” uses UCCL-EP, and “mok” uses Mixture-of-Kittens’ fused communication and SwiGLU megakernel.

dispatcher_async_dispatch
bool = False

Whether DeepEP/UCCL-EP dispatch and combine should return asynchronously and allocate their outputs on the communication stream.

dispatcher_capacity_factor
float | None = None

HybridEP only, dynamic routing. Run HybridEP dispatch in its non-blocking mode with output buffers sized to the first microbatch’s permuted row count times this factor (EP-group max, 4-token aligned) instead of letting every dispatch drain the compute stream to read the exact count. Overflow trips a device-side assert. None (default) keeps the blocking reference path. The capacity is calibrated once per communication manager from its first real batch and kept for the run; with dispatcher_share_token_dispatcher=True (default) that is one calibration, taken on the first local MoE layer’s first microbatch, for every layer, so choose a factor that covers the largest routing skew expected across layers and batches. Not supported for experts with expert_bias (GroupedExpertsDeepEP rejects the combination).

dispatcher_equal_token_counts
bool = False

HybridEP only. Declare that every EP rank dispatches the same row count, so the per-dispatch EP-group max all-reduce (and its host sync) that derives the pad size is skipped and the count is only aligned. Default False; keep it False for variable-length or in-batch-packed inputs.

dispatcher_hybridep_compact_routing
bool = False

Send compact top-k expert indices to HybridEP instead of the dense Boolean routing map. This can reduce metadata traffic, but remains opt-in because its latency depends on topology and model shape.

dispatcher_hybridep_num_blocks_permute
int | None = None

Optional number of dispatch permutation blocks. None preserves HybridEP’s default.

dispatcher_hybridep_num_blocks_unpermute
int | None = None

Optional number of combine unpermutation blocks. None preserves HybridEP’s default.

dispatcher_hybridep_num_sms_preprocessing
int | None = None

Optional number of SMs used by HybridEP routing-metadata preprocessing.

dispatcher_hybridep_permute_fusion
bool = False

Fuse HybridEP token permutation into dispatch and unpermutation into combine. Has no effect for other dispatcher backends.

dispatcher_num_sms
int = 20
dispatcher_recompute_cache
str = 'none'

HybridEP only; effective under full-block activation checkpointing (not installed for checkpoint_moe_only regions, which end at the combine the recompute never re-runs). "none" (default) recomputes the whole MoE layer in the backward. "combine" keeps a clone of every MoE layer’s combine output ([tokens, hidden] in the activation dtype) from the checkpoint forward and returns it in the recompute instead of re-running the unpermute, the EP-group barrier and the combine all-to-all. "moe_tail" additionally builds the down projection’s autograd node in the recompute without running its GEMM (its output only fed the skipped combine; torch_mm experts without expert bias and with the router weight applied inside the activation) (a post-expert callback that reads the output keeps the GEMM). Costs one [tokens, hidden] activation per MoE layer per in-flight microbatch; gradients are unchanged.

dispatcher_share_token_dispatcher
bool = True

Whether flex token dispatchers share a communication manager instance across MoE layers.

enable_deepep
bool | None = None

Removed and ignored. Logs a warning if set; configure “dispatcher” and “experts” explicitly instead.

enable_fsdp_optimizations
bool = False

Whether to enable FSDP2 optimizations.

enable_hf_state_dict_adapter
bool = True

Whether to enable HuggingFace state dict adapter.

experts
Literal['torch', 'te', 'gmm', 'torch_mm', 'torch_mm_mxfp8']

MoE expert GEMM backend. “torch” uses per-expert loop, “te” uses TE GroupedLinear, “torch_mm” uses torch._grouped_mm, “gmm” is a deprecated compatibility alias for “torch_mm”, and “torch_mm_mxfp8” uses torch._grouped_mm dispatch but routes the expert grouped GEMMs through torchao’s MXFP8 scaled grouped GEMM (training-only; GB200/sm_100+ with torchao installed, else falls back to torch._grouped_mm at runtime).

fake_balanced_gate
bool = False

If True, replace the learned Gate with FakeBalancedGate that assigns tokens to experts without learned routing weights.

fake_gate_noise
float = 0.0

Noise level [0, 1] for FakeBalancedGate. When > 0, uses biased topk selection seeded from the input content so routing varies dynamically across training steps (like real Gate) while remaining deterministic for activation checkpointing recompute (same input = same routing). Only used when fake_balanced_gate=True.

gate_precision
str | dtype | None = None

Optional dtype override for the gate computation. Accepts torch.dtype or string (e.g., “torch.float32”, “float32”).

linear
Literal['torch', 'te', 'quack']

Linear layer backend (“torch”, “te”, or “quack”).

mok
MoKBackendConfig = field(default_factory=MoKBackendConfig)

Mixture-of-Kittens tuning settings used when dispatcher="mok".

rms_norm
Literal['torch', 'torch_fp32', 'te', 'quack'] = 'torch_fp32'

RMSNorm backend (“torch”, “torch_fp32”, “te”, or “quack”).

rope
Literal['torch', 'quack', 'triton'] = 'torch'

Rotary embedding backend (“torch”, “quack” or “triton”). QuACK is currently integrated for Llama-family rotary embeddings; “triton” selects the fused Triton rotation of DeepSeek V4 / V4.1 (other models keep the torch path).

rope_fusion
bool = HAVE_TE and torch.cuda.is_available()

Whether to use fused RoPE (requires TE).

shared_expert_overlap
bool = False

run the shared experts of MoE models (the shared MoE class, e.g. DeepSeek V4 / V4.1, and Kimi K3) on a side CUDA stream so their GEMMs overlap the expert-parallel dispatch / combine communication of the routed path; numerics unchanged. Default False.

sparse_attn
Literal['generic', 'msa'] = 'generic'

Sparse-attention backend. “generic” preserves each model’s existing sparse mask plus attn path; “msa” selects the optional SM100 MSA kernels a model provides for its sparse layers only.

te_fp8
TEFp8Config | None = None
te_grouped_tensor
bool = False

Prefer TE 2.19 native grouped-tensor GEMMs for experts="te". Requires dispatcher="deepep", "hybridep", or "uccl_ep". Device and precision support follows TE’s GroupedLinear contract. Parameters and checkpoints retain their existing layouts.

nemo_automodel.components.models.common.utils.BackendConfig.__post_init__() -> None
class nemo_automodel.components.models.common.utils.CudaGraphConfig(
modules: list[typing.Literal['attn', 'te_dpa', 'moe_router', 'moe_preprocess']] = list()
)
Dataclass

Configuration for scoped partial CUDA graphs.

modules
list[Literal['attn', 'te_dpa', 'moe_router', 'moe_preprocess']] = field(default_factory=list)

Per-layer modules to capture with Transformer Engine, following Megatron Core’s module names. AutoModel supports whole attn, moe_router, and moe_preprocess scopes, plus the AutoModel-specific narrow te_dpa scope. An empty list disables CUDA graphs.

nemo_automodel.components.models.common.utils.CudaGraphConfig.__post_init__() -> None

Validate the declarative CUDA-graph configuration.

class nemo_automodel.components.models.common.hf_checkpointing_mixin.HFCheckpointingMixin()

Mixin providing HF-compatible API using NeMo’s checkpointing infrastructure.

Provides save_pretrained() and from_pretrained() methods that use Checkpointer for unified distributed/async support with HF format conversion.

Key design: We do NOT override state_dict() or load_state_dict() because PyTorch’s DCP expects these to behave like standard nn.Module methods.

For PreTrainedModel subclasses:

  • super().from_pretrained() handles: downloads, quantization config, meta device init
  • Checkpointer.load_base_model() handles: actual weight loading with format conversion

For nn.Module subclasses (no parent from_pretrained):

  • Falls back to manual config loading + Checkpointer
nemo_automodel.components.models.common.hf_checkpointing_mixin.HFCheckpointingMixin.save_pretrained(
save_directory: str,
tokenizer: typing.Optional[transformers.tokenization_utils.PreTrainedTokenizerBase] = None,
kwargs = {}
) -> None

Save model in HF-compatible format using Checkpointer infrastructure.

Supports distributed saving, sharding, and async checkpointing.

Parameters:

save_directory
str

Output path

checkpointer
Optional[Checkpointer]Defaults to None

Checkpointer instance. Uses self._checkpointer if not provided.

tokenizer
Optional[PreTrainedTokenizerBase]Defaults to None

Optional tokenizer to save alongside model

**kwargs
Defaults to {}

Additional arguments, including peft_config and is_final_checkpoint. Direct callers that do not have recipe step-scheduler context default is_final_checkpoint to False.

class nemo_automodel.components.models.common.utils.MoKBackendConfig(
precision: typing.Literal['bf16', 'mxfp8'] = 'bf16',
fwd_num_comm_sms: int = 40,
bwd_num_comm_sms: int = 28,
minibatch_size: int = 4096,
macrobatch_size: int = 131072,
schedule_capacity_multiplier: float = 0.5,
all_gather_top_experts_chunk_bytes: int = 2048
)
Dataclass

Configuration for the optional Mixture-of-Kittens MoE backend.

all_gather_top_experts_chunk_bytes
int = 2048

Routing all-gather chunk size in bytes.

bwd_num_comm_sms
int = 28

Communication SMs reserved during the backward megakernel.

fwd_num_comm_sms
int = 40

Communication SMs reserved during the forward megakernel.

macrobatch_size
int = 131072

Routed-token ring-buffer capacity and a multiple of minibatch_size.

minibatch_size
int = 4096

Routed-token overlap granularity, divisible by 256.

precision
Literal['bf16', 'mxfp8'] = 'bf16'

Routed-expert compute precision. "bf16" passes the materialized BF16 weights directly to MoK; "mxfp8" quantizes them with MoK’s E4M3 data and E8M0 block-scale format.

schedule_capacity_multiplier
float = 0.5

Worst-case fraction of routes sent to one EP rank.

nemo_automodel.components.models.common.utils.MoKBackendConfig.__post_init__() -> None

Validate settings that do not depend on a CUDA device or EP group.

nemo_automodel.components.models.common.utils.MoKBackendConfig.build() -> nemo_automodel.components.models.common.utils._MoKFunctionalConfig

Build the optional MoK runtime configuration.

Returns: _MoKFunctionalConfig

A mok.functional.MoKConfig with these declarative settings.

Raises:

  • UnavailableError: If Mixture-of-Kittens is not installed.
nemo_automodel.components.models.common.utils.compute_lm_head_logits(
lm_head: torch.nn.Module | None,
hidden_states: torch.Tensor,
logits_to_keep: int | torch.Tensor = 0,
is_thd: bool = False,
fp32_lm_head: bool = False,
output_hidden_states: bool = False
) -> transformers.modeling_outputs.CausalLMOutputWithPast

Project hidden states through lm_head and wrap the result.

Centralizes the lm_head projection and output packaging shared by every custom *ForCausalLM / *ForConditionalGeneration forward(). The returned CausalLMOutputWithPast carries the projected logits and, when requested, the final hidden_states; callers that also need loss, past_key_values, etc. read .logits and build their own output.

  • lm_head is None (e.g. a non-final pipeline-parallel stage that does not own the head): hidden_states is passed through as logits so the next stage receives it.
  • logits_to_keep == 0 (training default): every position is projected. The full range is deliberately not sliced, because slice(0, None) on a DTensor is unsupported (it raises on the aten.alias op under tensor parallel with sequence parallelism).
  • logits_to_keep as a positive int or a tensor of indices: only the requested positions are projected. Both 2D [T, H] (THD/packed) and 3D [B, S, H] (BSHD) hidden states are handled.
  • is_thd: THD/packed inputs yield 2D [T, V] logits; the leading batch dim is restored (unsqueeze(0) -> [1, T, V]) so downstream code sees a uniform [B, S, V] layout. The same restoration is applied to the hidden_states field. Only applied while the tensor is still 2D, so an inputs_embeds path that already produced [1, T, *] is left untouched.
  • fp32_lm_head: run the projection in fp32 and cast the logits back to the input dtype. Used by models whose lm_head.weight has been promoted to fp32 (e.g. via the MoE lm_head_precision setting). The matmul goes through lm_head (nn.Linear, DTensor-aware under FSDP2) rather than F.linear so DTensor redistribution is preserved.
  • Otherwise, cast the sliced input to the head’s effective compute dtype. For a standalone FSDP2 head this comes from its mixed-precision policy, since its sharded weight may still expose the master-copy dtype before the forward pre-hook runs. Other heads use their materialized weight dtype.
  • output_hidden_states: when set, the (full-sequence, THD-restored) hidden_states are attached to the output so the fused cross-entropy path can recompute logits over every position; otherwise the field is None.

Parameters:

lm_head
nn.Module | None

The language-model head module, or None on a pipeline stage that does not own it.

hidden_states
torch.Tensor

Final hidden states, shaped [T, H] or [B, S, H].

logits_to_keep
int | torch.TensorDefaults to 0

0 to project every position; a positive int to keep the last N positions; or a tensor of position indices.

is_thd
boolDefaults to False

Whether the inputs were THD/packed; if so, a 2D logits (and hidden-states) result is unsqueezed back to a leading batch dim of 1.

fp32_lm_head
boolDefaults to False

Project in fp32 and cast the result back to the input dtype. Ignored when lm_head is None.

output_hidden_states
boolDefaults to False

Attach the final hidden states to the output.

Returns: CausalLMOutputWithPast

A CausalLMOutputWithPast whose logits are the projected logits

nemo_automodel.components.models.common.utils.get_rope_config(
config
) -> tuple[float, dict, float]

Extract rope configuration from config.rope_parameters.

Parameters:

config

A HuggingFace model config object.

Returns: tuple[float, dict, float]

Tuple of (rope_theta, rope_parameters, partial_rotary_factor).

nemo_automodel.components.models.common.utils.initialize_linear_module(
linear_impl: str,
in_features: int,
out_features: int,
bias: bool = False,
device: torch.device | str | None = None,
dtype: torch.dtype = torch.bfloat16
) -> torch.nn.Module

Initialize Linear module with the specified backend.

For TE backend, creates TE module directly on specified device. Call reset_parameters() to materialize weights if created on meta device.

Parameters:

linear_impl
str

Backend implementation (“te”, “torch”, or “quack”)

in_features
int

Input features

out_features
int

Output features

bias
boolDefaults to False

Whether to use bias

device
torch.device | str | NoneDefaults to None

Device to create module on (None uses PyTorch default, typically CPU)

dtype
torch.dtypeDefaults to torch.bfloat16

Parameter dtype

Returns: nn.Module

Linear module

nemo_automodel.components.models.common.utils.initialize_rms_norm_module(
rms_norm_impl: str,
dim: int,
eps: float = 1e-05,
device: torch.device | str | None = None,
dtype: torch.dtype = torch.bfloat16
) -> torch.nn.Module

Initialize RMSNorm module with the specified backend.

For TE backend, creates TE module directly on specified device. Call reset_parameters() to materialize weights if created on meta device.

Parameters:

rms_norm_impl
str

Backend implementation (“te”, “torch”, “torch_fp32”, or “quack”)

  • “te”: Transformer Engine fused RMSNorm kernel
  • “torch”: PyTorch native nn.RMSNorm (computes in input dtype)
  • “torch_fp32”: torch.compiled fp32 RMSNorm for training stability
  • “quack”: QuACK CuTe DSL RMSNorm kernel
dim
int

Normalized dimension

eps
floatDefaults to 1e-05

Epsilon for numerical stability

device
torch.device | str | NoneDefaults to None

Device to create module on (None uses PyTorch default, typically CPU)

dtype
torch.dtypeDefaults to torch.bfloat16

Parameter dtype

Returns: nn.Module

RMSNorm module

nemo_automodel.components.models.common.utils.AttentionBackend = Literal['torch', 'te', 'sdpa', 'flex', 'eager', 'tilelang', 'cudnn', 'fa4', 'mag...
nemo_automodel.components.models.common.utils.HAVE_DEEP_EP = importlib.util.find_spec('deep_ep') is not None
nemo_automodel.components.models.common.utils.HAVE_TE = importlib.util.find_spec('transformer_engine') is not None