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...,
dispatcher_num_sms: int = 20,
dispatcher_share_token_dispatcher: bool = True,
dispatcher_async_dispatch: bool = False,
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. 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.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, 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_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_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.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
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