Features/amd qwen

[siluzhou]

This PR adds ROCm support and validation coverage for Qwen3.5/Qwen3-VL on MI308.

Main changes:

• Support MRoPE position ids on ROCm, including ROCm fused RoPE KV cache path, Qwen3-Next/Qwen3-VL model config propagation, and CUDA Graph replay handling for combo_position_ids.

• Harden ROCm Qwen3.5 runtime integration, including AITER JIT patching, flash-attn ROCm handling, AITER prefill test coverage, and Qwen VL checkpoint prefix compatibility.

[LLLLKKKK]

AI Code Review - PR #1175

Status: BLOCKING

Summary: P0/0 · P1/1 · P2/0 · P3/0

Blocking Issues

P1

• 新增 OSS smoke 用例依赖私有绝对模型路径 @ rtp_llm/test/smoke/data/model/qwen35/qwen35_bf16_rocm.json:3
  • 建议：将 model_path 改为可移植的模型资源别名、环境变量注入或 smoke macro 统一解析；如果这些用例只能在内部资源环境运行，就不要挂到 OSS suite，并加明确 skip/平台条件。同步修正同批新增的 qwen35 ROCm JSON。

Checklist Violations (1 fail / 51 total)

General Principles Checklist

• [6.1] Architecture — 兼容性：外部 HTTP/RPC API、持久数据、配置、环境迁移安全 → issue 新增 OSS smoke 用例依赖私有绝对模型路径
  新增 qwen35 ROCm smoke JSON 在 model_path 固定为绝对私有挂载路径；同批 6 个 qwen35 ROCm JSON 也相同，并在 smoke_rocm_qwen35 OSS suite 中被引用。通用 CI/外部环境没有该挂载会直接无法启动，用例也不应在公开评论中暴露该路径。

Strengths

• ROCm fused RoPE 对 MRoPE position_ids 做了显式校验，避免旧逻辑在缺失 position_ids 时静默算错。
• cache store copy 抽出统一 tensor 工具，并补了 CPU block read 单测，修复方向清晰。
• perf max_seq_len、ROCm 可见设备拆分、flash_attn 缺失 fallback 都有对应测试覆盖。

[LLLLKKKK]

AI Code Review - PR #1175

Status: BLOCKING

Summary: P0/0 · P1/1 · P2/0 · P3/0

Blocking Issues

P1

• ROCm MRoPE 新内核缺少数值回归覆盖 @ rtp_llm/models_py/modules/factory/attention/rocm_impl/test/test_aiter_prefill_op.py:1326
  • 建议：补一个 ROCm 数值回归：使用非一致的 per-axis position_ids（包含首 token 为 0）跑 prefill/decode fused RoPE+KV cache，并与 torch/HF 或 CUDA reference 比对 Q/K/attention 输出；CUDA Graph 路径也建议覆盖。

Checklist Violations (2 fail / 51 total)

General Principles Checklist

• [6.1] Tests — 新逻辑有聚焦单测 + 相关集成/smoke 测试 → issue ROCm MRoPE 新内核缺少数值回归覆盖
  新增的 _check_mrope_accepts_combo_position_ids 只构造 impl 并断言 position_ids.defined()/numel()；没有执行 forward 或与参考 RoPE 输出比对。即使 kernel 的 token-major 索引、axis 映射或 position_id=0 处理错误，这个测试仍会通过。
• [6.1] Tests — 边界 case 覆盖（空、单元素、最大值） → issue ROCm MRoPE 新内核缺少数值回归覆盖
  新增的 _check_mrope_accepts_combo_position_ids 只构造 impl 并断言 position_ids.defined()/numel()；没有执行 forward 或与参考 RoPE 输出比对。即使 kernel 的 token-major 索引、axis 映射或 position_id=0 处理错误，这个测试仍会通过。

Strengths

• cache_store 的 CPU/GPU byte tensor 拷贝被收敛到统一 helper，调用点比原来更一致。
• RDMA mode 传播、perf max_seq_len 保留用户显式配置等行为补了针对性单测。
• ROCm smoke 覆盖扩展到了 Qwen3.5/Qwen3-VL、PD、CUDA Graph 和 reuse cache 等关键组合。

