Skip to content

Commit c5bc7e5

Browse files
yisongbetterakoumpajgerhHuiyingLi
authored
perf(deepseek-v41): opt-in torch.compile cores for mHC and fp32 RMSNorm (NVIDIA-NeMo#4079)
* perf(deepseek-v41): opt-in torch.compile cores for mHC and fp32 RMSNorm The single-pass mHC math and the fp32 RMSNorm run as 4-5 full fp32 passes each in eager mode, and HyperConnection.expand materializes a [batch, sequence, streams, streams, hidden] fp32 temporary before its stream sum (1.3 GB per call at 4096 tokens). On the 16-node GB200 benchmark these elementwise kernels are 44 % of the step. The arithmetic now lives in module-level cores that torch.compile can fuse: - DeepseekV41TextConfig.hc_impl = "compile" (default "torch") wraps the projection, collapse and expand cores once per process; the Sinkhorn normalization keeps its TileKernels/torch backend (compiling the 20-iteration torch loop is slower than the fused kernel); - BackendConfig.compile_norm (the existing generic knob, Kimi K3 pattern) wraps the fp32 RMSNorm core. Modules call the dispatchers, so instances built before the hook still pick up the compiled core. Static shapes only (dynamic=False): with dynamic shapes inductor keeps the broadcast temporary and expand runs at eager speed. Compiled numerics are allclose to eager (bf16 streams within one ulp, fp32 coefficient gradients within 5e-4), not bitwise-identical; both knobs default off. 1-GPU micro-benchmark (GB200, 4096 / 8192 tokens, fwd+bwd per call): expand 5.98 -> 0.75 ms / 11.86 -> 1.43, collapse 1.24 -> 0.49 / 2.46 -> 0.42, projection 1.50 -> 1.06 / 2.73 -> 1.07, RMSNorm 0.82 -> 0.44 / 1.57 -> 0.44; transient peak of expand 2.97 -> 0.51 GiB. Signed-off-by: Yisong Li <yisongbetter@gmail.com> * Apply batched suggestions from code review Co-authored-by: jgerh <163925524+jgerh@users.noreply.github.com> Signed-off-by: Alexandros Koumparoulis <153118171+akoumpa@users.noreply.github.com> * test(deepseek-v41): runtime budget for the compiled-cores parity test The test compiles the three mHC cores and the fp32 RMSNorm core for one dtype and takes about 10 s on the CI runners, over the 5 s default budget that main now applies to changed tests. Declare the exact budget with its reason as the conftest asks. Signed-off-by: Yisong Li <yisongbetter@gmail.com> * refactor(deepseek-v41): compile_hc moves to BackendConfig Review follow-up: the knob moves from DeepseekV41TextConfig.hc_impl (a model config field, saved into config.json) to BackendConfig.compile_hc, next to compile_norm and compile_situ, which select kernels the same way. Same hooks, same compiled cores; recipes set model.backend.compile_hc: true instead of model.config.text_config.hc_impl: compile. Tests and the model-coverage page follow. Signed-off-by: Yisong Li <yisongbetter@gmail.com> * fix(deepseek-v41): compiled mHC cores run eager past the recompile limit With fullgraph=True a run with more than recompile_limit distinct sequence lengths raised FailOnRecompileLimitHit once a core passed the limit. The cores have no graph breaks (asserted in the parity test), so fullgraph adds nothing at runtime; without it dynamo skips the frame and the core runs eager. New test lowers the limit to one and checks the second shape runs eager with eager numerics. Signed-off-by: Yisong Li <yisongbetter@gmail.com> --------- Signed-off-by: Yisong Li <yisongbetter@gmail.com> Signed-off-by: Alexandros Koumparoulis <153118171+akoumpa@users.noreply.github.com> Co-authored-by: Alexandros Koumparoulis <153118171+akoumpa@users.noreply.github.com> Co-authored-by: jgerh <163925524+jgerh@users.noreply.github.com> Co-authored-by: Huiying <willwin.lee@gmail.com>
1 parent 3770b1e commit c5bc7e5

6 files changed

Lines changed: 304 additions & 35 deletions

File tree

‎docs/model-coverage/llm/deepseek-ai/deepseek-v41-flash.mdx‎

Lines changed: 21 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -4,7 +4,7 @@ description: "Fine-tune DeepSeek-V4.1-Flash with NeMo AutoModel using checked-in
44
slug: model-coverage/large-language-models/deepseek-ai/DeepSeek-V4.1-Flash
55
---
66

7-
[DeepSeek-V4.1 Flash](https://huggingface.co/deepseek-ai/DeepSeek-V4.1-Flash)
7+
[DeepSeek-V4.1-Flash](https://huggingface.co/deepseek-ai/DeepSeek-V4.1-Flash)
88
uses a 40-layer causal encoder/decoder Mixture-of-Experts (MoE) backbone with
99
384 routed experts, one shared expert, and six selected experts per token. Its CSA2 attention shares
1010
compressed KV and index state between Full, Reuse, and Reindex layers.
@@ -41,7 +41,7 @@ The configuration requests 100 updates with validation every 50 updates,
4141
activation checkpointing, and reshard-after-forward. It uses no gradient
4242
accumulation or additional expert FSDP shard axis, and requests a one-hour
4343
allocation. W&B logging is disabled by default; set `wandb.enable: true` to
44-
enable online logging. TE FusedAdam retains FP32 master-weight precision
44+
enable online logging. Transformer Engine (TE) FusedAdam retains FP32 master-weight precision
4545
through int16 remainders for BF16 parameters.
4646

4747
<Note>
@@ -57,6 +57,18 @@ The model and recipe use `attn: tilelang`, `linear: torch`,
5757
Set `model.backend.rms_norm: te` to use Transformer Engine RMSNorm for the
5858
text backbone, attention projections, compressor, and indexer. This preserves
5959
parameter names and dtypes; the default retains eager FP32 normalization.
60+
Set `model.backend.compile_hc: true` to run the single-pass mHC
61+
coefficient projection, stream collapse and stream expand through
62+
`torch.compile` (one fused kernel each instead of four to five FP32 passes;
63+
the eager `expand` materializes a `[batch, sequence, streams, streams, hidden]`
64+
FP32 temporary), and `model.backend.compile_norm: true` to compile the FP32
65+
RMSNorm chain the same way. Both are opt-in: compiled numerics are allclose to
66+
eager, not bitwise-identical (the 12-layer mini benchmark reproduced the eager
67+
loss to four decimals). The cores are specialized per shape: a new sequence
68+
length (or grad mode) compiles them again, so keep them off for runs with many
69+
distinct packed lengths.
70+
The Sinkhorn normalization keeps its TileKernels or torch backend. Both knobs are
71+
backend fields: they select kernels and are not written to `config.json`.
6072
TileLang and TileKernels are the optional dependencies also used by DeepSeek V4.
6173
HybridEP provides both token dispatch and expert-output combine. Router
6274
correction biases remain fixed during fine-tuning through
@@ -105,7 +117,8 @@ gradient-sized norm temporaries for contiguous CUDA owner shards. Set
105117
The GB200 recipes place 748B trainable parameters on 64 GPUs, about 117 GB per
106118
GPU for BF16 weights, BF16 gradients, BF16 moments, and int16 master
107119
remainders. On 80 GB H100 parts the same recipes need twice the ranks: with
108-
`ep_size: 64` on 128 or more GPUs the mesh derives an additional `ep_shard`
120+
remainders. On 80 GB H100 parts, the same recipes need twice the ranks: with
121+
`ep_size: 64` on 128 or more GPUs, the mesh derives an additional `ep_shard`
109122
FSDP axis over the routed experts, which halves the per-rank expert state.
110123
The variants below were validated on 8-GPU H100 nodes with InfiniBand, 100
111124
updates each, using the NGC `nemo-automodel:26.08` container plus this
@@ -130,7 +143,7 @@ Three points apply to every H100 variant:
130143
- **Set `NCCL_PXN_DISABLE=1`, or `NCCL_NET_PLUGIN=none NCCL_NET=IB`,** on
131144
multi-node InfiniBand systems. With the container's default NCCL network
132145
plugin and PXN both active, the Engram row-owner all-to-all received
133-
byte-shifted token IDs from the first node's ranks and the recipe stopped in
146+
byte-shifted token IDs from the first node's ranks, and the recipe stopped in
134147
`_validate_received_ids` during the first forward, with NCCL 2.30.5 and
135148
2.30.7 alike, and with both HPC-X 2.50 plugins the image ships (the default
136149
Spectrum-X `spcx` and IBext). Either switch alone avoids it: disabling PXN
@@ -144,7 +157,7 @@ Three points apply to every H100 variant:
144157
true`, torch 2.13 hits an activation-checkpoint recompute mismatch on the
145158
expert all-gather copy-out ([issue #3958](https://github.com/NVIDIA-NeMo/Automodel/issues/3958)).
146159
`moe.reshard_after_forward: false` avoids it at about 8.5 GB per rank and
147-
becomes unnecessary once [PR #3961](https://github.com/NVIDIA-NeMo/Automodel/pull/3961)
160+
becomes unnecessary after [PR #3961](https://github.com/NVIDIA-NeMo/Automodel/pull/3961)
148161
is merged.
149162
- **Point the JIT caches at node-local storage** (`HYBRID_EP_CACHE_DIR`,
150163
`TILELANG_CACHE_DIR`, `TRITON_CACHE_DIR`). HybridEP resolves its cache from
@@ -185,7 +198,7 @@ preserves released parameter names, required FP32 weights, and logical Engram
185198
rows while removing the original quantization metadata. Quantized re-export
186199
is not supported.
187200

188-
Recipe initialization uses the NeMo AutoModel wrapper's internal Checkpointer and DCP
201+
Recipe initialization uses the NeMo AutoModel wrapper's internal Checkpointer and PyTorch Distributed Checkpoint (DCP)
189202
path to load the base checkpoint into the distributed model.
190203

191204
Text inputs use unpacked `input_ids [batch, sequence]`, right padding, and
@@ -196,7 +209,7 @@ straight-through gradients; there is no quantization toggle.
196209

197210
The released vision tower is enabled by the default model configuration and
198211
disabled by the HellaSwag recipe. `DeepseekV41Processor` supports ordinary
199-
text and image conversations and save/reload. Image inputs require unpacked
212+
text and image conversations and saving and reloading. Image inputs require unpacked
200213
`input_ids [batch, sequence]`, `pixel_values [all_patches, 3, patch_size, patch_size]`,
201214
`image_grid_hws [images, 2]`, and `vision_token_types [batch, sequence]`.
202215
The image encoder and projector reuse DeepSeek V4 modules; image delimiters
@@ -210,7 +223,7 @@ covered all 40 layers, 4,096 positions, and all 129,280 vocabulary entries
210223
against the pinned checkpoint's unchanged official `inference/` implementation.
211224
Official TP1 was compared with native EP4 and 4 Engram owners using
212225
`torch_mm`, native HybridEP dispatch and combine, and TileLang. Dense and expert
213-
reference weights were decoded for BF16 compute, with the original cache/index quantization retained;
226+
reference weights were decoded for BF16 compute, with the original cache and index quantization retained;
214227
this was not a default FP8/FP4 GEMM comparison.
215228

216229
The native execution streamed one production block at a time while carrying

‎nemo_automodel/components/models/common/utils.py‎

Lines changed: 12 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -485,9 +485,14 @@ class BackendConfig:
485485
mixing chain is compiled as well. Compiled numerics are allclose to eager but
486486
not bitwise-identical.
487487
compile_norm: torch.compile the fp32 RMSNorm chain of models that opt in
488-
(currently Kimi K3), fusing cast/pow/mean/rsqrt/mul into one kernel.
488+
(currently Kimi K3 and DeepSeek V4.1), fusing cast/pow/mean/rsqrt/mul into one kernel.
489489
Same lazy once-per-process pattern as ``compile_situ``; numerics are
490490
allclose to eager but not bitwise-identical.
491+
compile_hc: torch.compile the mHC (hyper-connection) cores of models that opt in
492+
(currently DeepSeek V4.1): coefficient projection, stream collapse and stream expand,
493+
one fused kernel each, static shapes (a new sequence length compiles them again; past
494+
dynamo's recompile limit a core runs eager). Same lazy once-per-process pattern as
495+
``compile_situ``; numerics are allclose to eager but not bitwise-identical.
491496
shared_expert_overlap: run the shared experts of opted-in MoE models (currently Kimi K3)
492497
on a side CUDA stream so their GEMMs overlap the expert-parallel dispatch / combine
493498
communication of the routed path; numerics unchanged. Default False.
@@ -552,10 +557,14 @@ class BackendConfig:
552557
# Kimi K3): fused whole-tensor weighted-SiTU forward/backward, the dense SituAndMul core
553558
# and the attn-res chain. Numerics are allclose to eager, not bitwise-identical. Default False.
554559
compile_situ: bool = False
555-
# When True, torch.compile the fp32 RMSNorm chain of opted-in models (currently Kimi K3),
556-
# same lazy once-per-process pattern as compile_situ. Numerics are allclose to eager,
560+
# When True, torch.compile the fp32 RMSNorm chain of opted-in models (currently Kimi K3 and
561+
# DeepSeek V4.1), same lazy once-per-process pattern as compile_situ. Numerics are allclose to eager,
557562
# not bitwise-identical. Default False.
558563
compile_norm: bool = False
564+
# When True, torch.compile the mHC projection / collapse / expand cores of opted-in models
565+
# (currently DeepSeek V4.1), static shapes, same lazy once-per-process pattern as compile_situ.
566+
# Numerics are allclose to eager, not bitwise-identical. Default False.
567+
compile_hc: bool = False
559568
# When True, models that opt in (currently Kimi K3) run their shared experts on a side CUDA
560569
# stream, launched before the routed-expert path and joined after it, so the shared-expert
561570
# GEMMs overlap the expert-parallel dispatch / combine communication (Megatron-Core's

‎nemo_automodel/components/models/deepseek_v41/config.py‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -183,6 +183,9 @@ def __init__(
183183
self.hc_mult = hc_mult
184184
self.hc_sinkhorn_iters = hc_sinkhorn_iters
185185
self.hc_eps = hc_eps
186+
# "torch": eager fp32 mHC math (released reference); "compile": the projection, collapse and expand
187+
# cores wrapped once per process with torch.compile (layers.compile_hc_cores) — allclose, not bitwise,
188+
# static shapes. Persisted by save_pretrained like every config field.
186189
self.engram_layer_ids = [1, 14] if engram_layer_ids is None else list(engram_layer_ids)
187190
self.engram_num_embeddings = (
188191
([384006168, 384016682] if self.engram_layer_ids else [])

‎nemo_automodel/components/models/deepseek_v41/layers.py‎

Lines changed: 100 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,100 @@
3030
from nemo_automodel.components.models.deepseek_v41.config import DeepseekV41TextConfig
3131

3232

33+
# ----------------------------------------------------------------------------------------- fp32 cores
34+
# The mHC coefficient projection, the stream collapse/expand and the fp32 RMSNorm each run several full fp32 passes
35+
# in eager mode, and ``expand`` materialises a ``[batch, sequence, streams, streams, hidden]`` fp32 temporary before
36+
# its stream sum. The arithmetic lives in the plain functions below so ``torch.compile`` can fuse each into one or
37+
# two kernels: ``BackendConfig.compile_hc`` wraps the three mHC cores, ``BackendConfig.compile_norm``
38+
# the RMSNorm core (same lazy once-per-process pattern as Kimi K3's ``compile_norm``). The modules call the
39+
# module-level dispatchers, so instances built before the hook ran still pick up the compiled core. Compiled
40+
# numerics are allclose to eager, not bitwise-identical. The mHC cores are compiled for static shapes only: with
41+
# dynamic shapes inductor keeps the broadcast temporary of ``expand`` and the fusion is lost, so every new sequence
42+
# length compiles them again. The RMSNorm core is compiled per shape as well: the dynamic-shape variant lost 15 % per
43+
# call and shifted the mini benchmark's first-step loss by 9e-4, the static kernels reproduce the eager loss.
44+
def _rms_norm_core(hidden_states: torch.Tensor, weight: torch.Tensor, eps: float) -> torch.Tensor:
45+
value = hidden_states.float()
46+
value = value * torch.rsqrt(value.square().mean(-1, keepdim=True) + eps)
47+
return (weight.float() * value).to(hidden_states.dtype)
48+
49+
50+
def _hc_project_core(
51+
hidden_states: torch.Tensor,
52+
fn: torch.Tensor,
53+
base: torch.Tensor,
54+
scale: torch.Tensor,
55+
streams: int,
56+
norm_eps: float,
57+
eps: float,
58+
) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
59+
"""Coefficient projection up to the Sinkhorn input: pre, post and the [streams, streams] combination logits."""
60+
flat = hidden_states.flatten(2).float()
61+
mixes = F.linear(flat, fn.float()) * torch.rsqrt(flat.square().mean(-1, keepdim=True) + norm_eps)
62+
scales = torch.cat((scale[0].expand(streams), scale[1].expand(streams), scale[2].expand(streams * streams)))
63+
# The released kernel fuses the affine multiply/add before sigmoid and
64+
# Sinkhorn. Separate operations round its FP32 coefficients differently;
65+
# these differences can survive the subsequent BF16 stream collapse. This
66+
# holds on the eager path; the compiled core decomposes addcmul and is
67+
# allclose only (see the module note above).
68+
logits = torch.addcmul(base, mixes, scales)
69+
pre = torch.sigmoid(logits[..., :streams]) + eps
70+
post = 2 * torch.sigmoid(logits[..., streams : 2 * streams])
71+
return pre, post, logits[..., 2 * streams :].unflatten(-1, (streams, streams))
72+
73+
74+
def _hc_collapse_core(hidden_states: torch.Tensor, pre_mix: torch.Tensor) -> torch.Tensor:
75+
return (pre_mix.unsqueeze(-1) * hidden_states.float()).sum(2).to(hidden_states.dtype)
76+
77+
78+
def _hc_expand_core(
79+
output: torch.Tensor, residual: torch.Tensor, post: torch.Tensor, comb: torch.Tensor
80+
) -> torch.Tensor:
81+
update = post.unsqueeze(-1) * output.unsqueeze(-2)
82+
mixed = (comb.unsqueeze(-1) * residual.unsqueeze(-2)).sum(2)
83+
return (update + mixed).to(output.dtype)
84+
85+
86+
_rms_norm = _rms_norm_core
87+
_hc_project = _hc_project_core
88+
_hc_collapse = _hc_collapse_core
89+
_hc_expand = _hc_expand_core
90+
_NORM_COMPILED = False
91+
_HC_COMPILED = False
92+
93+
94+
def compile_norm_core() -> None:
95+
"""Wrap the fp32 RMSNorm core with ``torch.compile`` (``BackendConfig.compile_norm``); runs once per process.
96+
97+
Static shapes: one specialization per norm site (block, attention, compressor and indexer norms) and grad mode.
98+
The dynamic-shape variant (Kimi K3's setting) was measured 15 % slower per call and moved the first-step loss of
99+
the 12-layer mini benchmark by 9e-4, while the static kernels reproduced the eager loss to four decimals; no
100+
``fullgraph``, so a shape past dynamo's recompile limit falls back to eager instead of raising.
101+
"""
102+
global _rms_norm, _NORM_COMPILED
103+
if _NORM_COMPILED:
104+
return
105+
_rms_norm = torch.compile(_rms_norm_core, dynamic=False)
106+
_NORM_COMPILED = True
107+
108+
109+
def compile_hc_cores() -> None:
110+
"""Wrap the mHC projection, collapse and expand cores with ``torch.compile`` (``BackendConfig.compile_hc``).
111+
112+
Runs once per process, static shapes (``dynamic=False``: the fusion of ``expand`` needs concrete shapes; a new
113+
sequence length compiles the three cores again). No ``fullgraph``: past dynamo's recompile limit (many distinct
114+
sequence lengths) a core runs eager instead of raising ``FailOnRecompileLimitHit``; the unit tests assert the
115+
cores compile without graph breaks. The Sinkhorn normalisation keeps its own backend (TileKernels or torch):
116+
compiling the 20-iteration torch loop is slower than the fused TileKernels kernel.
117+
"""
118+
global _hc_project, _hc_collapse, _hc_expand, _HC_COMPILED
119+
if _HC_COMPILED:
120+
return
121+
_hc_project = torch.compile(_hc_project_core, dynamic=False)
122+
_hc_collapse = torch.compile(_hc_collapse_core, dynamic=False)
123+
_hc_expand = torch.compile(_hc_expand_core, dynamic=False)
124+
_HC_COMPILED = True
125+
126+
33127
class DeepseekV41RMSNorm(nn.Module):
34128
"""Normalize in FP32 and multiply the scale before casting the result."""
35129

@@ -47,9 +141,7 @@ def forward(self, hidden_states: torch.Tensor) -> torch.Tensor:
47141
Returns:
48142
Tensor of shape [..., hidden], with the input dtype.
49143
"""
50-
value = hidden_states.float()
51-
value = value * torch.rsqrt(value.square().mean(-1, keepdim=True) + self.eps)
52-
return (self.weight.float() * value).to(hidden_states.dtype)
144+
return _rms_norm(hidden_states, self.weight, self.eps)
53145

54146

55147
class DeepseekV41Mix(NamedTuple):
@@ -97,24 +189,10 @@ def forward(self, hidden_states: torch.Tensor) -> DeepseekV41Mix:
97189
Coefficients with pre/post tensors of shape [batch, sequence, streams]
98190
and comb of shape [batch, sequence, streams, streams]. All use FP32.
99191
"""
100-
flat = hidden_states.flatten(2).float()
101-
mixes = F.linear(flat, self.fn.float()) * torch.rsqrt(flat.square().mean(-1, keepdim=True) + self.norm_eps)
102-
streams = self.streams
103-
scales = torch.cat(
104-
(self.scale[0].expand(streams), self.scale[1].expand(streams), self.scale[2].expand(streams * streams))
105-
)
106-
# The released kernel fuses the affine multiply/add before sigmoid and
107-
# Sinkhorn. Separate operations round its FP32 coefficients differently;
108-
# these differences can survive the subsequent BF16 stream collapse.
109-
logits = torch.addcmul(self.base, mixes, scales)
110-
pre = torch.sigmoid(logits[..., :streams]) + self.eps
111-
post = 2 * torch.sigmoid(logits[..., streams : 2 * streams])
112-
comb = dsv4_sinkhorn_normalize(
113-
logits[..., 2 * streams :].unflatten(-1, (streams, streams)),
114-
backend=self.sinkhorn_backend,
115-
repeat=self.iterations,
116-
eps=self.eps,
192+
pre, post, comb_logits = _hc_project(
193+
hidden_states, self.fn, self.base, self.scale, self.streams, self.norm_eps, self.eps
117194
)
195+
comb = dsv4_sinkhorn_normalize(comb_logits, backend=self.sinkhorn_backend, repeat=self.iterations, eps=self.eps)
118196
return DeepseekV41Mix(pre, post, comb)
119197

120198
@staticmethod
@@ -128,7 +206,7 @@ def collapse(hidden_states: torch.Tensor, pre_mix: torch.Tensor) -> torch.Tensor
128206
Returns:
129207
Tensor of shape [batch, sequence, hidden], with the input dtype.
130208
"""
131-
return (pre_mix.unsqueeze(-1) * hidden_states.float()).sum(2).to(hidden_states.dtype)
209+
return _hc_collapse(hidden_states, pre_mix)
132210

133211
@staticmethod
134212
def expand(output: torch.Tensor, residual: torch.Tensor, mix: DeepseekV41Mix) -> torch.Tensor:
@@ -143,6 +221,4 @@ def expand(output: torch.Tensor, residual: torch.Tensor, mix: DeepseekV41Mix) ->
143221
Returns:
144222
Tensor of shape [batch, sequence, streams, hidden], with output's dtype.
145223
"""
146-
update = mix.post.unsqueeze(-1) * output.unsqueeze(-2)
147-
mixed = (mix.comb.unsqueeze(-1) * residual.unsqueeze(-2)).sum(2)
148-
return (update + mixed).to(output.dtype)
224+
return _hc_expand(output, residual, mix.post, mix.comb)

‎nemo_automodel/components/models/deepseek_v41/model.py‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -79,6 +79,8 @@
7979
from nemo_automodel.components.models.deepseek_v41.layers import (
8080
DeepseekV41HyperConnection,
8181
DeepseekV41RMSNorm,
82+
compile_hc_cores,
83+
compile_norm_core,
8284
)
8385
from nemo_automodel.components.models.deepseek_v41.packing import packed_layout
8486
from nemo_automodel.components.models.deepseek_v41.processing import (
@@ -237,6 +239,10 @@ def __init__(
237239
config = config.text_config
238240
self.config = config
239241
self.moe_config = moe_config
242+
if backend.compile_hc:
243+
compile_hc_cores()
244+
if backend.compile_norm and backend.rms_norm != "te":
245+
compile_norm_core()
240246
dtype = dtype_from_str(config.dtype, torch.bfloat16)
241247
self.embed_tokens = nn.Embedding(config.vocab_size, config.hidden_size, dtype=dtype)
242248
active_engram = any(i < config.num_hidden_layers for i in config.engram_layer_ids)

0 commit comments

Comments
 (0)