创建位置

https://gitcode.com/cann/asc-comm/issues/new/requirement
（或仓库 Issues 页 → 新建 → 选择「Requirement|需求建议」模板）

标题

正文

Backgroud（背景信息）

当前 asc-comm 仓中 jetty 层点对点写原语（include/aicore/jetty/hcomm_jetty.h）仅支持 AIV（SIMD）执行模型：WQE 构造、SQ 头推进与 DWQE 发布均以 AIV 主体 + RingDoorbell() 的方式完成。

SIMT VF kernel 尚无法直接驱动 URMA jetty 通道完成点对点写：

• SIMT kernel 需要 jetty 能力时缺少公开、受支持的接口，只能依赖混合执行或自行拼接 WQE；
• 现有接口没有 thread / warp / group 多协作粒度的提交语义，无法利用多 warp 并行构造 WQE、统一发布的吞吐优势；
• 缺少延迟提交（WQE 构造与发布分离）路径，SIMT 侧难以做批量提交的性能优化。

Origin（信息来源）

内部需求：SIMT kernel 直驱 URMA jetty 点对点通信能力（对应 PR #177）。

Benefit / Necessity （价值/作用）

• SIMT VF kernel 获得 jetty 层点对点写能力，无需依赖 AIV 主体或手写 WQE；
• 提供 thread / warp / group 三种协作粒度，支持多 warp 并行构造 WQE、group leader 统一提交的 dispatch 模式；
• 提供立即提交（Put）与延迟提交（Write<..., false> + AdvanceSq + PublishSq）两条路径，便于按场景权衡吞吐与延迟；
• 保留 SIMT/SIMD 共享 SQ 的混合执行场景（SIMT 构造 SQ 内容，SIMD RingDoorbell() 发布），覆盖 SIMT/SIMD 混合 kernel。

Design（设计方案）

在 include/aicore/jetty/ 下新增 SIMT 公共头 hcomm_jetty_simt.h，提供 AscendC::simt::jetty::HcommJetty：

• 复用 SIMD 侧 HcommJettyInfo（80B）/HcommJettyPeerInfo（48B）元数据与 GM jetty 表，kernel 侧直接从 GM jetty 表解析队列与远端信息，不需要静态或动态 UB workspace；
• WQE 布局：1 WQEBB = 64B，warp 协作 WRITE 固定预留 4 WQEBB 槽位，单次 WRITE 最多携带 12 个 SGE；地址或长度为 0 的 SGE 不计入 WRITE，有效 SGE 自动压紧，未使用的 WQEBB 以 NOP 对齐；
• 接口：Write/Put（立即/延迟提交）、WriteValue（立即数写）、WriteWithNotify（带远端通知写）、AdvanceSq/PublishSq（SQ 头推进与 DWQE 发布，缺一不可）、Drain（带超时的回收）；
• 协作组：对齐 cooperative_groups 的参数顺序，支持 warp 内固定槽位预留与 group 统一提交；同一 channel 的并发 Write 必须属于同一 batch，由 group leader 统一提交；
• 实现按平台分层：v310（dav-3510，Ascend 950PR/950DT，COMM_PROTOCOL_UB_CTP）实现位于 src/aicore/jetty/impl/platform_v310/；
• 配套 examples/simt_jetty/ 样例（thread / warp / warp-pad / group / dispatch / mixed 六种模式，含编译/正确性/压测/性能脚本）及 UT 用例。