[LLLLKKKK]

AI Code Review - PR #1175

Status: LGTM

Summary: P0/0 · P1/0 · P2/0 · P3/0

lgtm ready to ci

Checklist ✅ (51 items passed)

Strengths

• CACHE_STORE_RDMA_MODE=0 的 ACCL 默认环境变量行为增加了回归测试，覆盖了配置兼容性风险。
• cache store tensor copy 抽成统一工具后补了 CPU blockRead 用例，能覆盖此前容易遗漏的 CPU block 路径。
• ROCm Qwen3.5/MRoPE 相关 smoke、perf 和 attention wrapper 测试覆盖了 CUDA Graph、PD、TP2x2、BF16/FP8 等关键组合。

[LLLLKKKK]

AI Code Review - PR #1175

Status: BLOCKING

Summary: P0/1 · P1/2 · P2/0 · P3/0

Blocking Issues

P0

• MRoPE 分段索引使用线程号，导致后两个位置轴不参与旋转 @ rtp_llm/models_py/bindings/rocm/kernels/fused_rope_kvcache_kernel.cu:124
  • 建议：按实际频率对索引而非线程号选择轴，并处理一个向量跨越 section 边界的情况。增加三轴 position id 明显不同的 prefill/decode 数值测试，与 Torch MRoPE 参考结果及 KV cache 内容逐项比较。

P1

• 新增 MRoPE ROCm 用例会在构造阶段直接失败 @ rtp_llm/models_py/modules/factory/attention/rocm_impl/test/test_aiter_prefill_op.py:126
  • 建议：将支持路径配置为 mrope_interleaved=true、index_factor=3，使用 token-major 且三轴不同的 position ids 并执行 forward 数值比对；另加独立用例通过 factory 验证 false 配置确实回退，而不是直接构造不支持的实现。
• 导入 rtp_llm 会永久替换进程级标准库函数 @ rtp_llm/__init__.py:4
  • 建议：将修复限制在 ROCm AITER JIT 的具体调用点，优先 patch AITER 自身函数。若必须临时拦截标准库，应使用带 try/finally 恢复的上下文管理器、精确校验 AITER 路径，并保持原 API 的返回及 shell 语义。

Checklist Violations (5 fail / 124 total)

General Principles Checklist

• [6.1] Architecture — 分层边界：新概念在正确层级，不泄漏内部 → issue 导入 rtp_llm 会永久替换进程级标准库函数
  任意平台执行 import rtp_llm 都会替换全局 os.system 和 importlib.import_module，且没有恢复路径。文件名包含 /codegen.py 等通用片段的非 AITER 命令也会被改为 shlex.split 后无 shell 执行；非零退出还从返回状态变为抛出 RuntimeError。因此仅导入库就会改变宿主进程中无关代码的执行和错误语义。
• [6.1] Architecture — 错误语义：fail-fast/retry/fallback/silent 行为显式 → issue 导入 rtp_llm 会永久替换进程级标准库函数
  任意平台执行 import rtp_llm 都会替换全局 os.system 和 importlib.import_module，且没有恢复路径。文件名包含 /codegen.py 等通用片段的非 AITER 命令也会被改为 shlex.split 后无 shell 执行；非零退出还从返回状态变为抛出 RuntimeError。因此仅导入库就会改变宿主进程中无关代码的执行和错误语义。
• [6.1] Tests — 分布式/跨平台变更有对应覆盖 → issue 新增 MRoPE ROCm 用例会在构造阶段直接失败
_make_mrope_attn_configs 未设置 mrope_interleaved，其 C++ 默认值为 false；两个 impl_cls 构造器又会无条件创建 fused RoPE op，而新增基类构造器明确拒绝该组合。因此在满足 ROCm/AITER 条件时，测试会在第 1335 行抛异常，无法到达断言。该用例还使用 [4, tokens] 且各轴相同的数据，只检查 numel，未覆盖生产侧 token-major 三轴契约。
• [6.1] Tests — 新逻辑有聚焦单测 + 相关集成/smoke 测试 → issue 新增 MRoPE ROCm 用例会在构造阶段直接失败
_make_mrope_attn_configs 未设置 mrope_interleaved，其 C++ 默认值为 false；两个 impl_cls 构造器又会无条件创建 fused RoPE op，而新增基类构造器明确拒绝该组合。因此在满足 ROCm/AITER 条件时，测试会在第 1335 行抛异常，无法到达断言。该用例还使用 [4, tokens] 且各轴相同的数据，只检查 numel，未覆盖生产侧 token-major 三轴契约。

