diff --git a/grpo/nemotron-nano-9b-v2/README.md b/grpo/nemotron-nano-9b-v2/README.md
new file mode 100644
index 0000000..e8cc7cd
--- /dev/null
+++ b/grpo/nemotron-nano-9b-v2/README.md
@@ -0,0 +1,190 @@
+# NVIDIA Nemotron Nano 9B v2 GRPO Ascend Recipe
+
+对应任务：[verl-ascend-recipe #70](https://github.com/verl-project/verl-ascend-recipe/issues/70)
+
+本目录提供 NVIDIA Nemotron Nano 9B v2 在 Ascend NPU 上执行 GRPO 训练的可复现配置。训练侧使用 Megatron，rollout 侧使用 vLLM-Ascend。验收运行使用 8 张 Ascend 910B3，从初始权重连续训练 13 小时 00 分 09 秒，完成 77 个 optimizer steps。
+
+## 训练链路
+
+训练通过 `verl.trainer.main_ppo` 统一调度：
+
+```text
+DAPO-Math prompts -> vLLM-Ascend TP8 async rollout (16 responses/prompt)
+                  -> DAPO rule reward -> GRPO advantage
+                  -> Megatron TP8 actor update -> rollout weight synchronization
+```
+
+实现沿用 `verl` 的 [Nemotron Nano Megatron GRPO 示例](https://github.com/verl-project/verl/blob/main/examples/grpo_trainer/run_nemotron_nano_v3_30b_a3b_megatron.sh)。针对 Nemotron Nano 9B v2 的 Mamba2 层，脚本通过 `VERL_USE_EXTERNAL_MODULES` 加载运行时临时模块：训练与 prefill 使用 MindSpeed-LLM 的 NPU SSD 实现，单 token decode 使用可被 `FULL_DECODE_ONLY` 图模式捕获的 NPU tensor 路径。兼容逻辑全部包含在启动脚本中，不需要额外的 recipe 源码 patch。
+
+## 文件
+
+| 路径 | 内容 |
+| --- | --- |
+| `run_nemotron_nano_9b_v2_grpo_megatron_npu.sh` | 8 NPU 环境、Mamba2 运行时适配、Megatron/vLLM-Ascend 参数与训练入口 |
+| `assets/reward_curve.png` | reward 逐步值与 10-step moving average |
+| `assets/loss_curve.png` | actor loss 逐步值与 10-step moving average |
+| `README.md` | 环境、数据、复现配置、指标定义与验收结果 |
+
+## 环境
+
+验收环境使用 `verl/main` Ascend vLLM + Megatron 官方依赖组合：
+
+| 组件 | 版本或配置 |
+| --- | --- |
+| NPU | 8 x Ascend 910B3 |
+| OS | Ubuntu 22.04 |
+| HDK / CANN | 26.0.rc1 / 9.1.0 |
+| Python | 3.12 |
+| PyTorch / torch_npu | 2.10.0 / 2.10.0.post4 |
+| transformers | 5.10.4 |
+| vLLM / vLLM-Ascend | 0.23.0 / 0.23.0 |
+| Triton / Triton-Ascend | 3.5.0 / 3.2.2 |
+| Megatron Core / Megatron Bridge | 0.18.3 / 0.5.0 |
+| MindSpeed / MindSpeed-LLM | 0.18.0 / 26.0.0.dev0 |
+| verl | `main` |
+
+按照 [`verl` Ascend 安装指南](https://github.com/verl-project/verl/blob/main/docs/ascend_tutorial/zh/get_start/install_guidance.rst)创建独立 Conda 环境：
+
+```bash
+source /usr/local/Ascend/ascend-toolkit/set_env.sh
+source /usr/local/Ascend/nnal/atb/set_env.sh
+
+conda create -n verl-vllm-npu python=3.12 -y
+conda activate verl-vllm-npu
+
+git clone --recursive https://github.com/verl-project/verl.git
+bash verl/scripts/install_vllm_mcore_npu.sh
+```
+
+在同一环境按照 [MindSpeed-LLM 官方安装文档](https://github.com/Ascend/MindSpeed-LLM/blob/master/docs/zh/pytorch/training/install_guide.md)安装 MindSpeed-LLM，然后验证关键组件：
+
+```bash
+python -c 'import torch_npu, vllm, mindspeed, mindspeed_llm'
+```
+
+## 模型与数据
+
+下载模型：
+
+```bash
+hf download nvidia/NVIDIA-Nemotron-Nano-9B-v2 \
+  --local-dir /path/to/NVIDIA-Nemotron-Nano-9B-v2
+```
+
+下载训练集与验证集：
+
+```bash
+hf download BytedTsinghua-SIA/DAPO-Math-17k \
+  --repo-type dataset \
+  --local-dir /path/to/dapo-math-17k
+
+hf download BytedTsinghua-SIA/AIME-2024 \
+  --repo-type dataset \
+  --local-dir /path/to/aime-2024
+```
+
+本次运行使用：
+
+- 模型：[nvidia/NVIDIA-Nemotron-Nano-9B-v2](https://huggingface.co/nvidia/NVIDIA-Nemotron-Nano-9B-v2)
+- 训练集：[BytedTsinghua-SIA/DAPO-Math-17k](https://huggingface.co/datasets/BytedTsinghua-SIA/DAPO-Math-17k) 的 `data/dapo-math-17k.parquet`
+- 验证集：[BytedTsinghua-SIA/AIME-2024](https://huggingface.co/datasets/BytedTsinghua-SIA/AIME-2024) 的 `data/aime-2024.parquet`
+- reward manager：`dapo`
+
+## 运行
+
+从 `verl` 仓库根目录启动 8 卡训练：
+
+```bash
+ASCEND_RT_VISIBLE_DEVICES=0,1,2,3,4,5,6,7 \
+RAY_DATA_HOME=/path/to/workspace \
+MODEL_PATH=/path/to/NVIDIA-Nemotron-Nano-9B-v2 \
+TRAIN_FILE=/path/to/dapo-math-17k/data/dapo-math-17k.parquet \
+TEST_FILE=/path/to/aime-2024/data/aime-2024.parquet \
+CKPTS_DIR=/path/to/checkpoints/nemotron-nano-9b-v2-grpo \
+TOTAL_TRAINING_STEPS=100 \
+SAVE_FREQ=100 \
+TEST_FREQ=-1 \
+RESUME_MODE=disable \
+VERL_USE_UV=0 \
+bash /path/to/verl-ascend-recipe/grpo/nemotron-nano-9b-v2/run_nemotron_nano_9b_v2_grpo_megatron_npu.sh
+```
+
+脚本接受额外 Hydra overrides，可直接追加在命令末尾。
+
+## 关键配置
+
+| 参数 | 验收值 |
+| --- | ---: |
+| train batch size | 32 |
+| responses per prompt | 16 |
+| 每步 rollout trajectories | 512 |
+| PPO mini / micro batch size | 32 / 8 |
+| reference / rollout log-prob micro batch size | 4 / 4 |
+| max prompt / response length | 2048 / 1024 |
+| overlong buffer length | 512 |
+| actor / rollout | Megatron / vLLM-Ascend async |
+| train / rollout TP | 8 / 8 |
+| PP / EP / ETP | 1 / 1 / 1 |
+| rollout max sequences | 512 |
+| rollout memory utilization | 0.40 |
+| chunked prefill / prefix caching | true / false |
+| dtype | BF16 |
+| dynamic batch size | false |
+| parameter / optimizer offload | false / true |
+| actor learning rate / warmup steps | 1e-6 / 10 |
+| graph mode | `FULL_DECODE_ONLY` |
+| graph capture sizes | 1, 2, 4, 8, 16, 32, 64, 128, 256, 512 |
+| save / test frequency | 100 / -1 |
+| 目标训练步数 | 100 |
+
+周期验证关闭，因此 TPS 与 step time 不包含额外 evaluation 时间。
+
+## Checkpoint 与日志
+
+脚本默认 `SAVE_FREQ=100`、`RESUME_MODE=disable`，仅在第 100 步保存最终 checkpoint，不写入中间 checkpoint。调试续训时可以显式设置：
+
+```bash
+CKPTS_DIR=/path/to/checkpoints/nemotron-nano-9b-v2-grpo \
+SAVE_FREQ=100 \
+RESUME_MODE=auto \
+bash /path/to/verl-ascend-recipe/grpo/nemotron-nano-9b-v2/run_nemotron_nano_9b_v2_grpo_megatron_npu.sh
+```
+
+正式验收数据来自 `RESUME_MODE=disable` 的 fresh run。训练指标由 console logger 输出；完整脱敏训练日志：[nemotron_nano_9b_v2_grpo_8npu.log](https://gist.github.com/OnPathXD/cedd7535f7d5a996fca748e7e94cc7da)。
+
+## 指标定义
+
+| 指标 | 计算方式 |
+| --- | --- |
+| reward | `critic/rewards/mean`；比较首 10 步和末 10 步均值 |
+| actor loss | `actor/loss` 的逐步值、全程均值和范围 |
+| 单 NPU TPS | `perf/throughput` 的 77 步算术均值 |
+| 8 NPU TPS | 单 NPU TPS 乘以 8 |
+| step time | `timing_s/step` 的 77 步算术均值 |
+
+## 13 小时验收结果
+
+在 8 x Ascend 910B3 上从初始权重连续训练 13 小时 00 分 09 秒，共完成 77 个 optimizer steps：
+
+| 指标 | 结果 |
+| --- | ---: |
+| 训练时长 / 完成步数 | 13 小时 00 分 09 秒 / 77 |
+| reward 首 10 步均值 | -0.946484 |
+| reward 末 10 步均值 | -0.192188 |
+| reward 绝对增量 | +0.754297 |
+| reward 全程线性斜率 | +0.011753 / step |
+| actor loss 均值 / 范围 | 0.018983 / 0.001587 至 0.062556 |
+| TPS | 935.5607 tokens/s（8 NPU），116.9451 tokens/s/NPU |
+| 首 10 步单 NPU TPS 均值 / 最低值 | 124.4040 / 109.5558 tokens/s/NPU |
+| 平均 step 时间 / 范围 | 599.9649 s / 511.9716 至 693.4601 s |
+| 处理 token 总数 | 43,234,521 |
+
+reward 的首尾窗口增量和全程线性斜率均为正，平均单 NPU TPS 高于 100 tokens/s。
+
+![Reward curve](assets/reward_curve.png)
+
+Reward 曲线展示逐步原始值及 10-step moving average。
+
+![Actor loss curve](assets/loss_curve.png)
+
+Actor loss 曲线展示逐步原始值及 10-step moving average。
diff --git a/grpo/nemotron-nano-9b-v2/assets/loss_curve.png b/grpo/nemotron-nano-9b-v2/assets/loss_curve.png
new file mode 100644
index 0000000..9ee2092
Binary files /dev/null and b/grpo/nemotron-nano-9b-v2/assets/loss_curve.png differ
diff --git a/grpo/nemotron-nano-9b-v2/assets/reward_curve.png b/grpo/nemotron-nano-9b-v2/assets/reward_curve.png
new file mode 100644
index 0000000..57fbbb8
Binary files /dev/null and b/grpo/nemotron-nano-9b-v2/assets/reward_curve.png differ
diff --git a/grpo/nemotron-nano-9b-v2/run_nemotron_nano_9b_v2_grpo_megatron_npu.sh b/grpo/nemotron-nano-9b-v2/run_nemotron_nano_9b_v2_grpo_megatron_npu.sh
new file mode 100755
index 0000000..9c0208e
--- /dev/null
+++ b/grpo/nemotron-nano-9b-v2/run_nemotron_nano_9b_v2_grpo_megatron_npu.sh
@@ -0,0 +1,575 @@
+#!/usr/bin/env bash
+set -xeo pipefail
+
+################################################### environment ###################################################
+# Adapted from verl's canonical Nemotron Nano v3 Megatron recipe:
+# examples/grpo_trainer/run_nemotron_nano_v3_30b_a3b_megatron.sh
+# Run in an official verl NPU environment with CANN, torch-npu,
+# vLLM-Ascend, Megatron-LM, MindSpeed, and Megatron-Bridge installed.
+
+if [[ -n "${CONDA_PREFIX:-}" ]]; then
+    export LD_LIBRARY_PATH="${CONDA_PREFIX}/lib:${LD_LIBRARY_PATH:-}"
+fi
+if [[ -f /usr/local/Ascend/ascend-toolkit/set_env.sh ]]; then
+    source /usr/local/Ascend/ascend-toolkit/set_env.sh
+fi
+if [[ -f /usr/local/Ascend/nnal/atb/set_env.sh ]]; then
+    source /usr/local/Ascend/nnal/atb/set_env.sh
+fi
+if [[ -n "${CONDA_PREFIX:-}" ]]; then
+    export LD_LIBRARY_PATH="${CONDA_PREFIX}/lib:${LD_LIBRARY_PATH:-}"
+fi
+set -u
+
+export DEVICE=npu
+export VLLM_USE_V1=1
+export TOKENIZERS_PARALLELISM=false
+export PYTHONUNBUFFERED=1
+export RAY_EXPERIMENTAL_NOSET_ASCEND_RT_VISIBLE_DEVICES=1
+export ASCEND_RT_VISIBLE_DEVICES=${ASCEND_RT_VISIBLE_DEVICES:-"0,1,2,3,4,5,6,7"}
+export TRAINING_BACKEND=verl
+
+# vLLM 0.23.0's Mamba prefill Triton kernel is not supported by the official
+# Triton-Ascend 3.2.2 wheel. MindSpeed-LLM provides the official pure PyTorch
+# NPU SSD implementation. This temporary module bridges the prefill API and
+# provides a graph-safe NPU tensor implementation for one-token Mamba decode;
+# it does not modify an installed package.
+# Install MindSpeed-LLM from https://github.com/Ascend/MindSpeed-LLM first.
+mamba_npu_compat_dir=$(mktemp -d "${TMPDIR:-/tmp}/verl-mamba-npu.XXXXXX")
+cat > "${mamba_npu_compat_dir}/verl_mamba_npu_compat.py" <<'PY'
+import torch
+import torch_npu
+import triton
+
+if not hasattr(triton, "set_allocator"):
+    triton.set_allocator = lambda allocator: None
+torch.cuda.device = torch.npu.device
+torch.cuda.get_device_properties = torch.npu.get_device_properties
+
+import mindspeed.megatron_adaptor
+import mindspeed_llm.tasks.models.ssm.state_space_duality as ssd
+
+
+class _SingleContextParallel:
+    get_context_parallel_world_size = staticmethod(lambda: 1)
+    get_context_parallel_rank = staticmethod(lambda: 0)
+    get_context_parallel_group = staticmethod(lambda: None)
+
+
+# The recipe fixes Megatron context parallelism at one. Rollout workers do not
+# initialize Megatron parallel state, so isolate that fact to this SSD module.
+ssd.mpu = _SingleContextParallel()
+
+
+def _run_ssd(x, dt, A, B, C, chunk_size, D, dt_bias, initial_states,
+             dt_softplus, dt_limit, return_state):
+    if not dt_softplus or D is None or D.ndim != 1 or dt_bias is None:
+        raise NotImplementedError("This NPU adapter supports the Nemotron Mamba2 layout only")
+    options = ssd.StateOptions(
+        initial_states=initial_states, return_final_state=return_state
+    )
+    processor = ssd.StateSpaceProcessor(
+        {
+            "nheads_local": x.shape[2],
+            "ngroups_local": B.shape[2],
+            "dt_min": dt_limit[0],
+            "dt_max": dt_limit[1],
+            "dt_bias": dt_bias,
+            "headdim": x.shape[3],
+            "d_state": B.shape[3],
+            "chunk_size": chunk_size,
+            "D_has_hdim": False,
+        }
+    )
+    result = processor.process(
+        ssd.ProcessInputs(
+            x=x.flatten(2),
+            dt=dt,
+            A=A,
+            B=B.flatten(2),
+            C=C.flatten(2),
+            D=D,
+        ),
+        options,
+    )
+    if return_state:
+        y, state = result
+        return y.to(x.dtype), state
+    return result.to(x.dtype), None
+
+
+def _mcore_scan(x, dt, A, B, C, chunk_size, D=None, z=None, dt_bias=None,
+                initial_states=None, seq_idx=None, cu_seqlens=None,
+                dt_softplus=False, dt_limit=(0.0, float("inf")),
+                return_final_states=False, return_varlen_states=False,
+                state_dtype=None):
+    if z is not None or seq_idx is not None or cu_seqlens is not None:
+        raise NotImplementedError("Packed Megatron Mamba training is not enabled")
+    if return_varlen_states:
+        raise NotImplementedError("Variable-length Megatron state return is not enabled")
+    y, state = _run_ssd(
+        x, dt, A, B, C, chunk_size, D, dt_bias, initial_states,
+        dt_softplus, dt_limit, return_final_states,
+    )
+    return (y, state) if return_final_states else y
+
+
+_plan_input = None
+_plan = None
+
+
+def _varlen_plan(cu_seqlens):
+    global _plan_input, _plan
+    if cu_seqlens is _plan_input:
+        return _plan
+    boundaries = cu_seqlens.detach().cpu().tolist()
+    length_values = [
+        end - start for start, end in zip(boundaries, boundaries[1:])
+    ]
+    if not length_values or min(length_values) <= 0:
+        raise ValueError("Mamba prefill requires non-empty sequences")
+    starts = torch.tensor(boundaries[:-1], device=cu_seqlens.device)
+    lengths = torch.tensor(length_values, device=cu_seqlens.device)
+    positions = torch.arange(max(length_values), device=cu_seqlens.device)
+    gather = starts[:, None] + positions[None, :]
+    valid = positions[None, :] < lengths[:, None]
+    _plan_input, _plan = cu_seqlens, (gather, valid)
+    return _plan
+
+
+def _vllm_varlen(x, dt, A, B, C, chunk_size, cu_seqlens,
+                 cu_chunk_seqlens, last_chunk_indices, seq_idx, out, D=None,
+                 z=None, dt_bias=None, initial_states=None, dt_softplus=False,
+                 dt_limit=(0.0, float("inf")), return_intermediate_states=False,
+                 state_dtype=None):
+    if z is not None or return_intermediate_states:
+        raise NotImplementedError("Mamba prefix caching is disabled by this recipe")
+    gather, valid = _varlen_plan(cu_seqlens)
+    safe = gather.clamp(max=x.shape[0] - 1)
+    y, states = _run_ssd(
+        torch.where(valid[..., None, None], x[safe], 0),
+        torch.where(valid[..., None], dt[safe], -100.0),
+        A,
+        torch.where(valid[..., None, None], B[safe], 0),
+        torch.where(valid[..., None, None], C[safe], 0),
+        chunk_size, D, dt_bias, initial_states, dt_softplus, dt_limit, True,
+    )
+    out.index_copy_(0, gather[valid], y[valid].to(out.dtype))
+    return states.to(state_dtype or C.dtype)
+
+
+def _npu_selective_state_update(
+    state,
+    x,
+    dt,
+    A,
+    B,
+    C,
+    D,
+    dt_bias,
+    z=None,
+    dt_softplus=False,
+    state_batch_indices=None,
+    dst_state_batch_indices=None,
+    null_block_id=-1,
+    out=None,
+    num_accepted_tokens=None,
+    cu_seqlens=None,
+    is_blackwell=False,
+):
+    """NPU implementation of vLLM's non-speculative Mamba2 decode step."""
+    import torch.nn.functional as F
+
+    del is_blackwell
+    if state.dim() == 3:
+        state = state.unsqueeze(1)
+    if x.dim() == 2:
+        x = x.unsqueeze(1)
+    if dt.dim() == 2:
+        dt = dt.unsqueeze(1)
+    if A.dim() == 2:
+        A = A.unsqueeze(0)
+    if B.dim() == 2:
+        B = B.unsqueeze(1)
+    if C.dim() == 2:
+        C = C.unsqueeze(1)
+    if D.dim() == 1:
+        D = D.unsqueeze(0)
+    if z is not None and z.dim() == 2:
+        z = z.unsqueeze(1)
+    if dt_bias.dim() == 1:
+        dt_bias = dt_bias.unsqueeze(0)
+    if out is None:
+        raise ValueError("vLLM Mamba2 decode requires a preallocated output")
+    if out.dim() == 2:
+        out = out.unsqueeze(1)
+    if state_batch_indices is not None and state_batch_indices.dim() == 1:
+        state_batch_indices = state_batch_indices.unsqueeze(1)
+    if dst_state_batch_indices is not None and dst_state_batch_indices.dim() == 1:
+        dst_state_batch_indices = dst_state_batch_indices.unsqueeze(1)
+
+    cache_batch, nheads, headdim, dstate = state.shape
+    token_count = x.shape[0]
+    sequence_count = len(cu_seqlens) - 1 if cu_seqlens is not None else token_count
+    if token_count != sequence_count:
+        raise NotImplementedError(
+            "The issue-70 recipe supports vLLM's non-speculative one-token decode"
+        )
+    if x.shape != (token_count, nheads, headdim) or dt.shape != x.shape:
+        raise ValueError("Unexpected Mamba2 decode x/dt shapes")
+    if A.shape != (nheads, headdim, dstate):
+        raise ValueError("Unexpected Mamba2 decode A shape")
+    ngroups = B.shape[1]
+    if B.shape != (token_count, ngroups, dstate) or C.shape != B.shape:
+        raise ValueError("Unexpected Mamba2 decode B/C shapes")
+    if nheads % ngroups or D.shape != (nheads, headdim):
+        raise ValueError("Unexpected Mamba2 decode head/group shapes")
+    if dt_bias.shape != (nheads, headdim) or out.shape != x.shape:
+        raise ValueError("Unexpected Mamba2 decode bias/output shapes")
+    if z is not None and z.shape != x.shape:
+        raise ValueError("Unexpected Mamba2 decode gate shape")
+    if num_accepted_tokens is not None and num_accepted_tokens.shape != (sequence_count,):
+        raise ValueError("Unexpected Mamba2 accepted-token metadata shape")
+
+    if state_batch_indices is None:
+        source_indices = torch.arange(sequence_count, device=x.device)
+    else:
+        source_indices = state_batch_indices[:, 0].to(device=x.device, dtype=torch.long)
+    if dst_state_batch_indices is None:
+        destination_indices = source_indices
+    else:
+        destination_indices = dst_state_batch_indices[:, 0].to(
+            device=x.device, dtype=torch.long
+        )
+    if source_indices.numel() != sequence_count or destination_indices.numel() != sequence_count:
+        raise ValueError("Unexpected Mamba2 state-index metadata shape")
+
+    valid = (source_indices != null_block_id) & (destination_indices != null_block_id)
+    safe_sources = source_indices.clamp(min=0, max=cache_batch - 1)
+    current_state = state.index_select(0, safe_sources).float()
+    current_state = current_state * valid[:, None, None, None]
+
+    heads_per_group = nheads // ngroups
+    b_heads = B.repeat_interleave(heads_per_group, dim=1).float()
+    c_heads = C.repeat_interleave(heads_per_group, dim=1).float()
+
+    # Nemotron expands these tensors over the head dimension. Match vLLM's
+    # tied-head fast path without materializing the repeated values.
+    tie_hdim = (
+        A.stride(-1) == 0
+        and A.stride(-2) == 0
+        and dt.stride(-1) == 0
+        and dt_bias.stride(-1) == 0
+    )
+    if tie_hdim:
+        dt_scalar = dt[..., 0].float() + dt_bias[..., 0].float().unsqueeze(0)
+        if dt_softplus:
+            dt_scalar = F.softplus(dt_scalar)
+        decay = torch.exp(A[:, 0, 0].float().unsqueeze(0) * dt_scalar)
+        x_discrete = x.float() * dt_scalar.unsqueeze(-1)
+        next_state = current_state * decay[:, :, None, None]
+        next_state = next_state + x_discrete.unsqueeze(-1) * b_heads.unsqueeze(2)
+    else:
+        dt_fp32 = dt.float() + dt_bias.float().unsqueeze(0)
+        if dt_softplus:
+            dt_fp32 = F.softplus(dt_fp32)
+        decay = torch.exp(A.float().unsqueeze(0) * dt_fp32.unsqueeze(-1))
+        next_state = current_state * decay
+        next_state = next_state + (
+            x.float().unsqueeze(-1) * dt_fp32.unsqueeze(-1) * b_heads.unsqueeze(2)
+        )
+
+    output = (next_state * c_heads.unsqueeze(2)).sum(dim=-1)
+    output = output + x.float() * D.float().unsqueeze(0)
+    if z is not None:
+        output = output * F.silu(z.float())
+    output = output * valid[:, None, None]
+    out.copy_(output.to(out.dtype))
+
+    # Avoid aclnnNonzero, which synchronizes the stream and cannot be captured
+    # in an Ascend graph. Padded rows use cache row zero, restored after write.
+    safe_destinations = destination_indices.clamp(min=0, max=cache_batch - 1)
+    row_zero_before = state[0].clone()
+    row_zero_mask = valid & (destination_indices == 0)
+    row_zero_candidate = (
+        next_state * row_zero_mask[:, None, None, None]
+    ).sum(dim=0)
+    row_zero_after = torch.where(
+        row_zero_mask.any(), row_zero_candidate, row_zero_before
+    )
+    state.index_copy_(0, safe_destinations, next_state.to(state.dtype))
+    state[0].copy_(row_zero_after.to(state.dtype))
+
+
+from megatron.core.ssm import mamba_mixer as mcore_mamba
+from vllm.model_executor.layers.mamba import mamba_mixer2 as vllm_mamba
+from vllm.model_executor.layers.mamba.ops import ssd_combined as vllm_ssd
+from vllm.model_executor.layers.mamba.ops import ssu_dispatch as vllm_ssu
+
+mcore_mamba.mamba_chunk_scan_combined = _mcore_scan
+vllm_ssd.mamba_chunk_scan_combined_varlen = _vllm_varlen
+vllm_mamba.mamba_chunk_scan_combined_varlen = _vllm_varlen
+vllm_ssu.selective_state_update = _npu_selective_state_update
+vllm_mamba.selective_state_update = _npu_selective_state_update
+
+print("[issue-70] enabled official MindSpeed SSD and NPU tensor SSU on NPU")
+PY
+export PYTHONPATH="${mamba_npu_compat_dir}:${PYTHONPATH:-}"
+if [[ -n "${VERL_USE_EXTERNAL_MODULES:-}" ]]; then
+    export VERL_USE_EXTERNAL_MODULES="${VERL_USE_EXTERNAL_MODULES},verl_mamba_npu_compat"
+else
+    export VERL_USE_EXTERNAL_MODULES=verl_mamba_npu_compat
+fi
+
+IFS=',' read -r -a visible_devices <<< "${ASCEND_RT_VISIBLE_DEVICES}"
+if [[ ${#visible_devices[@]} -ne 8 ]]; then
+    echo "This recipe requires exactly 8 visible NPUs; got ${ASCEND_RT_VISIBLE_DEVICES}" >&2
+    exit 2
+fi
+
+################################################### quick config ###################################################
+
+# ---- user-adjustable ----
+rollout_mode=${ROLLOUT_MODE:-"async"}
+return_raw_chat="False"
+rollout_name="vllm"
+if [ "$rollout_mode" = "async" ]; then
+    export VLLM_USE_V1=1
+    return_raw_chat="True"
+fi
+dtype="bfloat16"
+
+project_name='DAPO'
+exp_name=${EXP_NAME:-'nemotron_nano_9b_v2_grpo_npu'}
+
+adv_estimator=grpo
+
+use_kl_in_reward=False
+kl_coef=0.0
+use_kl_loss=False
+kl_loss_coef=0.0
+
+clip_ratio_low=0.2
+clip_ratio_high=0.28
+
+max_prompt_length=${MAX_PROMPT_LENGTH:-2048}
+max_response_length=${MAX_RESPONSE_LENGTH:-1024}
+enable_overlong_buffer=True
+overlong_buffer_len=${OVERLONG_BUFFER_LEN:-512}
+overlong_penalty_factor=1.0
+
+loss_agg_mode="token-mean"
+
+train_prompt_bsz=${TRAIN_PROMPT_BSZ:-32}
+n_resp_per_prompt=${N_RESP_PER_PROMPT:-16}
+train_prompt_mini_bsz=${PPO_MINI_BATCH_SIZE:-32}
+
+# Ray
+RAY_ADDRESS=${RAY_ADDRESS:-"http://localhost:8265"}
+WORKING_DIR=${WORKING_DIR:-"${PWD}"}
+RUNTIME_ENV=${RUNTIME_ENV:-"${WORKING_DIR}/verl/verl/trainer/runtime_env.yaml"}
+NNODES=${NNODES:-1}
+# Paths
+RAY_DATA_HOME=${RAY_DATA_HOME:?Set RAY_DATA_HOME to the model, data, and checkpoint storage root}
+MODEL_PATH=${MODEL_PATH:-"${RAY_DATA_HOME}/models/NVIDIA-Nemotron-Nano-9B-v2"}
+CKPTS_DIR=${CKPTS_DIR:-"${RAY_DATA_HOME}/ckpts/${project_name}/${exp_name}"}
+TRAIN_FILE=${TRAIN_FILE:-"${RAY_DATA_HOME}/data/dapo-math-17k.parquet"}
+TEST_FILE=${TEST_FILE:-"${RAY_DATA_HOME}/data/aime-2024.parquet"}
+
+TOTAL_TRAINING_STEPS=${TOTAL_TRAINING_STEPS:-100}
+SAVE_FREQ=${SAVE_FREQ:-100}
+TEST_FREQ=${TEST_FREQ:--1}
+RESUME_MODE=${RESUME_MODE:-disable}
+
+# Algorithm
+temperature=1.0
+top_p=1.0
+top_k=-1 # 0 for HF rollout, -1 for vLLM rollout
+val_top_p=0.7
+
+# Performance Related Parameter
+use_dynamic_bsz=${USE_DYNAMIC_BSZ:-False}
+actor_ppo_max_token_len=$(((max_prompt_length + max_response_length) * 10 / 10))
+infer_ppo_max_token_len=$(((max_prompt_length + max_response_length) * 1))
+param_offload=${PARAM_OFFLOAD:-False}
+optimizer_offload=${OPTIMIZER_OFFLOAD:-True}
+actor_ppo_micro_bsz=${PPO_MICRO_BATCH_SIZE_PER_GPU:-8}
+rollout_gpu_memory_utilization=${ROLLOUT_GPU_MEMORY_UTILIZATION:-0.40}
+rollout_max_num_seqs=${ROLLOUT_MAX_NUM_SEQS:-512}
+cudagraph_capture_sizes=${CUDAGRAPH_CAPTURE_SIZES:-"[1,2,4,8,16,32,64,128,256,512]"}
+gen_tp=${ROLLOUT_TP:-8}
+train_tp=${TRAIN_TP:-8}
+train_pp=1
+EP=1
+ETP=1
+if (( 8 % gen_tp != 0 || 8 % train_tp != 0 )); then
+    echo "ROLLOUT_TP and TRAIN_TP must divide 8; got ${gen_tp} and ${train_tp}" >&2
+    exit 2
+fi
+# ---- end user-adjustable ----
+
+# ---- no user adjustment needed below ----
+################################################### start of config ###################################################
+
+FP8=(
+    # train
+    # +actor_rollout_ref.actor.megatron.override_transformer_config.fp8="e4m3" # e4m3 or hybrid
+    # +actor_rollout_ref.actor.megatron.override_transformer_config.fp8_recipe="blockwise"
+    # +actor_rollout_ref.actor.optim.override_optimizer_config.fp8_recipe="blockwise"
+    # rollout
+    actor_rollout_ref.actor.megatron.dtype=${dtype}
+    actor_rollout_ref.rollout.dtype=${dtype}
+    # +actor_rollout_ref.rollout.quantization="fp8"
+)
+
+DATA=(
+    data.train_files="${TRAIN_FILE}"
+    data.val_files="${TEST_FILE}"
+    data.prompt_key=prompt
+    data.return_raw_chat=$return_raw_chat
+    data.truncation='left'
+    data.max_prompt_length=${max_prompt_length}
+    data.max_response_length=${max_response_length}
+    data.train_batch_size=${train_prompt_bsz}
+)
+
+REWARD_MODEL=(
+    +reward_model.reward_kwargs.overlong_buffer_cfg.enable=${enable_overlong_buffer}
+    +reward_model.reward_kwargs.overlong_buffer_cfg.len=${overlong_buffer_len}
+    +reward_model.reward_kwargs.overlong_buffer_cfg.penalty_factor=${overlong_penalty_factor}
+    +reward_model.reward_kwargs.overlong_buffer_cfg.log=False
+    +reward_model.reward_kwargs.max_resp_len=${max_response_length}
+    reward_model.reward_manager=dapo
+)
+
+PERF_OPT=(
+    +actor_rollout_ref.actor.megatron.override_transformer_config.apply_rope_fusion=True
+    # The fused Mamba training path requires the CUDA-only causal-conv1d extension.
+    # Use the official unfused Mamba path on NPU.
+    +actor_rollout_ref.actor.megatron.override_transformer_config.use_mamba_mem_eff_path=False
+    actor_rollout_ref.model.use_fused_kernels=False
+    +actor_rollout_ref.actor.megatron.override_transformer_config.recompute_method=uniform
+    +actor_rollout_ref.actor.megatron.override_transformer_config.recompute_granularity=full
+    +actor_rollout_ref.actor.megatron.override_transformer_config.recompute_num_layers=1
+)
+
+ACTOR=(
+    actor_rollout_ref.actor.use_kl_loss=${use_kl_loss}
+    actor_rollout_ref.actor.kl_loss_coef=${kl_loss_coef}
+    actor_rollout_ref.actor.clip_ratio_low=${clip_ratio_low}
+    actor_rollout_ref.actor.clip_ratio_high=${clip_ratio_high}
+    actor_rollout_ref.actor.clip_ratio_c=10.0
+    actor_rollout_ref.actor.ppo_micro_batch_size_per_gpu=${actor_ppo_micro_bsz}
+    actor_rollout_ref.actor.use_dynamic_bsz=${use_dynamic_bsz}
+    actor_rollout_ref.actor.ppo_max_token_len_per_gpu=${actor_ppo_max_token_len}
+    actor_rollout_ref.actor.optim.lr=1e-6
+    actor_rollout_ref.actor.optim.lr_warmup_steps=10
+    actor_rollout_ref.actor.optim.weight_decay=0.1
+    actor_rollout_ref.actor.optim.clip_grad=1.0
+    actor_rollout_ref.actor.ppo_mini_batch_size=${train_prompt_mini_bsz}
+    actor_rollout_ref.actor.megatron.param_offload=${param_offload}
+    actor_rollout_ref.actor.megatron.optimizer_offload=${optimizer_offload}
+    actor_rollout_ref.actor.megatron.pipeline_model_parallel_size=${train_pp}
+    actor_rollout_ref.actor.megatron.tensor_model_parallel_size=${train_tp}
+    actor_rollout_ref.actor.megatron.expert_model_parallel_size=$EP
+    actor_rollout_ref.actor.megatron.expert_tensor_parallel_size=$ETP
+    actor_rollout_ref.actor.megatron.context_parallel_size=1
+    actor_rollout_ref.actor.megatron.use_remove_padding=False
+    actor_rollout_ref.actor.entropy_coeff=0
+    actor_rollout_ref.actor.loss_agg_mode=${loss_agg_mode}
+    actor_rollout_ref.actor.megatron.use_mbridge=True
+    actor_rollout_ref.actor.megatron.vanilla_mbridge=False
+    actor_rollout_ref.actor.use_torch_compile=False
+)
+
+ROLLOUT=(
+    actor_rollout_ref.rollout.name=${rollout_name}
+    actor_rollout_ref.rollout.mode=${rollout_mode}
+    actor_rollout_ref.rollout.gpu_memory_utilization=${rollout_gpu_memory_utilization}
+    actor_rollout_ref.rollout.tensor_model_parallel_size=${gen_tp}
+    actor_rollout_ref.rollout.enable_chunked_prefill=True
+    actor_rollout_ref.rollout.enable_prefix_caching=False
+    actor_rollout_ref.rollout.max_num_batched_tokens=$((max_prompt_length + max_response_length))
+    actor_rollout_ref.rollout.max_model_len=$((max_prompt_length + max_response_length))
+    actor_rollout_ref.rollout.max_num_seqs=${rollout_max_num_seqs}
+    actor_rollout_ref.rollout.cudagraph_capture_sizes="${cudagraph_capture_sizes}"
+    +actor_rollout_ref.rollout.engine_kwargs.vllm.compilation_config.cudagraph_mode=FULL_DECODE_ONLY
+    actor_rollout_ref.rollout.temperature=${temperature}
+    actor_rollout_ref.rollout.top_p=${top_p}
+    actor_rollout_ref.rollout.top_k=${top_k}
+    actor_rollout_ref.rollout.val_kwargs.temperature=${temperature}
+    actor_rollout_ref.rollout.val_kwargs.top_p=${val_top_p}
+    actor_rollout_ref.rollout.val_kwargs.top_k=${top_k}
+    actor_rollout_ref.rollout.val_kwargs.do_sample=True
+    actor_rollout_ref.rollout.val_kwargs.n=1
+    actor_rollout_ref.rollout.calculate_log_probs=True
+    actor_rollout_ref.rollout.n=${n_resp_per_prompt}
+)
+
+TRAINER=(
+    trainer.logger=['console']
+    trainer.project_name="${project_name}"
+    trainer.experiment_name="${exp_name}"
+    trainer.n_gpus_per_node=8
+    trainer.nnodes="${NNODES}"
+    trainer.val_before_train=False
+    trainer.test_freq=${TEST_FREQ}
+    trainer.save_freq=${SAVE_FREQ}
+    trainer.total_epochs=10
+    trainer.total_training_steps=${TOTAL_TRAINING_STEPS}
+    trainer.default_local_dir="${CKPTS_DIR}"
+    trainer.resume_mode=${RESUME_MODE}
+    trainer.log_val_generations=10
+    trainer.device=npu
+)
+
+FORWARD_ONLY_SETS=(
+    actor_rollout_ref.ref.log_prob_micro_batch_size_per_gpu=4
+    actor_rollout_ref.rollout.log_prob_micro_batch_size_per_gpu=4
+    actor_rollout_ref.ref.log_prob_use_dynamic_bsz=${use_dynamic_bsz}
+    actor_rollout_ref.rollout.log_prob_use_dynamic_bsz=${use_dynamic_bsz}
+    actor_rollout_ref.ref.log_prob_max_token_len_per_gpu=${infer_ppo_max_token_len}
+    actor_rollout_ref.rollout.log_prob_max_token_len_per_gpu=${infer_ppo_max_token_len}
+    actor_rollout_ref.ref.use_torch_compile=False
+)
+
+MODEL=(
+    actor_rollout_ref.model.path="${MODEL_PATH}"
+    actor_rollout_ref.model.trust_remote_code=True
+    actor_rollout_ref.model.use_remove_padding=False
+)
+
+ALGORITHM=(
+    algorithm.adv_estimator=${adv_estimator}
+    algorithm.use_kl_in_reward=${use_kl_in_reward}
+    algorithm.kl_ctrl.kl_coef=${kl_coef}
+)
+EXTRA=(
+    model_engine=megatron
+)
+
+################################################### start script ###################################################
+
+# uv (set VERL_USE_UV=0 for system python): GPU vllm/sglang × megatron run the driver and every Ray worker
+# (runtime_env.py_executable) through `uv run` on the matching extras of the committed uv.lock;
+# other backends / NPU fall back to ambient python. Run from the verl repo root.
+LAUNCH=(python3)
+RAY=(ray_kwargs.ray_init.runtime_env.py_executable=null)
+if [ "${VERL_USE_UV:-1}" != 0 ] && [ "${DEVICE:-gpu}" = gpu ] && { [ "${rollout_name}" = vllm ] || [ "${rollout_name}" = sglang ]; }; then
+    LAUNCH=(uv run --frozen --all-packages --extra "${rollout_name}" --extra megatron python3)
+    RAY=(ray_kwargs.ray_init.runtime_env.py_executable="uv -v run --frozen --all-packages --extra ${rollout_name} --extra megatron")
+fi
+"${LAUNCH[@]}" -m verl.trainer.main_ppo \
+    "${DATA[@]}" \
+    "${ALGORITHM[@]}" \
+    "${MODEL[@]}" \
+    "${ROLLOUT[@]}" \
+    "${ACTOR[@]}" \
+    "${REWARD_MODEL[@]}" \
+    "${FP8[@]}" \
+    "${PERF_OPT[@]}" \
+    "${TRAINER[@]}" \
+    "${FORWARD_ONLY_SETS[@]}" \
+    "${EXTRA[@]}" \
+    "${RAY[@]}" \
+    "$@"