Python Static-First Checklist

• [P.F] 语言陷阱 — 禁止模块级 import 副作用 → issue 导入 rtp_llm 会永久替换进程级标准库函数
  任意平台执行 import rtp_llm 都会替换全局 os.system 和 importlib.import_module，且没有恢复路径。文件名包含 /codegen.py 等通用片段的非 AITER 命令也会被改为 shlex.split 后无 shell 执行；非零退出还从返回状态变为抛出 RuntimeError。因此仅导入库就会改变宿主进程中无关代码的执行和错误语义。

Strengths

• mrope_interleaved 已贯通模型配置、pybind 和 ROCm 实现，并为不支持的布局提供了显式回退判断。
• CUDA Graph 对 combo_position_ids 增加了容量检查，并按元素大小计算复制字节数。
• cache-store、配置传播及 ROCm smoke 场景获得了较丰富的测试补充。

[LLLLKKKK]

AI Code Review - PR #1175

Status: BLOCKING

Summary: P0/0 · P1/3 · P2/1 · P3/0

Blocking Issues

P1

• MRoPE kernel 按线程而非 rotary 维度选择 position axis @ rtp_llm/models_py/bindings/rocm/kernels/fused_rope_kvcache_kernel.cu:140
  • 建议：按实际 scalar/rotary-pair 索引选择 axis；若一个 vector 跨越 section，需要逐 lane 处理或拆分 vector。增加包含不同 temporal/height/width position ids 的 prefill、decode 数值 reference 测试，并覆盖 section 非 vector 对齐场景。
• 新增 MRoPE 测试会被自身的 interleaved 校验拒绝 @ rtp_llm/models_py/modules/factory/attention/rocm_impl/test/test_aiter_prefill_op.py:126
  • 建议：将测试配置设为生产支持的 mrope_interleaved=True，并让 index_factor、position tensor 的 token-major 布局与真实模型一致。测试不应止于 prepare()，还应执行 forward 并与独立 MRoPE reference 比较。
• 导入顶层包会修改进程级标准库行为 @ rtp_llm/__init__.py:4
  • 建议：把修复放到 ROCm Aiter 初始化路径，并只包装具体 Aiter codegen/import 调用；优先修改 Aiter adapter，或使用有明确进入/退出边界的 context manager，避免在包 __init__ 中 monkey patch 全局标准库。

Non-blocking Suggestions

P2

• 通用 HIP 调试变量会开启逐层同步 tensor dump @ rtp_llm/models_py/triton_kernels/fla/utils.py:29
  • 建议：只允许专用诊断开关启用这些 dump；加入次数/请求范围限制，默认仅记录 shape、dtype 和状态，不回读 tensor 值或请求标识。需要完整 dump 时使用显式、短生命周期的诊断工具，而非常驻 forward 日志。

Checklist Violations (8 fail / 124 total)

General Principles Checklist

• [6.1] Architecture — 依赖方向：无循环依赖/跨层惊喜 → issue 导入顶层包会修改进程级标准库行为
  任意 import rtp_llm 都会执行 install_aiter_jit_patch()，进而替换进程全局的 os.system 和 importlib.import_module；特定 import 失败时还会删除 .so 与 build 目录。该副作用不受 ROCm、Aiter 是否启用或具体模型限制，会改变宿主应用、插件及测试进程的标准库语义。
• [6.1] Architecture — 分层边界：新概念在正确层级，不泄漏内部 → issue 导入顶层包会修改进程级标准库行为
  任意 import rtp_llm 都会执行 install_aiter_jit_patch()，进而替换进程全局的 os.system 和 importlib.import_module；特定 import 失败时还会删除 .so 与 build 目录。该副作用不受 ROCm、Aiter 是否启用或具体模型限制，会改变宿主应用、插件及测试进程的标准库语义。
• [6.1] Architecture — 可观测性：日志/指标/超时可操作、非噪声 → issue 通用 HIP 调试变量会开启逐层同步 tensor dump
HIP_LAUNCH_BLOCKING 被等同于专用诊断开关。启用后，model、每层 GDN、sampler 和 C++ wrapper 会逐 forward 输出 warning，并通过 .cpu()、.item()、min/max/isfinite 回读 GPU tensor；日志还包含 token、request 和 tensor 样本。仅为定位异步错误开启 launch blocking 时，会额外产生大量同步、日志和请求数据暴露。
• [6.1] Quality — 无 per-forward 调试日志 / 噪声热路径输出 → issue 通用 HIP 调试变量会开启逐层同步 tensor dump
HIP_LAUNCH_BLOCKING 被等同于专用诊断开关。启用后，model、每层 GDN、sampler 和 C++ wrapper 会逐 forward 输出 warning，并通过 .cpu()、.item()、min/max/isfinite 回读 GPU tensor；日志还包含 token、request 和 tensor 样本。仅为定位异步错误开启 launch blocking 时，会额外产生大量同步、日志和请求数据暴露。
• [6.1] Tests — 新逻辑有聚焦单测 + 相关集成/smoke 测试 → issue 新增 MRoPE 测试会被自身的 interleaved 校验拒绝
_make_mrope_attn_configs 设置 RopeStyle.Mrope，但没有把新增字段 mrope_interleaved 设为 True，其 C++ 默认值为 false。测试随后直接构造 AiterPrefillImpl*，而构造函数会对该组合抛出异常，因此 ROCm target 中两个新增用例都无法到达 position_ids 断言。

RTP-LLM Checklist

• [F] 跨平台 — ROCm/XPU/CUDA fallback 的 shape、stride、dtype、sentinel/mask、sampling/beam/logits 数值语义必须与 CUDA/CPU reference 对齐；不支持时必须 graceful reject → issue MRoPE kernel 按线程而非 rotary 维度选择 position axis
get_mrope_position_id 将 threadIdx.x 直接交给 section 映射，但每个线程处理一个 8 元素 Vec_t。mrope_dim* 来自 rotary 维度分段，例如首段为 16 时，处理后续维度的多个线程仍落入 axis 0；跨 section 的 vector 也只能获得一个 position id。新测试只检查 tensor 被保存，没有执行 kernel 或比较数值，图像/空间轴会使用错误位置。
• [H] 测试与 CI — 新增、迁移或删除测试必须证明目标行为仍在 CI 中执行；DISABLED_、#if 0、open_skip、导入即失败、未被 target 引用的用例不得计为覆盖 → issue 新增 MRoPE 测试会被自身的 interleaved 校验拒绝
_make_mrope_attn_configs 设置 RopeStyle.Mrope，但没有把新增字段 mrope_interleaved 设为 True，其 C++ 默认值为 false。测试随后直接构造 AiterPrefillImpl*，而构造函数会对该组合抛出异常，因此 ROCm target 中两个新增用例都无法到达 position_ids 断言。

Python Static-First Checklist

• [P.F] 语言陷阱 — 禁止模块级 import 副作用 → issue 导入顶层包会修改进程级标准库行为
  任意 import rtp_llm 都会执行 install_aiter_jit_patch()，进而替换进程全局的 os.system 和 importlib.import_module；特定 import 失败时还会删除 .so 与 build 目录。该副作用不受 ROCm、Aiter 是否启用或具体模型限制，会改变宿主应用、插件及测试进程的标准库语义。

Strengths

• Cache Store 的 CPU/GPU byte copy 已收敛到统一 helper，并补充了 CPU block read 覆盖。
• FLA 与 scatter kernel 的大 offset 改为 64 位计算，并增加了溢出边界测试。
• combo_position_ids 已贯通普通 forward 与 CUDA Graph capture/replay，ROCm smoke 也覆盖了 TP、PD、reuse 和 graph 组合。