增加jetty层simt接口-asc-comm-AtomGit

Mexamples/README.md+1-0

	@@ -14,6 +14,7 @@
14	\| [aicore/hcomm/03_one_multi_path](./aicore/hcomm/03_one_multi_path/README.md) \| 查询`UB_MEM`链路，为每个peer创建one path/multi path Channel，逐个处理peer，并通过双Stream并发搬运和校验该peer的远端数据。 \| Ascend 950PR/Ascend 950DT \|	14	\| [aicore/hcomm/03_one_multi_path](./aicore/hcomm/03_one_multi_path/README.md) \| 查询`UB_MEM`链路，为每个peer创建one path/multi path Channel，逐个处理peer，并通过双Stream并发搬运和校验该peer的远端数据。 \| Ascend 950PR/Ascend 950DT \|
15	\| [aicore/hcomm/04_simt_urma](./aicore/hcomm/04_simt_urma/README.md) \| 演示并验证全部 SIMT URMA 接口：`WriteNbi`、`WriteValueNbi`、`WriteWithNotifyNbi`、`AtomicFAA` 和 `AtomicCAS`，覆盖立即提交与批量提交。 \| Ascend 950PR / Ascend 950DT \|	15	\| [aicore/hcomm/04_simt_urma](./aicore/hcomm/04_simt_urma/README.md) \| 演示并验证全部 SIMT URMA 接口：`WriteNbi`、`WriteValueNbi`、`WriteWithNotifyNbi`、`AtomicFAA` 和 `AtomicCAS`，覆盖立即提交与批量提交。 \| Ascend 950PR / Ascend 950DT \|
16	\| [aicore/hcomm/05_simt_urma_perftest](./aicore/hcomm/05_simt_urma_perftest/README.md) \| 测量上述五个 SIMT URMA 接口的下发时延与完成带宽，覆盖立即提交/延迟提交两种发布策略。 \| Ascend 950PR / Ascend 950DT \|	16	\| [aicore/hcomm/05_simt_urma_perftest](./aicore/hcomm/05_simt_urma_perftest/README.md) \| 测量上述五个 SIMT URMA 接口的下发时延与完成带宽，覆盖立即提交/延迟提交两种发布策略。 \| Ascend 950PR / Ascend 950DT \|
		17	+\| [simt_jetty](./simt_jetty/README.md) \| 演示多卡场景下SIMT Kernel以thread/warp/group协作方式通过`HcommJetty::Write`直接构造并提交Jetty WQE，覆盖dispatch多warp批量提交与SIMT/SIMD混合发布模式，并附设备侧性能与压力测试脚本。 \| Ascend 950PR / Ascend 950DT \|
17	\| [hcomm_jetty_write](./hcomm_jetty_write/README.md) \| 演示多卡场景下AIV Kernel通过`HcommJetty::Write`和`HcommJetty::WriteValue`直接向Jetty SQ提交WQE，并通过`Drain`等待完成与校验结果。 \| Ascend 950PR/Ascend 950DT \|	18	\| [hcomm_jetty_write](./hcomm_jetty_write/README.md) \| 演示多卡场景下AIV Kernel通过`HcommJetty::Write`和`HcommJetty::WriteValue`直接向Jetty SQ提交WQE，并通过`Drain`等待完成与校验结果。 \| Ascend 950PR/Ascend 950DT \|
18	\| [ccu/ccu_direct](./ccu/ccu_direct/01_allgather/README.md) \| 演示基于HCCL通信域和CCU数据面接口，以直调`<<<>>>`方式实现AllGather等集合通信操作。 \| Ascend 950PR/Ascend 950DT \|	19	\| [ccu/ccu_direct](./ccu/ccu_direct/01_allgather/README.md) \| 演示基于HCCL通信域和CCU数据面接口，以直调`<<<>>>`方式实现AllGather等集合通信操作。 \| Ascend 950PR/Ascend 950DT \|
19		20

Mexamples/README_en.md+1-0

	@@ -14,6 +14,7 @@ This directory contains usage samples for asc-comm APIs.
14	\| [aicore/hcomm/03_one_multi_path](./aicore/hcomm/03_one_multi_path/README_en.md) \| Queries a `UB_MEM` link, creates one path/multi path channels for every peer, processes peers serially, and concurrently moves and verifies each peer's remote data over two streams. \| Ascend 950PR / Ascend 950DT \|	14	\| [aicore/hcomm/03_one_multi_path](./aicore/hcomm/03_one_multi_path/README_en.md) \| Queries a `UB_MEM` link, creates one path/multi path channels for every peer, processes peers serially, and concurrently moves and verifies each peer's remote data over two streams. \| Ascend 950PR / Ascend 950DT \|
15	\| [aicore/hcomm/04_simt_urma](./aicore/hcomm/04_simt_urma/README_en.md) \| Demonstrates and validates every SIMT URMA interface: `WriteNbi`, `WriteValueNbi`, `WriteWithNotifyNbi`, `AtomicFAA`, and `AtomicCAS`, with immediate and batched submission. \| Ascend 950PR / Ascend 950DT \|	15	\| [aicore/hcomm/04_simt_urma](./aicore/hcomm/04_simt_urma/README_en.md) \| Demonstrates and validates every SIMT URMA interface: `WriteNbi`, `WriteValueNbi`, `WriteWithNotifyNbi`, `AtomicFAA`, and `AtomicCAS`, with immediate and batched submission. \| Ascend 950PR / Ascend 950DT \|
16	\| [aicore/hcomm/05_simt_urma_perftest](./aicore/hcomm/05_simt_urma_perftest/README_en.md) \| Measures issue latency and completion bandwidth for those five SIMT URMA interfaces under immediate and deferred publication. \| Ascend 950PR / Ascend 950DT \|	16	\| [aicore/hcomm/05_simt_urma_perftest](./aicore/hcomm/05_simt_urma_perftest/README_en.md) \| Measures issue latency and completion bandwidth for those five SIMT URMA interfaces under immediate and deferred publication. \| Ascend 950PR / Ascend 950DT \|
		17	+\| [simt_jetty](./simt_jetty/README_en.md) \| Demonstrates a SIMT Kernel building and submitting Jetty WQEs directly through `HcommJetty::Write` with thread/warp/group cooperation, covering the multi-warp dispatch batching mode and the mixed SIMT/SIMD publication mode, with device-side performance and stress scripts. \| Ascend 950PR / Ascend 950DT \|
17	\| [hcomm_jetty_write](./hcomm_jetty_write/README_en.md) \| Demonstrates an AIV Kernel submitting WQEs directly to the Jetty SQ through `HcommJetty::Write` and `HcommJetty::WriteValue` in a multi-card scenario, with `Drain` used for completion and verification. \| Ascend 950PR / Ascend 950DT \|	18	\| [hcomm_jetty_write](./hcomm_jetty_write/README_en.md) \| Demonstrates an AIV Kernel submitting WQEs directly to the Jetty SQ through `HcommJetty::Write` and `HcommJetty::WriteValue` in a multi-card scenario, with `Drain` used for completion and verification. \| Ascend 950PR / Ascend 950DT \|
18	\| [ccu/ccu_direct](./ccu/ccu_direct/01_allgather/README_en.md) \| Demonstrates AllGather and other collective communication operations using HCCL communicators and CCU dataplane interfaces with direct `<<<>>>` invocation. \| Ascend 950PR/Ascend 950DT \|	19	\| [ccu/ccu_direct](./ccu/ccu_direct/01_allgather/README_en.md) \| Demonstrates AllGather and other collective communication operations using HCCL communicators and CCU dataplane interfaces with direct `<<<>>>` invocation. \| Ascend 950PR/Ascend 950DT \|
19		20

Aexamples/simt_jetty/CMakeLists.txt+126-0

@@ -0,0 +1,126 @@
	1	+# ----------------------------------------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------------------------------------
	10	+
	11	+cmake_minimum_required(VERSION 3.16)
	12	+
	13	+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
	14	+if(NOT CMAKE_ASC_ARCHITECTURES STREQUAL "dav-3510")
	15	+ message(FATAL_ERROR "simt_jetty only supports dav-3510, but got ${CMAKE_ASC_ARCHITECTURES}")
	16	+endif()
	17	+
	18	+find_package(ASC REQUIRED)
	19	+
	20	+project(simt_jetty LANGUAGES ASC CXX)
	21	+
	22	+option(HCOMM_CCE_RES_USAGE "Print CCE stack and register usage" OFF)
	23	+set(SIMT_JETTY_SGE_NUM "12" CACHE STRING "Compile-time SGE count for thread and warp modes (1 to 12)")
	24	+if(NOT SIMT_JETTY_SGE_NUM MATCHES "^[1-9]$\|^1[0-2]$")
	25	+ message(FATAL_ERROR "SIMT_JETTY_SGE_NUM must be an integer in [1, 12]")
	26	+endif()
	27	+
	28	+set(CMAKE_CXX_STANDARD 17)
	29	+set(CMAKE_CXX_STANDARD_REQUIRED ON)
	30	+set(CMAKE_POSITION_INDEPENDENT_CODE ON)
	31	+
	32	+set_source_files_properties(op_kernel.cpp PROPERTIES LANGUAGE ASC)
	33	+
	34	+# The public headers reach their implementation through package-layout relative paths
	35	+# (../../../impl/comm_api/...), which only resolve in an installed tree. Stage a
	36	+# package-shaped header tree from the working copy so this sample builds against the
	37	+# repository instead of whatever asc-comm version happens to be installed.
	38	+set(ASCCOMM_ROOT "${CMAKE_CURRENT_SOURCE_DIR}/../..")
	39	+set(ASCCOMM_STAGE_ROOT "${CMAKE_BINARY_DIR}/package_layout")
	40	+# Mirror the installed depth exactly. Installed, the public headers live at
	41	+# asc/include/adv_api/<mod>/ and are also reachable via the asc/include/comm_api/aicore/<mod>
	42	+# symlink; their "../../../impl/comm_api/..." includes resolve relative to the real adv_api
	43	+# location, i.e. to asc/impl/. Staging the headers directly under comm_api/aicore would make
	44	+# that same path resolve to asc/include/impl/ and fail, so replicate the adv_api layout.
	45	+set(ASCCOMM_STAGE_INCLUDE_DIR "${ASCCOMM_STAGE_ROOT}/asc/include/adv_api")
	46	+set(ASCCOMM_STAGE_IMPL_DIR "${ASCCOMM_STAGE_ROOT}/asc/impl/comm_api/aicore")
	47	+set(ASCCOMM_STAGE_STAMP "${ASCCOMM_STAGE_ROOT}/.header_layout.stamp")
	48	+
	49	+file(GLOB_RECURSE ASCCOMM_STAGE_INPUTS CONFIGURE_DEPENDS
	50	+ "${ASCCOMM_ROOT}/include/aicore/*"
	51	+ "${ASCCOMM_ROOT}/src/aicore/*"
	52	+)
	53	+
	54	+add_custom_command(
	55	+ OUTPUT ${ASCCOMM_STAGE_STAMP}
	56	+ COMMAND ${CMAKE_COMMAND} -E remove_directory "${ASCCOMM_STAGE_ROOT}/asc"
	57	+ COMMAND ${CMAKE_COMMAND} -E copy_directory
	58	+ "${ASCCOMM_ROOT}/include/aicore/hcomm" "${ASCCOMM_STAGE_INCLUDE_DIR}/hcomm"
	59	+ COMMAND ${CMAKE_COMMAND} -E copy_directory
	60	+ "${ASCCOMM_ROOT}/include/aicore/jetty" "${ASCCOMM_STAGE_INCLUDE_DIR}/jetty"
	61	+ COMMAND ${CMAKE_COMMAND} -E copy_directory
	62	+ "${ASCCOMM_ROOT}/src/aicore/hcomm" "${ASCCOMM_STAGE_IMPL_DIR}/hcomm"
	63	+ COMMAND ${CMAKE_COMMAND} -E copy_directory
	64	+ "${ASCCOMM_ROOT}/src/aicore/jetty" "${ASCCOMM_STAGE_IMPL_DIR}/jetty"
	65	+ COMMAND ${CMAKE_COMMAND} -E copy_directory
	66	+ "${ASCCOMM_ROOT}/src/aicore/utils" "${ASCCOMM_STAGE_IMPL_DIR}/utils"
	67	+ # The ASC driver prepends its own search paths (including the installed asc/include/adv_api)
	68	+ # ahead of every -I, so an "#include \"hcomm/...\"" always finds the installed copy first and
	69	+ # no -I ordering can override it. Expose the staged headers under a prefix that exists
	70	+ # nowhere in the driver's builtin paths. It is a symlink to the staged adv_api dir, so the
	71	+ # public headers' "../../../impl/..." includes still resolve from the real location.
	72	+ COMMAND ${CMAKE_COMMAND} -E create_symlink
	73	+ "adv_api" "${ASCCOMM_STAGE_ROOT}/asc/include/asccomm_staged"
	74	+ COMMAND ${CMAKE_COMMAND} -E touch ${ASCCOMM_STAGE_STAMP}
	75	+ DEPENDS ${ASCCOMM_STAGE_INPUTS}
	76	+ COMMENT "Preparing package-shaped asc-comm headers for simt_jetty"
	77	+)
	78	+add_custom_target(simt_jetty_header_layout DEPENDS ${ASCCOMM_STAGE_STAMP})
	79	+
	80	+add_executable(simt_jetty
	81	+ main.cpp
	82	+ op_kernel.cpp
	83	+)
	84	+
	85	+add_dependencies(simt_jetty simt_jetty_header_layout)
	86	+
	87	+target_compile_definitions(simt_jetty PRIVATE
	88	+ USE_SIMT
	89	+ SIMT_JETTY_SGE_NUM=${SIMT_JETTY_SGE_NUM}
	90	+)
	91	+
	92	+set_target_properties(simt_jetty PROPERTIES LINKER_LANGUAGE ASC)
	93	+
	94	+target_include_directories(simt_jetty BEFORE PRIVATE
	95	+ ${CMAKE_CURRENT_SOURCE_DIR}
	96	+ ${ASCCOMM_STAGE_ROOT}/asc/include
	97	+ $ENV{ASCEND_HOME_PATH}/include
	98	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc
	99	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc/include
	100	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc/include/basic_api
	101	+ # Deliberately NOT adding ${ASCEND_HOME_PATH}/asc/include/adv_api here: it carries an
	102	+ # installed copy of the Hcomm/Jetty headers that shadows the staged tree above regardless
	103	+ # of -I order. kernel_operator.h and friends resolve from asc/include instead.
	104	+)
	105	+
	106	+target_link_directories(simt_jetty PRIVATE
	107	+ $ENV{ASCEND_HOME_PATH}/lib64
	108	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/lib64
	109	+)
	110	+
	111	+target_link_libraries(simt_jetty PRIVATE
	112	+ hcomm
	113	+ hccl
	114	+ ascendcl
	115	+ acl_rt
	116	+ runtime
	117	+ c_sec
	118	+ dl
	119	+ pthread
	120	+)
	121	+
	122	+target_compile_options(simt_jetty PRIVATE
	123	+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
	124	+ $<$<COMPILE_LANGUAGE:ASC>:--cce-auto-infer-kernel-type=false>
	125	+ $<$<AND:$<COMPILE_LANGUAGE:ASC>,$<BOOL:${HCOMM_CCE_RES_USAGE}>>:--cce-res-usage>
	126	+)

Aexamples/simt_jetty/README.md+136-0

@@ -0,0 +1,136 @@
	1	+# SIMT Jetty WRITE样例
	2	+
	3	+## 概述
	4	+
	5	+本样例展示如何在AscendC SIMT Kernel中通过`jetty/hcomm_jetty_simt.h`接口使用URMA Jetty点对点通信能力。样例覆盖`thread / warp / group`三种协作方式，以及多个warp并行生成WQE、group统一提交的dispatch模式和SIMT/SIMD混合发布的mixed模式。
	6	+
	7	+Kernel侧通过`HcommPeer`和`HcommJetty`接口直接从GM Jetty表解析队列与远端信息并提交数据面通信任务，不需要静态或动态UB workspace。固定SGE模板参数范围为2到12；地址或长度为0的SGE不计入WRITE，有效SGE自动压紧，未使用的WQEBB使用NOP对齐。
	8	+
	9	+延迟提交路径由`Write<..., false>`、`AdvanceSq`和`PublishSq`组成：`AdvanceSq`完成SQ头推进与终止WQE标记，`PublishSq`写入128B DWQE窗口发布队列，二者缺一不可。立即提交路径`Put`在单次调用内完成WQE构造、SQ写入与DWQE发布。同一个channel的并发`Write`必须属于同一个batch，并在所有参与warp完成后由group leader统一提交；所有lane构造`HcommJetty`后、任一warp开始`Write`前必须先做一次group同步。
	10	+
	11	+`mixed`模式使用`__global__ __mix__`入口：SIMT VF构造并提交SQ内容，AIV主体随后使用共享的80B Jetty元数据调用SIMD `RingDoorbell()`发布队列，保留SIMT/SIMD共享SQ的场景。mixed是correctness场景，不参与`perf.sh`的设备计时模式。
	12	+
	13	+## 支持范围
	14	+
	15	+\| 项目 \| 说明 \|
	16	+\| --- \| --- \|
	17	+\| 产品 \| `Ascend 950PR / Ascend 950DT` \|
	18	+\| NPU架构 \| `dav-3510` \|
	19	+\| 通信协议 \| `COMM_PROTOCOL_UB_CTP` \|
	20	+\| 执行模型 \| SIMT VF（mixed模式额外包含AIV主体） \|
	21	+\| 部署形态 \| 单机多卡 \|
	22	+\| CANN包 \| 使用当前环境中`set_env.sh`指向的CANN包 \|
	23	+
	24	+> 单卡环境仅支持编译验证，实际运行需要至少2张NPU。
	25	+
	26	+## 目录结构
	27	+
	28	+```text
	29	+simt_jetty
	30	+├── CMakeLists.txt // 编译工程文件
	31	+├── README.md // 中文说明
	32	+├── README_en.md // 英文说明
	33	+├── build.sh // 编译脚本
	34	+├── run.sh // 多rank正确性执行脚本
	35	+├── stress.sh // 同channel重复压测脚本
	36	+├── perf.sh // 设备侧计时性能脚本
	37	+├── main.cpp // Host侧资源准备与结果校验
	38	+├── op_kernel.cpp // SIMT Kernel侧Jetty调用
	39	+└── simt_jetty_common.h // 样例共享定义
	40	+```
	41	+
	42	+## 模式说明
	43	+
	44	+\| 模式 \| 说明 \|
	45	+\| --- \| --- \|
	46	+\| `thread` \| 单线程构造12个有效SGE的4BB WQE，按2-SGE分段独立发布，验证标量路径 \|
	47	+\| `warp` \| warp协作构造WQE，验证固定槽位预留的立即与延迟提交 \|
	48	+\| `warp-pad` \| 1个有效SGE加零长度SGE，验证`1BB WRITE + 1BB NOP`填充路径 \|
	49	+\| `group` \| 多warp共享batch状态，group leader统一提交 \|
	50	+\| `dispatch` \| 每个warp各预留一个WQE，block统一提交实际数量 \|
	51	+\| `mixed` \| SIMT VF构造SQ内容，AIV侧SIMD `RingDoorbell()`发布 \|
	52	+
	53	+## 编译
	54	+
	55	+在本样例目录下执行：
	56	+
	57	+```bash
	58	+# cann包为默认路径安装时，以root用户为例（非root用户，将/usr/local替换为${HOME}）
	59	+source /usr/local/Ascend/cann/set_env.sh
	60	+# cann包为指定路径安装时
	61	+# source ${install_path}/cann/set_env.sh
	62	+
	63	+bash build.sh
	64	+```
	65	+
	66	+没有硬件时，可在编译阶段检查SIMT VF的栈和寄存器占用：
	67	+
	68	+```bash
	69	+CCE_RES_USAGE=1 bash build.sh
	70	+```
	71	+
	72	+默认SGE数量为12，可通过`SIMT_JETTY_SGE_NUM`重编译调整（范围1到12）：
	73	+
	74	+```bash
	75	+cmake -DSIMT_JETTY_SGE_NUM=6 .. && make -j
	76	+```
	77	+
	78	+## 运行
	79	+
	80	+编译完成后，在样例目录执行：
	81	+
	82	+```bash
	83	+./run.sh <nranks> [all\|thread\|warp\|warp-pad\|group\|dispatch\|mixed] [options]
	84	+```
	85	+
	86	+例如两卡执行全部模式：
	87	+
	88	+```bash
	89	+./run.sh 2
	90	+```
	91	+
	92	+也可以只执行一个模式，避免失败用例污染后续SQ/CQ状态：
	93	+
	94	+```bash
	95	+./run.sh 2 warp
	96	+./run.sh 2 mixed
	97	+```
	98	+
	99	+通用选项：`--iterations N`指定正确性校验轮数（默认1），`--warmup N`与`--timing`供性能模式使用。
	100	+
	101	+## 正确性压测
	102	+
	103	+`stress.sh`在同一个channel/SQ上重复执行，每一轮均等待接收端校验完成。接收端校验后会清零目标缓冲区，避免后一轮未实际下发时被前一轮结果掩盖。默认每个模式执行100次：
	104	+
	105	+```bash
	106	+./stress.sh 2
	107	+./stress.sh 2 warp-pad 1000
	108	+```
	109	+
	110	+## 性能测试
	111	+
	112	+`perf.sh`使用设备侧`clock()`打点。性能循环单独调用`Write`组WQE，在`Write`返回后记录`write_cycles`，再调用提交并单独记录`commit_cycles`，最后在`Drain`返回后记录完整完成耗时`completion_cycles`，因此`average_write_us`不包含提交时间。默认warmup 100次、计时1000次，每个warmup/计时样本单独launch一个kernel；每轮都会`Drain`，确保SQ tail持续推进。host launch和同步时间不计入上述设备周期：
	113	+
	114	+```bash
	115	+./perf.sh 2
	116	+./perf.sh 2 warp 10000 1000
	117	+```
	118	+
	119	+结果中的周期按当前设备1 GHz时钟换算为微秒，与`simt_urma_perftest`的统计方式一致。
	120	+
	121	+## 执行结果
	122	+
	123	+正确性模式下，rank 1校验接收缓冲区各槽位，例如：
	124	+
	125	+```text
	126	+[rank 1] simt_jetty mode 4 iteration 0 PASS
	127	+```
	128	+
	129	+所有请求的模式均输出`PASS`表示样例执行成功。
	130	+
	131	+## 注意事项
	132	+
	133	+- 编译和运行前必须先加载CANN环境变量，确保`ASCEND_CANN_PACKAGE_PATH`、ASC CMake模块和运行时动态库可用。
	134	+- 样例依赖CANN中的`hccl`、`hcomm`、`ascendcl`和`runtime`库。
	135	+- `nranks`必须大于等于2；样例只实际使用rank 0和rank 1，其余rank会直接退出。
	136	+- 样例仅支持单机多卡运行，不支持跨节点多机运行。

Aexamples/simt_jetty/README_en.md+137-0

@@ -0,0 +1,137 @@
	1	+# SIMT Jetty WRITE Sample
	2	+
	3	+## Overview
	4	+
	5	+This sample demonstrates how to use URMA Jetty point-to-point communication capabilities from an AscendC SIMT Kernel through the `jetty/hcomm_jetty_simt.h` interface. It covers the `thread / warp / group` cooperation modes, the dispatch mode where multiple warps build WQEs in parallel and the group commits them together, and the mixed mode where SIMT and SIMD share the publication.
	6	+
	7	+The Kernel side resolves queue and remote metadata directly from the GM Jetty table through the `HcommPeer` and `HcommJetty` interfaces and submits data-plane communication tasks, with no static or dynamic UB workspace. The fixed SGE template parameter ranges from 2 to 12; SGEs with a zero address or length do not count toward the WRITE, valid SGEs are compacted, and unused WQEBBs are padded with NOPs.
	8	+
	9	+The deferred path consists of `Write<..., false>`, `AdvanceSq`, and `PublishSq`: `AdvanceSq` advances the SQ head and marks the final WQE, while `PublishSq` writes the 128B DWQE window to publish the queue — both steps are required. The immediate path `Put` builds the WQE, writes the SQ, and publishes the DWQE within a single call. Concurrent `Write` calls on the same channel must belong to the same batch and are committed by the group leader after every participating warp finishes; a group synchronization must happen after all lanes construct `HcommJetty` and before any warp starts to `Write`.
	10	+
	11	+The `mixed` mode uses a `__global__ __mix__` entry: the SIMT VF builds and submits the SQ contents, and the AIV body then calls the SIMD `RingDoorbell()` through the shared 80B Jetty metadata to publish the queue, preserving the SIMT/SIMD shared-SQ scenario. Mixed is a correctness scenario and does not participate in the device-timed modes of `perf.sh`.
	12	+
	13	+## Supported Scope
	14	+
	15	+\| Item \| Description \|
	16	+\| --- \| --- \|
	17	+\| Product \| `Ascend 950PR / Ascend 950DT` \|
	18	+\| NPU architecture \| `dav-3510` \|
	19	+\| Communication protocol \| `COMM_PROTOCOL_UB_CTP` \|
	20	+\| Execution model \| SIMT VF (the mixed mode additionally contains an AIV body) \|
	21	+\| Deployment mode \| Single-node multi-card \|
	22	+\| CANN package \| The CANN package selected by the current `set_env.sh` environment \|
	23	+
	24	+> A single-card environment only supports build verification. Running this sample requires at least two NPUs.
	25	+
	26	+## Directory Structure
	27	+
	28	+```text
	29	+simt_jetty
	30	+├── CMakeLists.txt // Build project file
	31	+├── README.md // Chinese README
	32	+├── README_en.md // English README
	33	+├── build.sh // Build script
	34	+├── run.sh // Multi-rank correctness launch script
	35	+├── stress.sh // Repeated same-channel stress script
	36	+├── perf.sh // Device-side timing performance script
	37	+├── main.cpp // Host-side setup and result verification
	38	+├── op_kernel.cpp // Kernel-side SIMT Jetty calls
	39	+└── simt_jetty_common.h // Shared sample definitions
	40	+```
	41	+
	42	+## Mode Description
	43	+
	44	+\| Mode \| Description \|
	45	+\| --- \| --- \|
	46	+\| `thread` \| A single thread builds a 4BB WQE with 12 valid SGEs, published per 2-SGE segment; verifies the scalar path \|
	47	+\| `warp` \| Warp-cooperative WQE construction; verifies fixed slot reservation with immediate and deferred submission \|
	48	+\| `warp-pad` \| One valid SGE plus zero-length SGEs; verifies the `1BB WRITE + 1BB NOP` padding path \|
	49	+\| `group` \| Multiple warps share the batch state; the group leader commits \|
	50	+\| `dispatch` \| Each warp reserves one WQE; the block commits the actual count \|
	51	+\| `mixed` \| The SIMT VF builds the SQ contents; the AIV side publishes through the SIMD `RingDoorbell()` \|
	52	+
	53	+## Build
	54	+
	55	+Run the following commands in this sample directory:
	56	+
	57	+```bash
	58	+# If CANN is installed in the default path, for example as root.
	59	+# For a non-root user, replace /usr/local with ${HOME}.
	60	+source /usr/local/Ascend/cann/set_env.sh
	61	+# If CANN is installed in a custom path
	62	+# source ${install_path}/cann/set_env.sh
	63	+
	64	+bash build.sh
	65	+```
	66	+
	67	+Without hardware, the stack and register usage of the SIMT VF can still be checked at build time:
	68	+
	69	+```bash
	70	+CCE_RES_USAGE=1 bash build.sh
	71	+```
	72	+
	73	+The default SGE count is 12 and can be changed by rebuilding with `SIMT_JETTY_SGE_NUM` (range 1 to 12):
	74	+
	75	+```bash
	76	+cmake -DSIMT_JETTY_SGE_NUM=6 .. && make -j
	77	+```
	78	+
	79	+## Run
	80	+
	81	+After the build completes, run the following commands in this sample directory:
	82	+
	83	+```bash
	84	+./run.sh <nranks> [all\|thread\|warp\|warp-pad\|group\|dispatch\|mixed] [options]
	85	+```
	86	+
	87	+For example, to run every mode on two cards:
	88	+
	89	+```bash
	90	+./run.sh 2
	91	+```
	92	+
	93	+A single mode can also be run to avoid a failing case polluting the SQ/CQ state of subsequent cases:
	94	+
	95	+```bash
	96	+./run.sh 2 warp
	97	+./run.sh 2 mixed
	98	+```
	99	+
	100	+Common options: `--iterations N` sets the correctness rounds (default 1); `--warmup N` and `--timing` are for the performance mode.
	101	+
	102	+## Stress Verification
	103	+
	104	+`stress.sh` repeats on the same channel/SQ, and every round waits for the receiver to finish checking. After each check the receiver zeroes the target buffer so a later round that fails to submit cannot be masked by earlier results. Every mode runs 100 times by default:
	105	+
	106	+```bash
	107	+./stress.sh 2
	108	+./stress.sh 2 warp-pad 1000
	109	+```
	110	+
	111	+## Performance Test
	112	+
	113	+`perf.sh` uses the device-side `clock()`. The timed loop calls `Write` to build the WQE and records `write_cycles` when `Write` returns, then performs the commit and records `commit_cycles` separately, and finally records the full completion time `completion_cycles` after `Drain` returns, so `average_write_us` excludes the commit time. The default is a warmup of 100 followed by 1000 timed iterations, with one kernel launch per warmup/timed sample; every round drains so the SQ tail keeps advancing. Host launch and synchronization time is not included in the device cycles above:
	114	+
	115	+```bash
	116	+./perf.sh 2
	117	+./perf.sh 2 warp 10000 1000
	118	+```
	119	+
	120	+The cycles in the results are converted to microseconds on the assumption of a 1 GHz device clock, matching the statistics of `simt_urma_perftest`.
	121	+
	122	+## Expected Results
	123	+
	124	+In correctness modes, rank 1 verifies every slot of the receive buffer, for example:
	125	+
	126	+```text
	127	+[rank 1] simt_jetty mode 4 iteration 0 PASS
	128	+```
	129	+
	130	+The sample succeeds when every requested mode prints `PASS`.
	131	+
	132	+## Notes
	133	+
	134	+- Load the CANN environment variables before building or running so that `ASCEND_CANN_PACKAGE_PATH`, the ASC CMake modules, and the runtime libraries are available.
	135	+- This sample depends on the `hccl`, `hcomm`, `ascendcl`, and `runtime` libraries in the CANN package.
	136	+- `nranks` must be at least 2; the sample only actually uses rank 0 and rank 1, and the remaining ranks exit directly.
	137	+- This sample only supports single-node multi-card execution and does not support multi-node execution.

Aexamples/simt_jetty/build.sh+38-0

@@ -0,0 +1,38 @@
	1	+#!/usr/bin/env bash
	2	+# ----------------------------------------------------------------------------------------------------------
	3	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	5	+# CANN Open Software License Agreement Version 2.0 (the "License").
	6	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	7	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	8	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	9	+# See LICENSE in the root of the software repository for the full text of the License.
	10	+# ----------------------------------------------------------------------------------------------------------
	11	+
	12	+set -euo pipefail
	13	+
	14	+if [[ -z "${ASCEND_HOME_PATH:-}" ]]; then
	15	+ echo "ASCEND_HOME_PATH is not set. Source the CANN environment first:" >&2
	16	+ echo " source /usr/local/Ascend/cann/set_env.sh" >&2
	17	+ exit 1
	18	+fi
	19	+
	20	+script_dir="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
	21	+build_dir="${script_dir}/build"
	22	+build_jobs="${BUILD_JOBS:-$(nproc)}"
	23	+res_usage="${CCE_RES_USAGE:-0}"
	24	+
	25	+if [[ "${res_usage}" != "0" && "${res_usage}" != "1" ]]; then
	26	+ echo "CCE_RES_USAGE must be 0 or 1" >&2
	27	+ exit 1
	28	+fi
	29	+
	30	+rm -rf -- "${build_dir}"
	31	+
	32	+cmake \
	33	+ -S "${script_dir}" \
	34	+ -B "${build_dir}" \
	35	+ -DCMAKE_ASC_ARCHITECTURES=dav-3510 \
	36	+ -DHCOMM_CCE_RES_USAGE="${res_usage}"
	37	+
	38	+cmake --build "${build_dir}" --verbose --parallel "${build_jobs}"

Aexamples/simt_jetty/main.cpp+397-0

@@ -0,0 +1,397 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#include <chrono>
	12	+#include <cerrno>
	13	+#include <cstdint>
	14	+#include <cstdio>
	15	+#include <cstdlib>
	16	+#include <cstring>
	17	+#include <fstream>
	18	+#include <iomanip>
	19	+#include <iostream>
	20	+#include <string>
	21	+#include <thread>
	22	+
	23	+#include <sys/stat.h>
	24	+#include <sys/types.h>
	25	+
	26	+#include "hccl/hccl.h"
	27	+#include "hccl/hccl_rank_graph.h"
	28	+#include "hccl/hccl_res.h"
	29	+#include "hccl/hccl_types.h"
	30	+
	31	+#include "simt_jetty_common.h"
	32	+
	33	+#define ACLCHECK(ret) \
	34	+ do { \
	35	+ if ((ret) != ACL_SUCCESS) { \
	36	+ printf("acl interface return err %s:%d, retcode: %d \n", __FILE__, __LINE__, (ret)); \
	37	+ return (ret); \
	38	+ } \
	39	+ } while (0)
	40	+
	41	+#define HCCLCHECK(ret) \
	42	+ do { \
	43	+ if ((ret) != HCCL_SUCCESS) { \
	44	+ printf("hccl interface return err %s:%d, retcode: %d \n", __FILE__, __LINE__, (ret)); \
	45	+ return (ret); \
	46	+ } \
	47	+ } while (0)
	48	+
	49	+extern void LaunchSimtJetty(
	50	+ void* stream, uint64_t channel, void* jettyPtrTable, uint32_t sendBufIdx, uint32_t recvBufIdx, uint32_t mode,
	51	+ void* timing, uint32_t iterations, uint32_t warmup);
	52	+
	53	+namespace {
	54	+
	55	+constexpr uint32_t kBarrierWaitSeconds = 60U;
	56	+constexpr uint32_t kBarrierPollMs = 100U;
	57	+constexpr uint32_t kIterationPollMs = 1U;
	58	+constexpr CommProtocol kTargetCommProtocol = COMM_PROTOCOL_UBC_CTP;
	59	+constexpr const char* kSendBufTag = "simtJettySendBuf";
	60	+constexpr const char* kRecvBufTag = "simtJettyRecvBuf";
	61	+
	62	+const char* ModeName(uint32_t mode)
	63	+{
	64	+ constexpr const char* kModeNames[] = {"thread", "warp-pad", "group", "dispatch", "warp", "mixed"};
	65	+ return mode < sizeof(kModeNames) / sizeof(kModeNames[0]) ? kModeNames[mode] : "unknown";
	66	+}
	67	+
	68	+bool IsMarkerFileReady(const std::string& path)
	69	+{
	70	+ struct stat st {};
	71	+ return stat(path.c_str(), &st) == 0 && st.st_size >= 1;
	72	+}
	73	+
	74	+HcclResult WriteMarkerFile(const std::string& path)
	75	+{
	76	+ const std::string tmpPath = path + ".tmp";
	77	+ std::ofstream ofs(tmpPath, std::ios::binary \| std::ios::trunc);
	78	+ if (!ofs.is_open()) {
	79	+ return HCCL_E_OPEN_FILE_FAILURE;
	80	+ }
	81	+ ofs << '1';
	82	+ ofs.close();
	83	+ if (!ofs) {
	84	+ return HCCL_E_SYSCALL;
	85	+ }
	86	+ if (std::rename(tmpPath.c_str(), path.c_str()) != 0) {
	87	+ return HCCL_E_SYSCALL;
	88	+ }
	89	+ return HCCL_SUCCESS;
	90	+}
	91	+
	92	+HcclResult WaitMarkerFile(const std::string& path, uint32_t pollMs = kBarrierPollMs)
	93	+{
	94	+ const uint32_t maxPollTimes = kBarrierWaitSeconds * 1000U / pollMs;
	95	+ for (uint32_t i = 0; i < maxPollTimes; ++i) {
	96	+ if (IsMarkerFileReady(path)) {
	97	+ return HCCL_SUCCESS;
	98	+ }
	99	+ std::this_thread::sleep_for(std::chrono::milliseconds(pollMs));
	100	+ }
	101	+ return HCCL_E_TIMEOUT;
	102	+}
	103	+
	104	+HcclResult HostBarrier(int rank, int nranks, const std::string& syncDir, const std::string& tag)
	105	+{
	106	+ if (mkdir(syncDir.c_str(), 0700) != 0 && errno != EEXIST) {
	107	+ return HCCL_E_SYSCALL;
	108	+ }
	109	+ const std::string marker = syncDir + "/" + tag + "." + std::to_string(rank);
	110	+ HcclResult ret = WriteMarkerFile(marker);
	111	+ if (ret != HCCL_SUCCESS) {
	112	+ return ret;
	113	+ }
	114	+ const std::string prefix = syncDir + "/" + tag + ".";
	115	+ const uint32_t maxPollTimes = kBarrierWaitSeconds * 1000U / kBarrierPollMs;
	116	+ for (uint32_t poll = 0; poll < maxPollTimes; ++poll) {
	117	+ bool ready = true;
	118	+ for (int i = 0; i < nranks; ++i) {
	119	+ if (!IsMarkerFileReady(prefix + std::to_string(i))) {
	120	+ ready = false;
	121	+ break;
	122	+ }
	123	+ }
	124	+ if (ready) {
	125	+ return HCCL_SUCCESS;
	126	+ }
	127	+ std::this_thread::sleep_for(std::chrono::milliseconds(kBarrierPollMs));
	128	+ }
	129	+ return HCCL_E_TIMEOUT;
	130	+}
	131	+
	132	+HcclResult InitCommByRootInfo(uint32_t rank, uint32_t nranks, const std::string& rootInfoFile, HcclComm* comm)
	133	+{
	134	+ HcclRootInfo rootInfo{};
	135	+ HcclResult ret = HCCL_SUCCESS;
	136	+ if (rank == 0U) {
	137	+ HCCLCHECK(HcclGetRootInfo(&rootInfo));
	138	+ const std::string tmpPath = rootInfoFile + ".tmp";
	139	+ std::ofstream ofs(tmpPath, std::ios::binary \| std::ios::trunc);
	140	+ if (!ofs.is_open()) {
	141	+ return HCCL_E_OPEN_FILE_FAILURE;
	142	+ }
	143	+ ofs.write(reinterpret_cast<const char*>(&rootInfo), sizeof(rootInfo));
	144	+ ofs.close();
	145	+ if (!ofs) {
	146	+ return HCCL_E_SYSCALL;
	147	+ }
	148	+ if (std::rename(tmpPath.c_str(), rootInfoFile.c_str()) != 0) {
	149	+ return HCCL_E_SYSCALL;
	150	+ }
	151	+ } else {
	152	+ const uint32_t maxPollTimes = kBarrierWaitSeconds * 1000U / kBarrierPollMs;
	153	+ for (uint32_t i = 0; i < maxPollTimes; ++i) {
	154	+ struct stat st {};
	155	+ if (stat(rootInfoFile.c_str(), &st) == 0 && st.st_size >= static_cast<off_t>(sizeof(HcclRootInfo))) {
	156	+ break;
	157	+ }
	158	+ std::this_thread::sleep_for(std::chrono::milliseconds(kBarrierPollMs));
	159	+ }
	160	+ std::ifstream ifs(rootInfoFile, std::ios::binary);
	161	+ if (!ifs.is_open()) {
	162	+ return HCCL_E_OPEN_FILE_FAILURE;
	163	+ }
	164	+ ifs.read(reinterpret_cast<char*>(&rootInfo), sizeof(rootInfo));
	165	+ if (ifs.gcount() != static_cast<std::streamsize>(sizeof(rootInfo))) {
	166	+ return HCCL_E_SYSCALL;
	167	+ }
	168	+ }
	169	+ HCCLCHECK(HcclCommInitRootInfo(nranks, &rootInfo, rank, comm));
	170	+ return HCCL_SUCCESS;
	171	+}
	172	+
	173	+HcclResult AcquirePeerChannel(
	174	+ HcclComm comm, uint32_t rank, uint32_t peerRank, HcclMemHandle* memHandles, ChannelHandle& channel,
	175	+ uint32_t& remoteRecvIdx)
	176	+{
	177	+ HcclChannelDesc desc;
	178	+ HCCLCHECK(HcclChannelDescInit(&desc, 1));
	179	+ desc.remoteRank = peerRank;
	180	+ desc.channelProtocol = kTargetCommProtocol;
	181	+ desc.notifyNum = 0;
	182	+ desc.memHandles = memHandles;
	183	+ desc.memHandleNum = 2U;
	184	+
	185	+ uint32_t* layers = nullptr;
	186	+ uint32_t layerNum = 0U;
	187	+ HCCLCHECK(HcclRankGraphGetLayers(comm, &layers, &layerNum));
	188	+ if (layerNum == 0U \|\| layers == nullptr) {
	189	+ return HCCL_E_INTERNAL;
	190	+ }
	191	+ CommLink* links = nullptr;
	192	+ uint32_t linkNum = 0U;
	193	+ HCCLCHECK(HcclRankGraphGetLinks(comm, layers[0], rank, peerRank, &links, &linkNum));
	194	+ bool found = false;
	195	+ for (uint32_t i = 0U; i < linkNum; ++i) {
	196	+ if (links[i].linkAttr.linkProtocol == kTargetCommProtocol) {
	197	+ desc.localEndpoint = links[i].srcEndpointDesc;
	198	+ desc.remoteEndpoint = links[i].dstEndpointDesc;
	199	+ found = true;
	200	+ break;
	201	+ }
	202	+ }
	203	+ if (!found) {
	204	+ return HCCL_E_NOT_SUPPORT;
	205	+ }
	206	+ HCCLCHECK(HcclChannelAcquire(comm, COMM_ENGINE_AIV, &desc, 1U, &channel));
	207	+
	208	+ uint32_t memNum = 0U;
	209	+ CommMem* remoteMems = nullptr;
	210	+ char** tags = nullptr;
	211	+ HCCLCHECK(HcclChannelGetRemoteMems(comm, channel, &memNum, &remoteMems, &tags));
	212	+ for (uint32_t i = 0U; i < memNum; ++i) {
	213	+ if (tags[i] != nullptr && std::strcmp(tags[i], kRecvBufTag) == 0) {
	214	+ remoteRecvIdx = i;
	215	+ return HCCL_SUCCESS;
	216	+ }
	217	+ }
	218	+ return HCCL_E_INTERNAL;
	219	+}
	220	+
	221	+bool CheckWords(void* deviceBuf, const uint64_t expected[simt_jetty::kSlotCount])
	222	+{
	223	+ uint64_t actual[simt_jetty::kSlotCount] = {};
	224	+ if (aclrtMemcpy(actual, simt_jetty::kBufferBytes, deviceBuf, simt_jetty::kBufferBytes, ACL_MEMCPY_DEVICE_TO_HOST) !=
	225	+ ACL_SUCCESS) {
	226	+ return false;
	227	+ }
	228	+ for (uint32_t i = 0; i < simt_jetty::kSlotCount; ++i) {
	229	+ if (actual[i] != expected[i]) {
	230	+ std::cout << "slot " << i << " actual=" << actual[i] << " expected=" << expected[i] << std::endl;
	231	+ return false;
	232	+ }
	233	+ }
	234	+ return true;
	235	+}
	236	+
	237	+} // namespace
	238	+
	239	+int main(int argc, char* argv[])
	240	+{
	241	+ if (argc != 6 && argc != 9) {
	242	+ std::cout << "Usage: " << argv[0]
	243	+ << " <rank> <nranks> <root_info_file> <sync_dir> <mode> [iterations warmup timing]" << std::endl;
	244	+ return 1;
	245	+ }
	246	+
	247	+ const int rank = std::atoi(argv[1]);
	248	+ const int nranks = std::atoi(argv[2]);
	249	+ const std::string rootInfoFile = argv[3];
	250	+ const std::string syncDir = argv[4];
	251	+ const uint32_t mode = static_cast<uint32_t>(std::strtoul(argv[5], nullptr, 10));
	252	+ const uint32_t iterations = argc == 9 ? static_cast<uint32_t>(std::strtoul(argv[6], nullptr, 10)) : 1U;
	253	+ const uint32_t warmup = argc == 9 ? static_cast<uint32_t>(std::strtoul(argv[7], nullptr, 10)) : 0U;
	254	+ const bool deviceTiming = argc == 9 && std::strtoul(argv[8], nullptr, 10) != 0U;
	255	+
	256	+ if (nranks < 2 \|\| rank < 0 \|\| rank >= nranks \|\| syncDir.empty() \|\| mode > 5U \|\| iterations == 0U) {
	257	+ std::cout << "invalid rank/nranks/sync_dir/mode/iterations" << std::endl;
	258	+ return 1;
	259	+ }
	260	+ if (rank >= 2) {
	261	+ std::cout << "SKIP \| Rank=" << rank << " \| Only rank 0 and rank 1 run simt_jetty." << std::endl;
	262	+ return 0;
	263	+ }
	264	+
	265	+ ACLCHECK(aclInit(nullptr));
	266	+ ACLCHECK(aclrtSetDevice(rank));
	267	+ aclrtStream stream = nullptr;
	268	+ ACLCHECK(aclrtCreateStream(&stream));
	269	+
	270	+ HcclComm comm = nullptr;
	271	+ HCCLCHECK(InitCommByRootInfo(rank, nranks, rootInfoFile, &comm));
	272	+
	273	+ void* sendBuf = nullptr;
	274	+ void* recvBuf = nullptr;
	275	+ ACLCHECK(aclrtMalloc(&sendBuf, simt_jetty::kBufferBytes, ACL_MEM_MALLOC_HUGE_FIRST));
	276	+ ACLCHECK(aclrtMalloc(&recvBuf, simt_jetty::kBufferBytes, ACL_MEM_MALLOC_HUGE_FIRST));
	277	+ uint64_t sendInit[simt_jetty::kSlotCount] = {};
	278	+ for (uint32_t slot = 0; slot < simt_jetty::kSlotCount; ++slot) {
	279	+ sendInit[slot] = simt_jetty::SlotValue(slot);
	280	+ }
	281	+ uint64_t recvInit[simt_jetty::kSlotCount] = {};
	282	+ ACLCHECK(
	283	+ aclrtMemcpy(sendBuf, simt_jetty::kBufferBytes, sendInit, simt_jetty::kBufferBytes, ACL_MEMCPY_HOST_TO_DEVICE));
	284	+ ACLCHECK(
	285	+ aclrtMemcpy(recvBuf, simt_jetty::kBufferBytes, recvInit, simt_jetty::kBufferBytes, ACL_MEMCPY_HOST_TO_DEVICE));
	286	+
	287	+ CommMem sendMem{COMM_MEM_TYPE_DEVICE, sendBuf, simt_jetty::kBufferBytes};
	288	+ CommMem recvMem{COMM_MEM_TYPE_DEVICE, recvBuf, simt_jetty::kBufferBytes};
	289	+ HcclMemHandle memHandles[2] = {nullptr, nullptr};
	290	+ HCCLCHECK(HcclCommMemReg(comm, kSendBufTag, &sendMem, &memHandles[0]));
	291	+ HCCLCHECK(HcclCommMemReg(comm, kRecvBufTag, &recvMem, &memHandles[1]));
	292	+
	293	+ const uint32_t peerRank = static_cast<uint32_t>(rank == 0 ? 1 : 0);
	294	+ ChannelHandle channel = 0U;
	295	+ uint32_t remoteRecvIdx = 0U;
	296	+ HCCLCHECK(AcquirePeerChannel(comm, static_cast<uint32_t>(rank), peerRank, memHandles, channel, remoteRecvIdx));
	297	+
	298	+ void* jettyPtrTable = nullptr;
	299	+ ACLCHECK(aclrtMalloc(&jettyPtrTable, sizeof(uint64_t), ACL_MEM_MALLOC_HUGE_FIRST));
	300	+ const uint64_t channelValue = static_cast<uint64_t>(channel);
	301	+ ACLCHECK(aclrtMemcpy(
	302	+ jettyPtrTable, sizeof(channelValue), &channelValue, sizeof(channelValue), ACL_MEMCPY_HOST_TO_DEVICE));
	303	+
	304	+ HCCLCHECK(HostBarrier(rank, nranks, syncDir, "ready"));
	305	+
	306	+ constexpr uint32_t kLocalSendIdx = 1U;
	307	+ bool pass = true;
	308	+ void* timingBuf = nullptr;
	309	+ uint64_t timing[4] = {};
	310	+ if (rank == 0) {
	311	+ if (deviceTiming) {
	312	+ ACLCHECK(aclrtMalloc(&timingBuf, sizeof(timing), ACL_MEM_MALLOC_HUGE_FIRST));
	313	+ ACLCHECK(aclrtMemset(timingBuf, sizeof(timing), 0, sizeof(timing)));
	314	+ LaunchSimtJetty(
	315	+ stream, static_cast<uint64_t>(channel), jettyPtrTable, kLocalSendIdx, remoteRecvIdx, mode, timingBuf,
	316	+ iterations, warmup);
	317	+ ACLCHECK(aclrtSynchronizeStream(stream));
	318	+ ACLCHECK(aclrtMemcpy(timing, sizeof(timing), timingBuf, sizeof(timing), ACL_MEMCPY_DEVICE_TO_HOST));
	319	+ pass = timing[3] == iterations;
	320	+ const double averageWriteCycles = static_cast<double>(timing[0]) / iterations;
	321	+ const double averageCommitCycles = static_cast<double>(timing[1]) / iterations;
	322	+ const double averageCompletionCycles = static_cast<double>(timing[2]) / iterations;
	323	+ std::cout << std::fixed << std::setprecision(3) << "PERF \| mode=" << ModeName(mode)
	324	+ << " \| iterations=" << iterations << " \| warmup=" << warmup << " \| write_cycles=" << timing[0]
	325	+ << " \| average_write_cycles=" << averageWriteCycles
	326	+ << " \| average_write_us=" << averageWriteCycles / 1000.0 << " \| commit_cycles=" << timing[1]
	327	+ << " \| average_commit_cycles=" << averageCommitCycles
	328	+ << " \| average_commit_us=" << averageCommitCycles / 1000.0 << " \| completion_cycles=" << timing[2]
	329	+ << " \| average_completion_cycles=" << averageCompletionCycles
	330	+ << " \| average_completion_us=" << averageCompletionCycles / 1000.0 << " \| completed=" << timing[3]
	331	+ << std::endl;
	332	+ } else {
	333	+ for (uint32_t iteration = 0U; iteration < iterations; ++iteration) {
	334	+ LaunchSimtJetty(
	335	+ stream, static_cast<uint64_t>(channel), jettyPtrTable, kLocalSendIdx, remoteRecvIdx, mode, nullptr,
	336	+ 1U, 0U);
	337	+ ACLCHECK(aclrtSynchronizeStream(stream));
	338	+ HCCLCHECK(WriteMarkerFile(syncDir + "/done." + std::to_string(iteration)));
	339	+ HCCLCHECK(WaitMarkerFile(syncDir + "/checked." + std::to_string(iteration), kIterationPollMs));
	340	+ }
	341	+ }
	342	+ if (deviceTiming) {
	343	+ HCCLCHECK(WriteMarkerFile(syncDir + "/done.0"));
	344	+ HCCLCHECK(WaitMarkerFile(syncDir + "/checked.0", kIterationPollMs));
	345	+ }
	346	+ } else {
	347	+ uint64_t expected[simt_jetty::kSlotCount] = {};
	348	+ if (mode == 0U \|\| mode == 3U) {
	349	+ for (uint32_t slot = 0; slot < simt_jetty::kSlotCount; ++slot) {
	350	+ expected[slot] = simt_jetty::SlotValue(slot);
	351	+ }
	352	+ } else if (mode == 1U) {
	353	+ expected[0] = simt_jetty::SlotValue(0);
	354	+ } else if (mode == 2U \|\| mode == 4U \|\| mode == 5U) {
	355	+ // group/warp/mixed posts carry kConfiguredSgeNum SGEs; warp-pad keeps slot 0 only.
	356	+ for (uint32_t slot = 0; slot < simt_jetty::kConfiguredSgeNum; ++slot) {
	357	+ expected[slot] = simt_jetty::SlotValue(slot);
	358	+ }
	359	+ } else {
	360	+ expected[0] = simt_jetty::SlotValue(0);
	361	+ expected[1] = simt_jetty::SlotValue(1);
	362	+ }
	363	+ const uint32_t checks = deviceTiming ? 1U : iterations;
	364	+ for (uint32_t iteration = 0U; iteration < checks; ++iteration) {
	365	+ HCCLCHECK(
	366	+ WaitMarkerFile(syncDir + "/done." + std::to_string(deviceTiming ? 0U : iteration), kIterationPollMs));
	367	+ const bool iterationPass = CheckWords(recvBuf, expected);
	368	+ pass = pass && iterationPass;
	369	+ if (!iterationPass) {
	370	+ std::cout << "[rank 1] simt_jetty mode " << mode << " iteration " << iteration << " FAIL"
	371	+ << " checks="
	372	+ << (mode == 0U \|\| mode == 3U ? simt_jetty::kSlotCount :
	373	+ mode == 1U ? 1U :
	374	+ mode == 2U \|\| mode == 4U \|\| mode == 5U ? simt_jetty::kConfiguredSgeNum :
	375	+ 2U)
	376	+ << std::endl;
	377	+ }
	378	+ if (!deviceTiming && iteration + 1U < checks) {
	379	+ ACLCHECK(aclrtMemset(recvBuf, simt_jetty::kBufferBytes, 0, simt_jetty::kBufferBytes));
	380	+ }
	381	+ HCCLCHECK(WriteMarkerFile(syncDir + "/checked." + std::to_string(deviceTiming ? 0U : iteration)));
	382	+ }
	383	+ std::cout << "[rank 1] simt_jetty mode " << mode << " iterations " << checks << (pass ? " PASS" : " FAIL")
	384	+ << std::endl;
	385	+ }
	386	+
	387	+ if (timingBuf != nullptr) {
	388	+ ACLCHECK(aclrtFree(timingBuf));
	389	+ }
	390	+ ACLCHECK(aclrtFree(jettyPtrTable));
	391	+ ACLCHECK(aclrtFree(sendBuf));
	392	+ ACLCHECK(aclrtFree(recvBuf));
	393	+ HCCLCHECK(HcclCommDestroy(comm));
	394	+ ACLCHECK(aclrtDestroyStream(stream));
	395	+ ACLCHECK(aclFinalize());
	396	+ return pass ? 0 : 2;
	397	+}

Aexamples/simt_jetty/op_kernel.cpp+516-0

@@ -0,0 +1,516 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#include <cstdio>
	12	+
	13	+#include "kernel_operator.h"
	14	+#include "simt_api/asc_simt.h"
	15	+
	16	+#include "asccomm_staged/hcomm/hcomm_simt.h"
	17	+#include "asccomm_staged/jetty/hcomm_jetty.h"
	18	+#include "asccomm_staged/jetty/hcomm_jetty_simt.h"
	19	+#include "simt_jetty_common.h"
	20	+
	21	+namespace {
	22	+
	23	+using namespace simt_jetty;
	24	+
	25	+struct alignas(16) SimtJettyUbLayout {
	26	+ AscendC::simt::jetty::HcommJettyInfo jetty_info;
	27	+ AscendC::simt::jetty::HcommPeerInfo peer_info;
	28	+ alignas(16) uint64_t dwqe_stage[2 * AscendC::simt::jetty::HcommJetty::kNumWQEBBLanes] = {};
	29	+ uint32_t mixed_ready = 0U;
	30	+};
	31	+
	32	+static_assert(
	33	+ sizeof(AscendC::simt::jetty::HcommJettyInfo) == sizeof(AscendC::HcommJettyInfo),
	34	+ "SIMT and SIMD Jetty metadata must share an ABI");
	35	+static_assert(
	36	+ offsetof(AscendC::simt::jetty::HcommJettyInfo, packedHead) == offsetof(AscendC::HcommJettyInfo, packedHead),
	37	+ "SIMT/SIMD packed head offset mismatch");
	38	+static_assert(
	39	+ sizeof(AscendC::simt::jetty::HcommPeerInfo) == sizeof(AscendC::HcommJettyPeerInfo),
	40	+ "SIMT and SIMD peer metadata must share an ABI");
	41	+
	42	+__simt_callee__ __forceinline__ __ubuf__ SimtJettyUbLayout* GetSimtJettyUbLayout()
	43	+{
	44	+ return reinterpret_cast<__ubuf__ SimtJettyUbLayout*>(0);
	45	+}
	46	+
	47	+template <typename Coop>
	48	+__simt_callee__ __forceinline__ void PrepareJetty(
	49	+ const Coop& coop, uint64_t channel, __gm__ void* jettyPtrTable, __ubuf__ SimtJettyUbLayout* ub_layout,
	50	+ AscendC::simt::jetty::HcommPeer& peer, AscendC::simt::jetty::HcommJetty& jetty)
	51	+{
	52	+ if (coop.thread_rank() == 0) {
	53	+ AscendC::simt::jetty::HcommPeer peer_init(&ub_layout->peer_info, jettyPtrTable, 0);
	54	+ AscendC::simt::jetty::HcommJetty jetty_init(
	55	+ &ub_layout->jetty_info, reinterpret_cast<__ubuf__ uint8_t*>(ub_layout->dwqe_stage), jettyPtrTable, 0);
	56	+ }
	57	+ coop.sync();
	58	+ peer.Attach(&ub_layout->peer_info, channel);
	59	+ jetty.Attach(&ub_layout->jetty_info, reinterpret_cast<__ubuf__ uint8_t*>(ub_layout->dwqe_stage), channel);
	60	+}
	61	+
	62	+template <int kNumSGEs, bool kPadSecondSge, bool kMeasure, bool kStandaloneWrite = false>
	63	+__simt_callee__ inline void SimtJettyWriteThreadImpl(
	64	+ uint64_t channel, uint32_t sendBufIdx, uint32_t recvBufIdx, const AscendC::simt::jetty::HcommPeer& peer,
	65	+ AscendC::simt::jetty::HcommJetty& jetty, uint64_t& writeCycles, uint64_t& commitCycles, uint64_t& completionCycles)
	66	+{
	67	+ __gm__ uint8_t* localSrc = AscendC::simt::LocalBufferAddr(channel, sendBufIdx);
	68	+ __gm__ uint8_t* remoteDst = AscendC::simt::RemoteBufferAddr(channel, recvBufIdx);
	69	+
	70	+ static_assert(kNumSGEs >= 2 && kNumSGEs <= 12, "thread SGE count must be in [2, 12]");
	71	+ AscendC::simt::jetty::HcommSge sges[kNumSGEs] = {};
	72	+#pragma unroll
	73	+ for (int sge_idx = 0; sge_idx < kNumSGEs; ++sge_idx) {
	74	+ sges[sge_idx].addr = reinterpret_cast<uint64_t>(localSrc) + static_cast<uint64_t>(sge_idx) * kSlotBytes;
	75	+ sges[sge_idx].len = static_cast<int>(kSlotBytes);
	76	+ }
	77	+ if constexpr (kPadSecondSge) {
	78	+ sges[1].addr = 0;
	79	+ sges[1].len = 0;
	80	+ }
	81	+
	82	+ uint64_t begin = 0U;
	83	+ if constexpr (kMeasure) {
	84	+ begin = clock();
	85	+ }
	86	+ // The SIMT DWQE hardware path is limited to a two-BB WRITE. Split larger
	87	+ // thread payloads into 2-SGE segments while preserving the contiguous
	88	+ // remote layout expected by the sample.
	89	+ constexpr int kSegmentSges = 2;
	90	+ for (int segment = 0; segment < kNumSGEs; segment += kSegmentSges) {
	91	+ AscendC::simt::jetty::HcommSge segmentSges[kSegmentSges] = {sges[segment], sges[segment + 1]};
	92	+ __gm__ void* segmentDst = remoteDst + static_cast<uint64_t>(segment) * kSlotBytes;
	93	+ if constexpr (kStandaloneWrite) {
	94	+ jetty.template Write<kSegmentSges, false>(peer, segmentDst, segmentSges);
	95	+ jetty.AdvanceSq();
	96	+ // Each split segment is its own implicit batch; publish it so the NIC
	97	+ // fetches before the next segment overwrites the two-BB DWQE window.
	98	+ jetty.PublishSq();
	99	+ } else {
	100	+ jetty.template Put<kSegmentSges>(peer, segmentDst, segmentSges);

Aexamples/simt_jetty/perf.sh+25-0

@@ -0,0 +1,25 @@
	1	+#!/usr/bin/env bash
	2	+# ----------------------------------------------------------------------------------------------------------
	3	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	5	+# CANN Open Software License Agreement Version 2.0 (the "License").
	6	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	7	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	8	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	9	+# See LICENSE in the root of the software repository for the full text of the License.
	10	+# ----------------------------------------------------------------------------------------------------------
	11	+
	12	+set -euo pipefail
	13	+
	14	+if [ "$#" -lt 1 ] \|\| [ "$#" -gt 4 ]; then
	15	+ echo "Usage: $0 <nranks> [all\|thread\|warp\|warp-pad\|group\|dispatch] [iterations] [warmup]"
	16	+ exit 1
	17	+fi
	18	+
	19	+script_dir="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
	20	+nranks="$1"
	21	+mode="${2:-all}"
	22	+iterations="${3:-1000}"
	23	+warmup="${4:-100}"
	24	+
	25	+exec "${script_dir}/run.sh" "$nranks" "$mode" --iterations "$iterations" --warmup "$warmup" --timing

Aexamples/simt_jetty/run.sh+168-0

@@ -0,0 +1,168 @@
	1	+#!/usr/bin/env bash
	2	+# ----------------------------------------------------------------------------------------------------------
	3	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	5	+# CANN Open Software License Agreement Version 2.0 (the "License").
	6	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	7	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	8	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	9	+# See LICENSE in the root of the software repository for the full text of the License.
	10	+# ----------------------------------------------------------------------------------------------------------
	11	+
	12	+set -euo pipefail
	13	+
	14	+usage() {
	15	+ cat <<'EOF'
	16	+Usage: run.sh <nranks> [all\|thread\|warp\|warp-pad\|group\|dispatch\|mixed] [options]
	17	+
	18	+Options:
	19	+ --iterations N Number of correctness iterations, default: 1
	20	+ --warmup N Warmup iterations for --timing, default: 0
	21	+ --timing Run all iterations in one kernel and report device clock cycles
	22	+EOF
	23	+}
	24	+
	25	+if [ "$#" -lt 1 ]; then
	26	+ usage
	27	+ exit 1
	28	+fi
	29	+
	30	+nranks="$1"
	31	+shift
	32	+requested_mode="all"
	33	+if [ "$#" -gt 0 ] && [[ "$1" != --* ]]; then
	34	+ requested_mode="$1"
	35	+ shift
	36	+fi
	37	+iterations=1
	38	+warmup=0
	39	+timing=0
	40	+while [ "$#" -gt 0 ]; do
	41	+ case "$1" in
	42	+ --iterations)
	43	+ [ "$#" -ge 2 ] \|\| { echo "--iterations requires a value" >&2; exit 1; }
	44	+ iterations="$2"
	45	+ shift 2
	46	+ ;;
	47	+ --warmup)
	48	+ [ "$#" -ge 2 ] \|\| { echo "--warmup requires a value" >&2; exit 1; }
	49	+ warmup="$2"
	50	+ shift 2
	51	+ ;;
	52	+ --timing) timing=1; shift ;;
	53	+ -h\|--help) usage; exit 0 ;;
	54	+ *) echo "unknown option: $1" >&2; usage; exit 1 ;;
	55	+ esac
	56	+done
	57	+
	58	+if ! [[ "$nranks" =~ ^[0-9]+$ ]] \|\| [ "$nranks" -lt 2 ]; then
	59	+ echo "nranks must be an integer greater than or equal to 2"
	60	+ exit 1
	61	+fi
	62	+if ! [[ "$iterations" =~ ^[0-9]+$ ]] \|\| [ "$iterations" -lt 1 ]; then
	63	+ echo "--iterations must be a positive integer"
	64	+ exit 1
	65	+fi
	66	+if ! [[ "$warmup" =~ ^[0-9]+$ ]]; then
	67	+ echo "--warmup must be a non-negative integer"
	68	+ exit 1
	69	+fi
	70	+if [ "$timing" -eq 0 ] && [ "$warmup" -ne 0 ]; then
	71	+ echo "--warmup requires --timing"
	72	+ exit 1
	73	+fi
	74	+if [ "$timing" -eq 1 ] && [ "$requested_mode" = "mixed" ]; then
	75	+ echo "mixed mode does not support device timing"
	76	+ exit 1
	77	+fi
	78	+
	79	+script_dir="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
	80	+binary="${script_dir}/build/simt_jetty"
	81	+
	82	+if [ ! -x "$binary" ]; then
	83	+ echo "binary does not exist or is not executable: $binary"
	84	+ echo "build it first: bash ${script_dir}/build.sh"
	85	+ exit 1
	86	+fi
	87	+
	88	+pids=()
	89	+active_dirs=()
	90	+cleanup() {
	91	+ if [ ${#pids[@]} -ne 0 ]; then
	92	+ kill "${pids[@]}" 2>/dev/null \|\| true
	93	+ fi
	94	+ if [ ${#active_dirs[@]} -ne 0 ]; then
	95	+ rm -rf "${active_dirs[@]}"
	96	+ fi
	97	+}
	98	+trap cleanup EXIT
	99	+trap 'exit 1' SIGINT SIGTERM
	100	+
	101	+mode_id() {
	102	+ case "$1" in
	103	+ thread) echo 0 ;;
	104	+ warp-pad) echo 1 ;;
	105	+ group) echo 2 ;;
	106	+ dispatch) echo 3 ;;
	107	+ warp) echo 4 ;;
	108	+ mixed) echo 5 ;;
	109	+ *) return 1 ;;
	110	+ esac
	111	+}
	112	+
	113	+run_mode() {
	114	+ local mode_name="$1"
	115	+ local id
	116	+ id="$(mode_id "$mode_name")"
	117	+ local root_info_dir
	118	+ local sync_dir
	119	+ root_info_dir="$(mktemp -d "${TMPDIR:-/tmp}/simt_jetty_root_info.XXXXXX")"
	120	+ sync_dir="$(mktemp -d "${TMPDIR:-/tmp}/simt_jetty_sync.XXXXXX")"
	121	+ active_dirs=("$root_info_dir" "$sync_dir")
	122	+ local root_info_file="${root_info_dir}/root_info.bin"
	123	+
	124	+ echo "[Running] mode=${mode_name} iterations=${iterations} warmup=${warmup} timing=${timing}"
	125	+ pids=()
	126	+ for ((rank = 0; rank < nranks; ++rank)); do
	127	+ "$binary" "$rank" "$nranks" "$root_info_file" "$sync_dir" "$id" "$iterations" "$warmup" "$timing" &
	128	+ pids+=("$!")
	129	+ done
	130	+
	131	+ local mode_status=0
	132	+ for pid in "${pids[@]}"; do
	133	+ if ! wait "$pid"; then
	134	+ mode_status=1
	135	+ fi
	136	+ done
	137	+ pids=()
	138	+ rm -rf "$root_info_dir" "$sync_dir"
	139	+ active_dirs=()
	140	+ return "$mode_status"
	141	+}
	142	+
	143	+status=0
	144	+if [ "$requested_mode" = "all" ]; then
	145	+ modes=(thread warp warp-pad group dispatch)
	146	+ if [ "$timing" -eq 0 ]; then
	147	+ modes+=(mixed)
	148	+ fi
	149	+else
	150	+ if ! mode_id "$requested_mode" >/dev/null; then
	151	+ usage
	152	+ exit 1
	153	+ fi
	154	+ modes=("$requested_mode")
	155	+fi
	156	+
	157	+for mode in "${modes[@]}"; do
	158	+ if ! run_mode "$mode"; then
	159	+ status=1
	160	+ fi
	161	+done
	162	+
	163	+if [ "$status" -eq 0 ]; then
	164	+ echo "RESULT \| Status=PASS"
	165	+else
	166	+ echo "RESULT \| Status=FAIL"
	167	+fi
	168	+exit "$status"

Aexamples/simt_jetty/simt_jetty_common.h+41-0

@@ -0,0 +1,41 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#ifndef EXAMPLES_SIMT_JETTY_SIMT_JETTY_COMMON_H
	12	+#define EXAMPLES_SIMT_JETTY_SIMT_JETTY_COMMON_H
	13	+
	14	+#include <cstdint>
	15	+
	16	+namespace simt_jetty {
	17	+
	18	+constexpr uint64_t kSlotBytes = sizeof(uint64_t);
	19	+constexpr uint32_t kSlotCount = 12U;
	20	+constexpr uint64_t kBufferBytes = kSlotCount * kSlotBytes;
	21	+
	22	+#ifndef SIMT_JETTY_SGE_NUM
	23	+#define SIMT_JETTY_SGE_NUM 12
	24	+#endif
	25	+
	26	+static_assert(
	27	+ SIMT_JETTY_SGE_NUM >= 1 && SIMT_JETTY_SGE_NUM <= static_cast<int>(kSlotCount),
	28	+ "SIMT_JETTY_SGE_NUM must be in the range [1, 12]");
	29	+constexpr uint32_t kConfiguredSgeNum = static_cast<uint32_t>(SIMT_JETTY_SGE_NUM);
	30	+
	31	+#if defined(__CCE_KT_TEST__) \|\| defined(__DAV_C310__) \|\| defined(__NPU_ARCH__)
	32	+#define SIMT_JETTY_DEVICE_FN __aicore__
	33	+#else
	34	+#define SIMT_JETTY_DEVICE_FN
	35	+#endif
	36	+
	37	+SIMT_JETTY_DEVICE_FN constexpr uint64_t SlotValue(uint32_t slot) { return static_cast<uint64_t>(slot) + 1U; }
	38	+
	39	+} // namespace simt_jetty
	40	+
	41	+#endif // EXAMPLES_SIMT_JETTY_SIMT_JETTY_COMMON_H

Aexamples/simt_jetty/stress.sh+24-0

@@ -0,0 +1,24 @@
	1	+#!/usr/bin/env bash
	2	+# ----------------------------------------------------------------------------------------------------------
	3	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	5	+# CANN Open Software License Agreement Version 2.0 (the "License").
	6	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	7	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	8	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	9	+# See LICENSE in the root of the software repository for the full text of the License.
	10	+# ----------------------------------------------------------------------------------------------------------
	11	+
	12	+set -euo pipefail
	13	+
	14	+if [ "$#" -lt 1 ] \|\| [ "$#" -gt 3 ]; then
	15	+ echo "Usage: $0 <nranks> [all\|thread\|warp\|warp-pad\|group\|dispatch\|mixed] [iterations]"
	16	+ exit 1
	17	+fi
	18	+
	19	+script_dir="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
	20	+nranks="$1"
	21	+mode="${2:-all}"
	22	+iterations="${3:-100}"
	23	+
	24	+exec "${script_dir}/run.sh" "$nranks" "$mode" --iterations "$iterations"

Minclude/aicore/jetty/hcomm_jetty.h+32-20

	@@ -12,41 +12,53 @@
12	* \file hcomm_jetty.h	12	* \file hcomm_jetty.h
13	* \brief Low-level AIV Jetty primitives for URMA point-to-point writes.	13	* \brief Low-level AIV Jetty primitives for URMA point-to-point writes.
14	*/	14	*/
15	-#ifndef INCLUDE_ADV_API_HCOMM_HCOMM_JETTY_H	15	+#ifndef INCLUDE_ADV_API_JETTY_HCOMM_JETTY_H
16	-#define INCLUDE_ADV_API_HCOMM_HCOMM_JETTY_H	16	+#define INCLUDE_ADV_API_JETTY_HCOMM_JETTY_H
		17	+
		18	+#include <cstddef>
17		19
18	#include "kernel_basic_intf.h"	20	#include "kernel_basic_intf.h"
19	#include "../hcomm/hcomm_common.h"	21	#include "../hcomm/hcomm_common.h"
20		22
21	namespace AscendC {	23	namespace AscendC {
22		24
		25	+// Shared Jetty metadata. Both execution models describe the same hardware queues and
		26	+// exchange these descriptors through UB, so the layout is defined once here and the
		27	+// SIMT implementation reuses it through hcomm_jetty_simt.h.
23	struct alignas(8) HcommJettyPeerInfo {	28	struct alignas(8) HcommJettyPeerInfo {
24	- uint32_t tpId;	29	+ // Packed to match the SQE wire layout so building a WQE header is a single 64-bit read:
25	- uint32_t remoteTokenId;	30	+ // tpId in bits [23:0], numSges in bits [31:24] (filled per-post, not stored here) and
		31	+ // remoteTokenId in bits [51:32].
		32	+ uint64_t tpIdNumSgesRemoteTokenId;
26	uint64_t remoteEid[2];	33	uint64_t remoteEid[2];
27	- uint64_t remoteTokenValue;	34	+ uint64_t remoteTokenValueUdf;
28	uint64_t remoteBaseAddr;	35	uint64_t remoteBaseAddr;
29	uint64_t remoteBufferSize;	36	uint64_t remoteBufferSize;
30	};	37	};
31	static_assert(sizeof(HcommJettyPeerInfo) == 48U, "HcommJettyPeerInfo must be 48 bytes");	38	static_assert(sizeof(HcommJettyPeerInfo) == 48U, "HcommJettyPeerInfo must be 48 bytes");
		39	+static_assert(offsetof(HcommJettyPeerInfo, remoteEid) == sizeof(uint64_t), "unexpected peer metadata layout");
		40	+static_assert(
		41	+ offsetof(HcommJettyPeerInfo, remoteTokenValueUdf) == 3U * sizeof(uint64_t), "unexpected peer metadata layout");
32		42
33	struct alignas(8) HcommJettyInfo {	43	struct alignas(8) HcommJettyInfo {
34	- uint64_t sqBaseAddr; // SQ base address in device memory.	44	+ uint64_t sqBaseAddr; // SQ base address in device memory.
35	- uint64_t sqHeadAddr; // Address of the packed 64-bit SQ head: low 32 bits are the head in	45	+ uint64_t sqHeadAddr; // Address of the packed 64-bit SQ head in GM.
36	- // WQEBBs, high 32 bits are the expected CQE count.	46	+ uint64_t completionTailAddr; // Address of the 32-bit SQ completion tail published by the hardware.
37	- uint64_t sqTailAddr; // Address of the 32-bit SQ completion tail published by the hardware.	47	+ uint64_t sqDoorbellAddr; // SQ doorbell address.
38	- uint64_t sqDoorbellAddr; // SQ doorbell address.	48	+ uint64_t cqBaseAddr; // CQ base address in device memory.
39	- uint64_t cqBaseAddr; // CQ base address in device memory.	49	+ uint64_t cqTailAddr; // Address of the 32-bit CQ tail.
40	- uint64_t cqTailAddr; // Address of the 32-bit CQ tail.	50	+ uint64_t cqDoorbellAddr; // CQ doorbell address.
41	- uint64_t cqDoorbellAddr; // CQ doorbell address.	51	+ uint64_t packedHead; // Producer state mirroring the GM SQ head: low 32 bits are the head
42	- uint32_t sqHead; // SQ producer head in WQEBBs.	52	+ // in WQEBBs, high 32 bits the expected CQE count. Kept packed so
43	- uint32_t expectedCqeCnt; // Number of CQEs expected from the submitted WQEs.	53	+ // publishing is a single 64-bit store.
44	- uint32_t numWqebbBytes; // WQEBB size in bytes.	54	+ uint32_t numWqebbBytes; // WQEBB size in bytes.
45	- uint32_t numCqeBytes; // CQE size in bytes.	55	+ uint32_t numCqeBytes; // CQE size in bytes.
46	- uint32_t sqDepth; // SQ depth in WQEBBs.	56	+ uint32_t sqDepth; // SQ depth in WQEBBs.
47	- uint32_t cqDepth; // CQ depth in CQEs.	57	+ uint32_t cqDepth; // CQ depth in CQEs.
48	};	58	};
49	static_assert(sizeof(HcommJettyInfo) == 80U, "HcommJettyInfo must be 80 bytes");	59	static_assert(sizeof(HcommJettyInfo) == 80U, "HcommJettyInfo must be 80 bytes");
		60	+static_assert(offsetof(HcommJettyInfo, packedHead) == 56U, "SIMT/SIMD packed head offset mismatch");
		61	+static_assert(offsetof(HcommJettyInfo, numWqebbBytes) == 64U, "HcommJettyInfo scalar fields offset mismatch");
50		62
51	struct HcommPeer {	63	struct HcommPeer {
52	__ubuf__ HcommJettyPeerInfo* peerInfo = nullptr;	64	__ubuf__ HcommJettyPeerInfo* peerInfo = nullptr;
	@@ -176,4 +188,4 @@ private:
176	#undef HCOMM_UNDEF_INCLUDE_INTERNAL_HEADERS_HCOMM_JETTY_H	188	#undef HCOMM_UNDEF_INCLUDE_INTERNAL_HEADERS_HCOMM_JETTY_H
177	#endif	189	#endif
178		190
179	-#endif // INCLUDE_ADV_API_HCOMM_HCOMM_JETTY_H	191	+#endif // INCLUDE_ADV_API_JETTY_HCOMM_JETTY_H

Ainclude/aicore/jetty/hcomm_jetty_simt.h+118-0

@@ -0,0 +1,118 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file hcomm_jetty_simt.h
	13	+ * \brief Low-level SIMT Jetty primitives for URMA point-to-point writes.
	14	+ */
	15	+#ifndef INCLUDE_ADV_API_JETTY_HCOMM_JETTY_SIMT_H
	16	+#define INCLUDE_ADV_API_JETTY_HCOMM_JETTY_SIMT_H
	17	+
	18	+#include <cstddef>
	19	+#include <cstdint>
	20	+
	21	+// The Jetty metadata structs are defined once in the SIMD public header and shared by
	22	+// both execution models; including it first makes them visible to the SIMT impl chain.
	23	+#include "hcomm_jetty.h"
	24	+#include "../../../impl/comm_api/aicore/jetty/impl/hcomm_jetty_simt_impl_def.h"
	25	+
	26	+namespace AscendC::simt::jetty {
	27	+
	28	+using HcommPeerInfo = AscendC::HcommJettyPeerInfo;
	29	+using HcommSge = JettySge;
	30	+
	31	+/**
	32	+ * Public SIMT Jetty API. A call to Write publishes by default; use
	33	+ * kDoCommit=false for a group of posts followed by AdvanceSq and PublishSq.
	34	+ */
	35	+class HcommJetty {
	36	+public:
	37	+ // Number of 64-bit words in one WQEBB.
	38	+ static constexpr uint32_t kNumWQEBBLanes = JettyImpl::kNumWQEBBLanes;
	39	+ // Number of WQEBBs in one SQE slot reservation.
	40	+ static constexpr uint32_t kNumWQEBBsPerSQESlot = JettyImpl::kNumWQEBBsPerSQESlot;
	41	+ // Maximum number of SGEs one Write can carry.
	42	+ static constexpr uint32_t kNumMaxSGEsPerWrite = JettyImpl::kNumMaxSGEsPerWrite;
	43	+ // Number of WQEBBs occupied by a Write WQE reservation.
	44	+ static constexpr uint32_t kNumWriteWqebbs = JettyImpl::kNumWriteWqebbs;
	45	+ // Number of WQEBBs occupied by a WriteWithNotify WQE.
	46	+ static constexpr uint32_t kNumWriteWithNotifyWqebbs = JettyImpl::kNumWriteWithNotifyWqebbs;
	47	+ // Size in bytes of the DWQE staging buffer callers must allocate.
	48	+ static constexpr uint32_t kNumMaxSQEBytes = JettyImpl::kNumMaxSQEBytes;
	49	+ // Number of WQEBBs one Write spans in the SQ.
	50	+ static constexpr uint32_t kNumWQEBBsPerWrite = JettyImpl::kNumWQEBBsPerWrite;
	51	+ // Number of 64-bit words in the two-BB DWQE window.
	52	+ static constexpr uint32_t kNumDwqeWords = JettyImpl::kNumDwqeWords;
	53	+
	54	+ __simt_callee__ __forceinline__ HcommJetty();
	55	+ __simt_callee__ __forceinline__
	56	+ HcommJetty(__ubuf__ HcommJettyInfo* info, __ubuf__ uint8_t* stage, __gm__ void* table, uint32_t index);
	57	+ __simt_callee__ __forceinline__
	58	+ HcommJetty(__ubuf__ HcommJettyInfo* info, __ubuf__ uint8_t* stage, uint64_t channelAddr);
	59	+
	60	+ __simt_callee__ __forceinline__ __ubuf__ HcommJettyInfo* Info() const;
	61	+
	62	+ __simt_callee__ __forceinline__ void Attach(
	63	+ __ubuf__ HcommJettyInfo* info, __ubuf__ uint8_t* stage, uint64_t channel);
	64	+
	65	+ template <uint32_t kSgeNum, bool kDoCommit = true, typename Coop = HcommCoopThread>
	66	+ __simt_callee__ __forceinline__ int32_t
	67	+ Write(const HcommPeer& peer, __gm__ void* dst, const JettySge* sges, const Coop& coop = Coop{}) const;
	68	+
	69	+ // Cooperative-group-first spellings, mirroring the cooperative_groups argument order.
	70	+ template <uint32_t kSgeNum, bool kDoCommit = true, typename Coop = HcommCoopThread>
	71	+ __simt_callee__ __forceinline__ int32_t
	72	+ Write(const Coop& coop, const HcommPeer& peer, __gm__ void* dst, const HcommSge& sge) const;
	73	+
	74	+ template <uint32_t kSgeNum, bool kDoCommit = true, typename Coop = HcommCoopThread>
	75	+ __simt_callee__ __forceinline__ int32_t
	76	+ Write(const Coop& coop, const HcommPeer& peer, __gm__ void* dst, const HcommSge* sges) const;
	77	+
	78	+ template <uint32_t kSgeNum, bool kDoCommit = true, typename Coop = HcommCoopThread>
	79	+ __simt_callee__ __forceinline__ int32_t
	80	+ Put(const HcommPeer& peer, __gm__ void* dst, const HcommSge* sges, const Coop& coop = Coop{}) const;
	81	+
	82	+ template <uint32_t kSgeNum, bool kDoCommit = true, typename Coop = HcommCoopThread>
	83	+ __simt_callee__ __forceinline__ int32_t
	84	+ Put(const Coop& coop, const HcommPeer& peer, __gm__ void* dst, const HcommSge& sge) const;
	85	+
	86	+ template <int64_t kTimeout, typename T, bool kDoCommit = true>
	87	+ __simt_callee__ __forceinline__ int32_t WriteValue(const HcommPeer& peer, __gm__ void* dst, T value);
	88	+
	89	+ template <int64_t kTimeout, bool kDoCommit = true>
	90	+ __simt_callee__ __forceinline__ int32_t WriteWithNotify(
	91	+ const HcommPeer& peer, __gm__ void* dst, __gm__ void* src, uint64_t len, __gm__ void* notifyAddr,
	92	+ uint64_t notifyValue);
	93	+
	94	+ __simt_callee__ __forceinline__ void AdvanceSq() const;
	95	+ template <typename Coop>
	96	+ __simt_callee__ __forceinline__ void AdvanceSq(const Coop& coop) const;
	97	+ template <typename Coop>
	98	+ __simt_callee__ __forceinline__ void AdvanceSq(const Coop& coop, uint32_t a, uint32_t b) const;
	99	+
	100	+ __simt_callee__ __forceinline__ void PublishSq() const;
	101	+ template <typename Coop>
	102	+ __simt_callee__ __forceinline__ void PublishSq(const Coop& coop) const;
	103	+
	104	+ template <int64_t kTimeout>
	105	+ __simt_callee__ __forceinline__ int32_t Drain() const;
	106	+ template <int64_t kTimeout, typename Coop>
	107	+ __simt_callee__ __forceinline__ int32_t Drain(const Coop& coop) const;
	108	+
	109	+private:
	110	+ JettyImpl impl_;
	111	+};
	112	+
	113	+} // namespace AscendC::simt::jetty
	114	+
	115	+// Must stay after class HcommJetty: it defines the member functions declared above.
	116	+#include "../../../impl/comm_api/aicore/jetty/impl/hcomm_jetty_simt_impl.h"
	117	+
	118	+#endif // INCLUDE_ADV_API_JETTY_HCOMM_JETTY_SIMT_H

Asrc/aicore/jetty/impl/hcomm_jetty_simt_impl.h+140-0

@@ -0,0 +1,140 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file hcomm_jetty_simt_impl.h
	13	+ * \brief Definitions of the public SIMT Jetty member functions.
	14	+ */
	15	+
	16	+#ifndef IMPL_ADV_API_DETAIL_HCOMM_IMPL_HCOMM_JETTY_SIMT_IMPL_H
	17	+#define IMPL_ADV_API_DETAIL_HCOMM_IMPL_HCOMM_JETTY_SIMT_IMPL_H
	18	+
	19	+#include "hcomm_jetty_simt_impl_def.h"
	20	+
	21	+// SIMT Jetty is only available on v310. The ASC driver compiles each translation unit twice:
	22	+// __NPU_ARCH__ is 3510 in the device pass and undefined in the host pass. The public HcommJetty
	23	+// declares its members out-of-line, so definitions are only emitted on platforms that have the
	24	+// implementation; other platforms fail at link time if the API is used.
	25	+#if !defined(__NPU_ARCH__) \|\| __NPU_ARCH__ == 3510
	26	+#include "platform_v310/hcomm_simt_urma_jetty.h"
	27	+#endif
	28	+
	29	+namespace AscendC::simt::jetty {
	30	+
	31	+#if !defined(__NPU_ARCH__) \|\| __NPU_ARCH__ == 3510
	32	+
	33	+__simt_callee__ __forceinline__ HcommJetty::HcommJetty() = default;
	34	+
	35	+__simt_callee__ __forceinline__
	36	+HcommJetty::HcommJetty(__ubuf__ HcommJettyInfo* info, __ubuf__ uint8_t* stage, __gm__ void* table, uint32_t index)
	37	+ : impl_(info, stage, table, index)
	38	+{}
	39	+
	40	+__simt_callee__ __forceinline__
	41	+HcommJetty::HcommJetty(__ubuf__ HcommJettyInfo* info, __ubuf__ uint8_t* stage, uint64_t channelAddr)
	42	+ : impl_(info, stage, channelAddr)
	43	+{}
	44	+
	45	+__simt_callee__ __forceinline__ __ubuf__ HcommJettyInfo* HcommJetty::Info() const { return impl_.Info(); }
	46	+
	47	+__simt_callee__ __forceinline__ void HcommJetty::Attach(
	48	+ __ubuf__ HcommJettyInfo* info, __ubuf__ uint8_t* stage, uint64_t channel)
	49	+{
	50	+ impl_.Attach(info, stage, channel);
	51	+}
	52	+
	53	+template <uint32_t kSgeNum, bool kDoCommit, typename Coop>
	54	+__simt_callee__ __forceinline__ int32_t
	55	+HcommJetty::Write(const HcommPeer& peer, __gm__ void* dst, const JettySge* sges, const Coop& coop) const
	56	+{
	57	+ return impl_.template Write<kSgeNum, kDoCommit>(peer, dst, sges, coop);
	58	+}
	59	+
	60	+template <uint32_t kSgeNum, bool kDoCommit, typename Coop>
	61	+__simt_callee__ __forceinline__ int32_t
	62	+HcommJetty::Write(const Coop& coop, const HcommPeer& peer, __gm__ void* dst, const HcommSge& sge) const
	63	+{
	64	+ return impl_.template Write<kSgeNum, kDoCommit>(peer, dst, &sge, coop);
	65	+}
	66	+
	67	+template <uint32_t kSgeNum, bool kDoCommit, typename Coop>
	68	+__simt_callee__ __forceinline__ int32_t
	69	+HcommJetty::Write(const Coop& coop, const HcommPeer& peer, __gm__ void* dst, const HcommSge* sges) const
	70	+{
	71	+ return impl_.template Write<kSgeNum, kDoCommit>(peer, dst, sges, coop);
	72	+}
	73	+
	74	+template <uint32_t kSgeNum, bool kDoCommit, typename Coop>
	75	+__simt_callee__ __forceinline__ int32_t
	76	+HcommJetty::Put(const HcommPeer& peer, __gm__ void* dst, const HcommSge* sges, const Coop& coop) const
	77	+{
	78	+ return Write<kSgeNum, kDoCommit>(peer, dst, sges, coop);
	79	+}
	80	+
	81	+template <uint32_t kSgeNum, bool kDoCommit, typename Coop>
	82	+__simt_callee__ __forceinline__ int32_t
	83	+HcommJetty::Put(const Coop& coop, const HcommPeer& peer, __gm__ void* dst, const HcommSge& sge) const
	84	+{
	85	+ return impl_.template Write<kSgeNum, kDoCommit>(peer, dst, &sge, coop);
	86	+}
	87	+
	88	+template <int64_t kTimeout, typename T, bool kDoCommit>
	89	+__simt_callee__ __forceinline__ int32_t HcommJetty::WriteValue(const HcommPeer& peer, __gm__ void* dst, T value)
	90	+{
	91	+ return impl_.template WriteValue<kTimeout, T, kDoCommit>(peer, dst, value);
	92	+}
	93	+
	94	+template <int64_t kTimeout, bool kDoCommit>
	95	+__simt_callee__ __forceinline__ int32_t HcommJetty::WriteWithNotify(
	96	+ const HcommPeer& peer, __gm__ void* dst, __gm__ void* src, uint64_t len, __gm__ void* notifyAddr,
	97	+ uint64_t notifyValue)
	98	+{
	99	+ return impl_.template WriteWithNotify<kTimeout, kDoCommit>(peer, dst, src, len, notifyAddr, notifyValue);
	100	+}
	101	+
	102	+__simt_callee__ __forceinline__ void HcommJetty::AdvanceSq() const { impl_.AdvanceSq(); }
	103	+
	104	+template <typename Coop>
	105	+__simt_callee__ __forceinline__ void HcommJetty::AdvanceSq(const Coop& coop) const
	106	+{
	107	+ impl_.AdvanceSq(coop);
	108	+}
	109	+
	110	+template <typename Coop>
	111	+__simt_callee__ __forceinline__ void HcommJetty::AdvanceSq(const Coop& coop, uint32_t a, uint32_t b) const
	112	+{
	113	+ impl_.AdvanceSq(coop, a, b);
	114	+}
	115	+
	116	+__simt_callee__ __forceinline__ void HcommJetty::PublishSq() const { impl_.PublishSq(); }
	117	+
	118	+template <typename Coop>
	119	+__simt_callee__ __forceinline__ void HcommJetty::PublishSq(const Coop& coop) const
	120	+{
	121	+ impl_.PublishSq(coop);
	122	+}
	123	+
	124	+template <int64_t kTimeout>
	125	+__simt_callee__ __forceinline__ int32_t HcommJetty::Drain() const
	126	+{
	127	+ return impl_.template Drain<kTimeout>();
	128	+}
	129	+
	130	+template <int64_t kTimeout, typename Coop>
	131	+__simt_callee__ __forceinline__ int32_t HcommJetty::Drain(const Coop& coop) const
	132	+{
	133	+ return impl_.template Drain<kTimeout>(coop);
	134	+}
	135	+
	136	+#endif // !defined(__NPU_ARCH__) \|\| __NPU_ARCH__ == 3510
	137	+
	138	+} // namespace AscendC::simt::jetty
	139	+
	140	+#endif // IMPL_ADV_API_DETAIL_HCOMM_IMPL_HCOMM_JETTY_SIMT_IMPL_H

Asrc/aicore/jetty/impl/hcomm_jetty_simt_impl_def.h+29-0

@@ -0,0 +1,29 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file hcomm_jetty_simt_impl_def.h
	13	+ * \brief Declarations the public SIMT Jetty header needs before class HcommJetty.
	14	+ */
	15	+
	16	+#ifndef IMPL_ADV_API_DETAIL_HCOMM_IMPL_HCOMM_JETTY_SIMT_IMPL_DEF_H
	17	+#define IMPL_ADV_API_DETAIL_HCOMM_IMPL_HCOMM_JETTY_SIMT_IMPL_DEF_H
	18	+
	19	+#include "../../hcomm/common/hcomm_simt_utils.h"
	20	+
	21	+// SIMT Jetty is only available on v310. The ASC driver compiles each translation unit twice:
	22	+// __NPU_ARCH__ is 3510 in the device pass and undefined in the host pass. The public HcommJetty
	23	+// holds a JettyImpl member, so the class must stay visible in both passes; gating on
	24	+// "not some other arch" keeps v310 and the host pass while still excluding other platforms.
	25	+#if !defined(__NPU_ARCH__) \|\| __NPU_ARCH__ == 3510
	26	+#include "platform_v310/hcomm_simt_urma_jetty_def.h"
	27	+#endif
	28	+
	29	+#endif // IMPL_ADV_API_DETAIL_HCOMM_IMPL_HCOMM_JETTY_SIMT_IMPL_DEF_H

Msrc/aicore/jetty/impl/platform_v310/hcomm_aiv_urma_jetty.h+28-23

	@@ -47,10 +47,12 @@ __aicore__ inline HcommPeer::HcommPeer(__ubuf__ HcommJettyPeerInfo* peerInfo, GM
47	"jetty info is nullptr\n");	47	"jetty info is nullptr\n");
48	// The last registered buffer carries the shared remote token metadata.	48	// The last registered buffer carries the shared remote token metadata.
49	auto* remoteBuffer = jetty->remoteBufferAddr + (jetty->remoteBufferNum - 1U);	49	auto* remoteBuffer = jetty->remoteBufferAddr + (jetty->remoteBufferNum - 1U);
50	- peerInfo->remoteTokenValue = remoteBuffer->bufferInfo.rma.protectionInfo.memInfo.ub.tokenValue;	50	+ peerInfo->remoteTokenValueUdf = remoteBuffer->bufferInfo.rma.protectionInfo.memInfo.ub.tokenValue;
51	- peerInfo->remoteTokenId = remoteBuffer->bufferInfo.rma.protectionInfo.memInfo.ub.tokenId;
52	auto* sqContext = jetty->sqContextAddr + HCOMM_URMA_DEFAULT_QP_IDX;	51	auto* sqContext = jetty->sqContextAddr + HCOMM_URMA_DEFAULT_QP_IDX;
53	- peerInfo->tpId = sqContext->contextInfo.ubJfs.tpID;	52	+ // tpId occupies bits [23:0] and remoteTokenId bits [51:32]; numSges is OR-ed in per post.
		53	+ peerInfo->tpIdNumSgesRemoteTokenId =
		54	+ static_cast<uint64_t>(sqContext->contextInfo.ubJfs.tpID) \|
		55	+ (static_cast<uint64_t>(remoteBuffer->bufferInfo.rma.protectionInfo.memInfo.ub.tokenId) << 32U);
54	auto* remoteEid = reinterpret_cast<const uint32_t*>(sqContext->contextInfo.ubJfs.remoteEID);	56	auto* remoteEid = reinterpret_cast<const uint32_t*>(sqContext->contextInfo.ubJfs.remoteEID);
55	peerInfo->remoteEid[0] = static_cast<uint64_t>(remoteEid[0]) \| (static_cast<uint64_t>(remoteEid[1]) << 32U);	57	peerInfo->remoteEid[0] = static_cast<uint64_t>(remoteEid[0]) \| (static_cast<uint64_t>(remoteEid[1]) << 32U);
56	peerInfo->remoteEid[1] = static_cast<uint64_t>(remoteEid[2]) \| (static_cast<uint64_t>(remoteEid[3]) << 32U);	58	peerInfo->remoteEid[1] = static_cast<uint64_t>(remoteEid[2]) \| (static_cast<uint64_t>(remoteEid[3]) << 32U);
	@@ -71,14 +73,12 @@ __aicore__ inline HcommJettyImpl::HcommJettyImpl(
71	auto* cqContext = jetty->cqContextAddr + HCOMM_URMA_DEFAULT_QP_IDX;	73	auto* cqContext = jetty->cqContextAddr + HCOMM_URMA_DEFAULT_QP_IDX;
72	jettyInfo_->sqBaseAddr = sqContext->contextInfo.ubJfs.sqVa;	74	jettyInfo_->sqBaseAddr = sqContext->contextInfo.ubJfs.sqVa;
73	jettyInfo_->sqHeadAddr = sqContext->contextInfo.ubJfs.headAddr;	75	jettyInfo_->sqHeadAddr = sqContext->contextInfo.ubJfs.headAddr;
74	- jettyInfo_->sqTailAddr = sqContext->contextInfo.ubJfs.tailAddr;	76	+ jettyInfo_->completionTailAddr = sqContext->contextInfo.ubJfs.tailAddr;
75	jettyInfo_->sqDoorbellAddr = sqContext->contextInfo.ubJfs.dbVa;	77	jettyInfo_->sqDoorbellAddr = sqContext->contextInfo.ubJfs.dbVa;
76	jettyInfo_->cqBaseAddr = cqContext->contextInfo.ubJfc.scqVa;	78	jettyInfo_->cqBaseAddr = cqContext->contextInfo.ubJfc.scqVa;
77	jettyInfo_->cqTailAddr = cqContext->contextInfo.ubJfc.tailAddr;	79	jettyInfo_->cqTailAddr = cqContext->contextInfo.ubJfc.tailAddr;
78	jettyInfo_->cqDoorbellAddr = cqContext->contextInfo.ubJfc.dbVa;	80	jettyInfo_->cqDoorbellAddr = cqContext->contextInfo.ubJfc.dbVa;
79	- uint64_t packedHead = ReadGmBypassDCache(reinterpret_cast<__gm__ uint64_t*>(jettyInfo_->sqHeadAddr));	81	+ jettyInfo_->packedHead = ReadGmBypassDCache(reinterpret_cast<__gm__ uint64_t*>(jettyInfo_->sqHeadAddr));
80	- jettyInfo_->sqHead = static_cast<uint32_t>(packedHead);
81	- jettyInfo_->expectedCqeCnt = static_cast<uint32_t>(packedHead >> 32U);
82	jettyInfo_->numWqebbBytes = sqContext->contextInfo.ubJfs.wqeSize;	82	jettyInfo_->numWqebbBytes = sqContext->contextInfo.ubJfs.wqeSize;
83	jettyInfo_->numCqeBytes = cqContext->contextInfo.ubJfc.cqeSize;	83	jettyInfo_->numCqeBytes = cqContext->contextInfo.ubJfc.cqeSize;
84	jettyInfo_->sqDepth = sqContext->contextInfo.ubJfs.sqDepth;	84	jettyInfo_->sqDepth = sqContext->contextInfo.ubJfs.sqDepth;
	@@ -91,17 +91,18 @@ __aicore__ inline void HcommJettyImpl::FillWriteSqeHeader(
91	const HcommPeer& peer, GM_ADDR dstAddr, uint32_t inlineMsgLen, uint32_t sgeNum)	91	const HcommPeer& peer, GM_ADDR dstAddr, uint32_t inlineMsgLen, uint32_t sgeNum)
92	{	92	{
93	auto* info = peer.peerInfo;	93	auto* info = peer.peerInfo;
94	- const uint32_t owner = (jettyInfo_->sqHead & jettyInfo_->sqDepth) == 0U ? 1U : 0U;	94	+ const uint32_t owner = (static_cast<uint32_t>(jettyInfo_->packedHead) & jettyInfo_->sqDepth) == 0U ? 1U : 0U;
95	// Fill HcommUrmaSqeCtx with one 64-bit store per word: word0/1 pack the bitfields, word5 is rmtAddr.	95	// Fill HcommUrmaSqeCtx with one 64-bit store per word: word0/1 pack the bitfields, word5 is rmtAddr.
96	__ubuf__ uint64_t* sqeWords = reinterpret_cast<__ubuf__ uint64_t*>(ubufSqe_);	96	__ubuf__ uint64_t* sqeWords = reinterpret_cast<__ubuf__ uint64_t*>(ubufSqe_);
97	sqeWords[0] = (static_cast<uint64_t>(HcommJettyPackSqeFlag(config) & 0xFFU) << 16U) \| (1U << 28U) \| (1U << 29U) \|	97	sqeWords[0] = (static_cast<uint64_t>(HcommJettyPackSqeFlag(config) & 0xFFU) << 16U) \| (1U << 28U) \| (1U << 29U) \|
98	(static_cast<uint64_t>(owner) << 31U) \| (static_cast<uint64_t>(opCode) << 40U) \|	98	(static_cast<uint64_t>(owner) << 31U) \| (static_cast<uint64_t>(opCode) << 40U) \|
99	(static_cast<uint64_t>(inlineMsgLen) << 54U);	99	(static_cast<uint64_t>(inlineMsgLen) << 54U);
100	- sqeWords[1] = (static_cast<uint64_t>(info->tpId) & 0xFFFFFFU) \| (static_cast<uint64_t>(sgeNum) << 24U) \|	100	+ // tpId and remoteTokenId already share one word in the peer descriptor, so a single load
101	- ((static_cast<uint64_t>(info->remoteTokenId) & 0xFFFFFU) << 32U);	101	+ // covers both; numSges is not stored and is OR-ed in here.
		102	+ sqeWords[1] = info->tpIdNumSgesRemoteTokenId \| (static_cast<uint64_t>(sgeNum) << 24U);
102	sqeWords[2] = info->remoteEid[0];	103	sqeWords[2] = info->remoteEid[0];
103	sqeWords[3] = info->remoteEid[1];	104	sqeWords[3] = info->remoteEid[1];
104	- sqeWords[4] = static_cast<uint64_t>(info->remoteTokenValue);	105	+ sqeWords[4] = info->remoteTokenValueUdf;
105	sqeWords[5] = reinterpret_cast<uint64_t>(dstAddr);	106	sqeWords[5] = reinterpret_cast<uint64_t>(dstAddr);
106	}	107	}
107		108
	@@ -115,8 +116,8 @@ __aicore__ inline void HcommJettyImpl::BuildWriteSqe(
115	__ubuf__ uint64_t* sgeWords = reinterpret_cast<__ubuf__ uint64_t*>(ubufSqe_ + sizeof(HcommUrmaSqeCtx));	116	__ubuf__ uint64_t* sgeWords = reinterpret_cast<__ubuf__ uint64_t*>(ubufSqe_ + sizeof(HcommUrmaSqeCtx));
116	if constexpr (withNotify) {	117	if constexpr (withNotify) {
117	// Fill HcommUrmaNotifyCtx with one 64-bit store per word: word1 is notifyAddr, word2 is notifyData.	118	// Fill HcommUrmaNotifyCtx with one 64-bit store per word: word1 is notifyAddr, word2 is notifyData.
118	- sgeWords[0] = (static_cast<uint64_t>(peer.peerInfo->remoteTokenId) & 0xFFFFFU) \|	119	+ sgeWords[0] = ((peer.peerInfo->tpIdNumSgesRemoteTokenId >> 32U) & 0xFFFFFU) \|
119	- (static_cast<uint64_t>(peer.peerInfo->remoteTokenValue) << 32U);	120	+ ((peer.peerInfo->remoteTokenValueUdf & 0xFFFFFFFFULL) << 32U);
120	sgeWords[1] = reinterpret_cast<uint64_t>(notifyAddr);	121	sgeWords[1] = reinterpret_cast<uint64_t>(notifyAddr);
121	sgeWords[2] = notifyValue;	122	sgeWords[2] = notifyValue;
122	sgeWords[3] = 0U;	123	sgeWords[3] = 0U;
	@@ -155,17 +156,21 @@ __aicore__ inline void HcommJettyImpl::CopyWqeToSq(uint32_t currentHead, uint32_
155		156
156	__aicore__ inline void HcommJettyImpl::AdvanceSq(uint32_t numWqebbs)	157	__aicore__ inline void HcommJettyImpl::AdvanceSq(uint32_t numWqebbs)
157	{	158	{
158	- jettyInfo_->sqHead += numWqebbs;	159	+ // packedHead keeps the head in its low 32 bits and the expected CQE count in the high 32.
159	- ++jettyInfo_->expectedCqeCnt;	160	+ // The two counters wrap independently, so they cannot be advanced with one 64-bit add: a
160	- uint64_t packedHead =	161	+ // head wrapping past 0xFFFFFFFF would carry into the CQE count.
161	- static_cast<uint64_t>(jettyInfo_->sqHead) \| (static_cast<uint64_t>(jettyInfo_->expectedCqeCnt) << 32U);	162	+ const uint32_t head = static_cast<uint32_t>(jettyInfo_->packedHead) + numWqebbs;
162	- WriteGmBypassDCache(reinterpret_cast<__gm__ uint64_t*>(jettyInfo_->sqHeadAddr), packedHead);	163	+ const uint32_t expectedCqeCnt = static_cast<uint32_t>(jettyInfo_->packedHead >> 32U) + 1U;
		164	+ jettyInfo_->packedHead = static_cast<uint64_t>(head) \| (static_cast<uint64_t>(expectedCqeCnt) << 32U);
		165	+ WriteGmBypassDCache(reinterpret_cast<__gm__ uint64_t*>(jettyInfo_->sqHeadAddr), jettyInfo_->packedHead);
163	AscendC::DataSyncBarrier<AscendC::MemDsbT::DDR>();	166	AscendC::DataSyncBarrier<AscendC::MemDsbT::DDR>();
164	}	167	}
165		168
166	__aicore__ inline void HcommJettyImpl::RingDoorbell()	169	__aicore__ inline void HcommJettyImpl::RingDoorbell()
167	{	170	{
168	- WriteGmBypassDCache(reinterpret_cast<__gm__ uint32_t*>(jettyInfo_->sqDoorbellAddr), jettyInfo_->sqHead);	171	+ // The doorbell only takes the head, which lives in the low 32 bits of packedHead.
		172	+ WriteGmBypassDCache(
		173	+ reinterpret_cast<__gm__ uint32_t*>(jettyInfo_->sqDoorbellAddr), static_cast<uint32_t>(jettyInfo_->packedHead));
169	}	174	}
170		175
171	template <int64_t timeoutCycles>	176	template <int64_t timeoutCycles>
	@@ -176,8 +181,8 @@ __aicore__ inline void HcommJettyImpl::PollCqWhenSqOverflow()
176	// distance can jump past an exact value and an equality test would never fire.	181	// distance can jump past an exact value and an equality test would never fire.
177	constexpr uint32_t nearFullThreshold = 10U;	182	constexpr uint32_t nearFullThreshold = 10U;
178	constexpr uint32_t reclaimBatch = 100U;	183	constexpr uint32_t reclaimBatch = 100U;
179	- uint32_t completions = jettyInfo_->expectedCqeCnt;	184	+ uint32_t completions = static_cast<uint32_t>(jettyInfo_->packedHead >> 32U);
180	- uint32_t tail = ReadGmBypassDCache(reinterpret_cast<__gm__ uint32_t*>(jettyInfo_->sqTailAddr));	185	+ uint32_t tail = ReadGmBypassDCache(reinterpret_cast<__gm__ uint32_t*>(jettyInfo_->completionTailAddr));
181	if (completions - tail + nearFullThreshold >= jettyInfo_->cqDepth) {	186	if (completions - tail + nearFullThreshold >= jettyInfo_->cqDepth) {
182	PollCq<timeoutCycles>(tail + reclaimBatch < completions ? tail + reclaimBatch : completions);	187	PollCq<timeoutCycles>(tail + reclaimBatch < completions ? tail + reclaimBatch : completions);
183	}	188	}
	@@ -216,7 +221,7 @@ __aicore__ inline int32_t HcommJettyImpl::PollCq(uint32_t expectedTail)
216	WriteGmBypassDCache(reinterpret_cast<__gm__ uint32_t*>(jettyInfo_->cqTailAddr), currentTail);	221	WriteGmBypassDCache(reinterpret_cast<__gm__ uint32_t*>(jettyInfo_->cqTailAddr), currentTail);
217	WriteGmBypassDCache(	222	WriteGmBypassDCache(
218	reinterpret_cast<__gm__ uint32_t*>(jettyInfo_->cqDoorbellAddr), currentTail & HCOMM_JETTY_CQ_DOORBELL_MASK);	223	reinterpret_cast<__gm__ uint32_t*>(jettyInfo_->cqDoorbellAddr), currentTail & HCOMM_JETTY_CQ_DOORBELL_MASK);
219	- WriteGmBypassDCache(reinterpret_cast<__gm__ uint32_t*>(jettyInfo_->sqTailAddr), currentTail);	224	+ WriteGmBypassDCache(reinterpret_cast<__gm__ uint32_t*>(jettyInfo_->completionTailAddr), currentTail);
220	AscendC::DataSyncBarrier<AscendC::MemDsbT::DDR>();	225	AscendC::DataSyncBarrier<AscendC::MemDsbT::DDR>();
221	return HCOMM_SUCCESS;	226	return HCOMM_SUCCESS;
222	}	227	}
	@@ -230,7 +235,7 @@ __aicore__ inline bool HcommJettyImpl::Ready()
230	template <int64_t timeoutCycles, bool doCommit>	235	template <int64_t timeoutCycles, bool doCommit>
231	__aicore__ inline int32_t HcommJettyImpl::PostSqe(uint32_t numWqebbs)	236	__aicore__ inline int32_t HcommJettyImpl::PostSqe(uint32_t numWqebbs)
232	{	237	{
233	- uint32_t head = jettyInfo_->sqHead;	238	+ uint32_t head = static_cast<uint32_t>(jettyInfo_->packedHead);
234	PollCqWhenSqOverflow<timeoutCycles>();	239	PollCqWhenSqOverflow<timeoutCycles>();
235	AdvanceSq(numWqebbs);	240	AdvanceSq(numWqebbs);
236	CopyWqeToSq(head, numWqebbs);	241	CopyWqeToSq(head, numWqebbs);

Asrc/aicore/jetty/impl/platform_v310/hcomm_simt_urma_jetty.h+730-0

Asrc/aicore/jetty/impl/platform_v310/hcomm_simt_urma_jetty_def.h+209-0

@@ -0,0 +1,209 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file hcomm_simt_urma_jetty_def.h
	13	+ * \brief JettyImpl class skeleton shared by the public header and the v310 implementation.
	14	+ */
	15	+
	16	+#ifndef IMPL_ADV_API_DETAIL_HCOMM_IMPL_PLATFORM_V310_HCOMM_SIMT_URMA_JETTY_DEF_H
	17	+#define IMPL_ADV_API_DETAIL_HCOMM_IMPL_PLATFORM_V310_HCOMM_SIMT_URMA_JETTY_DEF_H
	18	+
	19	+#include <cstdint>
	20	+
	21	+#include "../../../hcomm/impl/platform_v310/hcomm_simt_urma_def.h"
	22	+#include "../../../utils/simt_utils.h"
	23	+
	24	+namespace AscendC::simt {
	25	+namespace jetty {
	26	+
	27	+// Geometry of the SIMT SQ layout: one WQEBB is 64B, a warp-cooperative WRITE
	28	+// reserves a fixed four-WQEBB slot for up to twelve SGEs.
	29	+constexpr uint32_t HCOMM_SIMT_WQEBB_LANES = HCOMM_URMA_WQE_BB_SIZE / sizeof(uint64_t);
	30	+constexpr uint32_t HCOMM_SIMT_WARP_SIZE = 32U;
	31	+constexpr uint32_t HCOMM_SIMT_MAX_SGES = 12U;
	32	+constexpr uint32_t HCOMM_SIMT_SLOT_WQEBBS = 4U;
	33	+
	34	+struct JettySge {
	35	+ uint64_t addr = 0U;
	36	+ int len = 0;
	37	+};
	38	+
	39	+enum class HcommPostOpcode : uint8_t {
	40	+ kWrite,
	41	+ kWriteValue,
	42	+ kWriteWithNotify,
	43	+};
	44	+
	45	+namespace detail {
	46	+template <bool kWithNotify = false>
	47	+struct PostSendArgs {
	48	+ __gm__ void* dst;
	49	+ const JettySge* sges;
	50	+};
	51	+
	52	+template <>
	53	+struct PostSendArgs<true> {
	54	+ __gm__ void* dst;
	55	+ const JettySge* sges;
	56	+ __gm__ void* notifyAddr;
	57	+ uint64_t notifyValue;
	58	+};
	59	+
	60	+struct InlinePostSendArgs {
	61	+ __gm__ void* dst;
	62	+ const uint8_t* value;
	63	+};
	64	+} // namespace detail
	65	+
	66	+class HcommPeer;
	67	+
	68	+class JettyImpl {
	69	+public:
	70	+ static constexpr uint32_t kNumWQEBBLanes = HCOMM_SIMT_WQEBB_LANES;
	71	+
	72	+ static constexpr uint32_t kNumWQEBBsPerSQESlot = HCOMM_SIMT_SLOT_WQEBBS;
	73	+
	74	+ static constexpr uint32_t kNumMaxSGEsPerWrite = HCOMM_SIMT_MAX_SGES;
	75	+
	76	+ static constexpr uint32_t kNumWriteWqebbs = HCOMM_SIMT_SLOT_WQEBBS;
	77	+
	78	+ static constexpr uint32_t kNumWriteWithNotifyWqebbs = 2U;
	79	+
	80	+ static constexpr uint32_t kNumMaxSQEBytes = HCOMM_SIMT_SLOT_WQEBBS * HCOMM_URMA_WQE_BB_SIZE;
	81	+
	82	+ static constexpr uint32_t kNumWQEBBsPerWrite = HCOMM_SIMT_SLOT_WQEBBS;
	83	+
	84	+ static constexpr uint32_t kNumDwqeWords = 2U * HCOMM_SIMT_WQEBB_LANES;
	85	+
	86	+ __simt_callee__ __forceinline__ JettyImpl() = default;
	87	+
	88	+ __simt_callee__ __forceinline__
	89	+ JettyImpl(__ubuf__ HcommJettyInfo* info, __ubuf__ uint8_t* stage, __gm__ void* table, uint32_t index);
	90	+
	91	+ __simt_callee__ __forceinline__
	92	+ JettyImpl(__ubuf__ HcommJettyInfo* info, __ubuf__ uint8_t* stage, uint64_t channelAddr);
	93	+
	94	+ __simt_callee__ __forceinline__ __ubuf__ HcommJettyInfo* Info() const;
	95	+
	96	+ template <uint32_t kSgeNum, bool kDoCommit = true, typename Coop = HcommCoopThread>
	97	+ __simt_callee__ __forceinline__ int32_t
	98	+ Write(const HcommPeer& peer, __gm__ void* dst, const JettySge* sges, const Coop& coop = Coop{}) const;
	99	+
	100	+ template <int64_t, typename T, bool kDoCommit = true>
	101	+ __simt_callee__ __forceinline__ int32_t WriteValue(const HcommPeer& peer, __gm__ void* dst, T value);
	102	+
	103	+ template <int64_t, bool kDoCommit = true>
	104	+ __simt_callee__ __forceinline__ int32_t WriteWithNotify(
	105	+ const HcommPeer& peer, __gm__ void* dst, __gm__ void* src, uint64_t len, __gm__ void* notifyAddr,
	106	+ uint64_t notifyValue);
	107	+
	108	+ __simt_callee__ __forceinline__ void AdvanceSq() const;
	109	+
	110	+ template <typename Coop>
	111	+ __simt_callee__ __forceinline__ void AdvanceSq(const Coop& coop) const;
	112	+
	113	+ template <typename Coop>
	114	+ __simt_callee__ __forceinline__ void AdvanceSq(const Coop& coop, uint32_t, uint32_t) const;
	115	+
	116	+ __simt_callee__ __forceinline__ void AdvanceImplicitSq() const;
	117	+
	118	+ __simt_callee__ __forceinline__ void PublishSq() const;
	119	+
	120	+ template <typename Coop>
	121	+ __simt_callee__ __forceinline__ void PublishSq(const Coop& coop) const;
	122	+
	123	+ template <int64_t kTimeout>
	124	+ __simt_callee__ __forceinline__ int32_t Drain() const;
	125	+
	126	+ template <int64_t kTimeout, typename Coop>
	127	+ __simt_callee__ __forceinline__ int32_t Drain(const Coop& coop) const;
	128	+
	129	+ __simt_callee__ __forceinline__ void Attach(
	130	+ __ubuf__ HcommJettyInfo* info, __ubuf__ uint8_t* stage, uint64_t channelAddr);
	131	+
	132	+private:
	133	+ struct BatchState {
	134	+ uint32_t wqebbCount;
	135	+ uint32_t lastOffset;
	136	+ uint32_t failed;
	137	+ };
	138	+
	139	+ __simt_callee__ __forceinline__ void Init(
	140	+ __ubuf__ HcommJettyInfo* info, __ubuf__ uint8_t* stage, __gm__ void* table, uint32_t index);
	141	+
	142	+ __simt_callee__ __forceinline__ __ubuf__ BatchState* Batch() const;
	143	+
	144	+ __simt_callee__ __forceinline__ uint32_t ReserveImplicitSq(const uint32_t reserveWqebbs) const;
	145	+
	146	+ __simt_callee__ __forceinline__ uint32_t ReserveBatchSq(const uint32_t reserveWqebbs) const;
	147	+
	148	+ template <
	149	+ uint32_t kSgeNum, HcommPostOpcode kOpcode, bool kDoCommit, uint32_t kInlineBytes = 0U, typename Coop,
	150	+ typename Args>
	151	+ __simt_callee__ __forceinline__ int32_t PostSend(const Coop& coop, const HcommPeer& peer, const Args& args) const;
	152	+
	153	+ __simt_callee__ __forceinline__ void AdvanceBatch() const;
	154	+
	155	+ uint64_t infoAddr_ = 0U;
	156	+
	157	+ uint64_t stageAddr_ = 0U;
	158	+
	159	+ uint64_t channel_ = 0U;
	160	+
	161	+ mutable uint32_t implicitBeginHead_ = 0U;
	162	+
	163	+ mutable uint32_t implicitWqebbs_ = 0U;
	164	+
	165	+ mutable uint32_t implicitPayloadWqebbs_ = 0U;
	166	+
	167	+ mutable bool implicitActive_ = false;
	168	+
	169	+ static constexpr int32_t kPollCqSuccess = 0;
	170	+
	171	+ static constexpr int32_t kPollCqUnavailable = 1;
	172	+
	173	+ // Poll one CQE and advance the local SQ completion tail. The caller owns CQ tail and
	174	+ // doorbell updates so this helper can serve both Drain and the incremental reservation path.
	175	+ template <int64_t kTimeout = HCOMM_SIMT_MAX_CQ_RETRY>
	176	+ __simt_callee__ __forceinline__ int32_t PollCqEntry(const uint32_t current, uint32_t& released) const;
	177	+
	178	+ // Reap one completed CQE, matching the SIMD PullCq behavior. This is used only by
	179	+ // the single-producer path, where no other producer can race the CQ tail update.
	180	+ __simt_callee__ __forceinline__ bool PullCqOne() const;
	181	+
	182	+ __simt_callee__ __forceinline__ static bool IsValidSge(const JettySge& sge);
	183	+
	184	+ template <uint32_t kSgeNum>
	185	+ __simt_callee__ __forceinline__ static uint32_t CountScalarSges(const JettySge* sges);
	186	+
	187	+ template <
	188	+ uint32_t kSgeNum, HcommPostOpcode kOpcode, uint32_t kInlineBytes, uint32_t kWqebbCount, typename Coop,
	189	+ typename Args>
	190	+ __simt_callee__ __forceinline__ void FillSqWarpLane(
	191	+ const Coop& coop, const HcommPeer& peer, const Args& args, const uint32_t head,
	192	+ const uint32_t runtimeWqebbCount, const uint32_t scalarSgeNum) const;
	193	+
	194	+ template <uint32_t kPayloadWords, uint32_t kInlineBytes>
	195	+ __simt_callee__ __forceinline__ static uint64_t BuildValueWord(
	196	+ const HcommPeer& peer, const detail::InlinePostSendArgs& args, const uint32_t word, const uint32_t sqIdx);
	197	+
	198	+ __simt_callee__ __forceinline__ static uint64_t BuildNotifyWord(
	199	+ const HcommPeer& peer, const detail::PostSendArgs<true>& args, const uint32_t word, const uint32_t sqIdx);
	200	+
	201	+ __simt_callee__ __forceinline__ static void StoreSqePair(
	202	+ __gm__ uint64_t* sq, const uint32_t depth, const uint32_t head, const uint32_t firstWord, const uint64_t first,
	203	+ const uint64_t second);
	204	+};
	205	+
	206	+} // namespace jetty
	207	+} // namespace AscendC::simt
	208	+
	209	+#endif // IMPL_ADV_API_DETAIL_HCOMM_IMPL_PLATFORM_V310_HCOMM_SIMT_URMA_JETTY_DEF_H

Asrc/aicore/utils/simt_utils.h+108-0

@@ -0,0 +1,108 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file simt_utils.h
	13	+ * \brief SIMT warp primitives and cooperative-group helpers shared by the SIMT paths.
	14	+ */
	15	+#ifndef SRC_AICORE_UTILS_SIMT_UTILS_H
	16	+#define SRC_AICORE_UTILS_SIMT_UTILS_H
	17	+
	18	+#include <cstdint>
	19	+#include <simt_api/device_warp_functions.h>
	20	+#include <simt_api/cooperative_groups.h>
	21	+
	22	+#ifdef __CCE__
	wanghan9月15日为什么要加这个条件编译？ suqwe 9月20日评论： 9月20日评论：
	23	+namespace AscendC::simt {
	24	+
	25	+template <uint32_t kNumCycles = 300>
	26	+__simt_callee__ __forceinline__ void nop()
	27	+{
	28	+ static_assert(kNumCycles % 15 == 0, "SIMT nop cycles must be a multiple of 15");
	29	+#pragma unroll
	30	+ for (int i = 0; i < kNumCycles / 15; ++i)
	31	+ asm volatile("NOP wait:0b0000000 stall:15" ::);
	32	+}
	33	+
	34	+__simt_callee__ __forceinline__ int exchange(const int& value, const int& lane_idx)
	35	+{
	36	+ return asc_shfl(value, lane_idx);
	37	+}
	38	+
	39	+template <int kMaxValue>
	40	+__simt_callee__ __forceinline__ bool deduplicate(const int& value, const int& lane_idx)
	41	+{
	42	+ auto peers = 0xffffffffu;
	43	+#pragma unroll
	44	+ for (int i = 0; (1 << i) <= kMaxValue; ++i) {
	45	+ const auto current_bit = (static_cast<unsigned>(value) >> i) & 1u;
	46	+ const auto ones = asc_ballot(current_bit);
	47	+ peers &= ones ^ (current_bit - 1u);
	48	+ }
	49	+ return (peers >> lane_idx) == 1u;
	50	+}
	51	+
	52	+__simt_callee__ __forceinline__ int warp_inclusive_sum(int value, const int& lane_idx)
	53	+{
	54	+#pragma unroll
	55	+ for (int offset = 1; offset < warpSize; offset <<= 1) {
	56	+ const auto synced = asc_shfl_up(value, offset);
	57	+ if (lane_idx >= offset)
	58	+ value += synced;
	59	+ }
	60	+ return value;
	61	+}
	62	+
	63	+__simt_callee__ __forceinline__ int warp_exclusive_sum(const int& value, const int& lane_idx)
	64	+{
	65	+ return warp_inclusive_sum(value, lane_idx) - value;
	66	+}
	67	+
	68	+struct HcommCoopThread {
	69	+ __simt_callee__ __forceinline__ void sync() const {}
	70	+ __simt_callee__ __forceinline__ int size() const { return 1; }
	71	+ __simt_callee__ __forceinline__ int thread_rank() const { return 0; }
	72	+};
	73	+struct HcommCoopWarp {
	74	+ cooperative_groups::thread_block_tile<32, cooperative_groups::thread_block> group_;
	75	+ __simt_callee__ __forceinline__ HcommCoopWarp()
	76	+ : group_(cooperative_groups::tiled_partition<32>(cooperative_groups::this_thread_block()))
	77	+ {}
	78	+ __simt_callee__ __forceinline__ void sync() const { group_.sync(); }
	79	+ __simt_callee__ __forceinline__ int size() const { return 32; }
	80	+ __simt_callee__ __forceinline__ int thread_rank() const { return static_cast<int>(group_.thread_rank()); }
	81	+};
	82	+struct HcommCoopGroup {
	83	+ cooperative_groups::thread_block group_;
	84	+ __simt_callee__ __forceinline__ HcommCoopGroup() : group_(cooperative_groups::this_thread_block()) {}
	85	+ __simt_callee__ __forceinline__ void sync() const { group_.sync(); }
	86	+ __simt_callee__ __forceinline__ int size() const { return static_cast<int>(group_.size()); }
	87	+ __simt_callee__ __forceinline__ int thread_rank() const { return static_cast<int>(group_.thread_rank()); }
	88	+};
	89	+template <typename Coop>
	90	+struct HcommCoopTraits {
	91	+ static constexpr bool kWarpCapable = false;
	92	+ static constexpr bool kIsWarp = false;
	93	+};
	94	+template <>
	95	+struct HcommCoopTraits<HcommCoopWarp> {
	96	+ static constexpr bool kWarpCapable = true;
	97	+ static constexpr bool kIsWarp = true;
	98	+};
	99	+template <>
	100	+struct HcommCoopTraits<HcommCoopGroup> {
	101	+ static constexpr bool kWarpCapable = true;
	102	+ static constexpr bool kIsWarp = false;
	103	+};
	104	+
	105	+} // namespace AscendC::simt
	106	+#endif
	107	+
	108	+#endif // SRC_AICORE_UTILS_SIMT_UTILS_H

Mtests/ut/CMakeLists.txt+3-0

	@@ -165,6 +165,9 @@ add_custom_command(
165	COMMAND ${CMAKE_COMMAND} -E copy_directory	165	COMMAND ${CMAKE_COMMAND} -E copy_directory
166	"${ASCCOMM_ROOT}/src/aicore/jetty"	166	"${ASCCOMM_ROOT}/src/aicore/jetty"
167	"${ASCCOMM_UT_STAGE_COMM_API_IMPL_DIR}/jetty"	167	"${ASCCOMM_UT_STAGE_COMM_API_IMPL_DIR}/jetty"
		168	+ COMMAND ${CMAKE_COMMAND} -E copy_directory
		169	+ "${ASCCOMM_ROOT}/src/aicore/utils"
		170	+ "${ASCCOMM_UT_STAGE_COMM_API_IMPL_DIR}/utils"
168	COMMAND ${CMAKE_COMMAND} -E touch ${ASCCOMM_UT_STAGE_STAMP}	171	COMMAND ${CMAKE_COMMAND} -E touch ${ASCCOMM_UT_STAGE_STAMP}
169	DEPENDS ${ASCCOMM_UT_LAYOUT_INPUTS}	172	DEPENDS ${ASCCOMM_UT_LAYOUT_INPUTS}
170	COMMENT "Preparing shared Ascend package-shaped headers for UT"	173	COMMENT "Preparing shared Ascend package-shaped headers for UT"

Mtests/ut/aicore/hcomm/test_hcomm_jetty_950.cpp+10-9

	@@ -109,9 +109,10 @@ TEST_F(HcommJettyTestSuite, Aiv_Jetty_PeerResolvesRemoteMetadata)
109	HcommPeer peer(&peerInfo_, channel_.GetJettyTable(), 0U);	109	HcommPeer peer(&peerInfo_, channel_.GetJettyTable(), 0U);
110		110
111	EXPECT_TRUE(peer.valid);	111	EXPECT_TRUE(peer.valid);
112	- EXPECT_EQ(peerInfo_.tpId, 77U);	112	+ // tpId occupies bits [23:0] and remoteTokenId bits [51:32] of the packed word.
113	- EXPECT_EQ(peerInfo_.remoteTokenId, 0x123456U);	113	+ EXPECT_EQ(peerInfo_.tpIdNumSgesRemoteTokenId & 0xFFFFFFU, 77U);
114	- EXPECT_EQ(peerInfo_.remoteTokenValue, 0x654321U);	114	+ EXPECT_EQ((peerInfo_.tpIdNumSgesRemoteTokenId >> 32U) & 0xFFFFFU, 0x123456U & 0xFFFFFU);
		115	+ EXPECT_EQ(peerInfo_.remoteTokenValueUdf, 0x654321U);
115	EXPECT_EQ(peerInfo_.remoteEid[0], 0x5566778811223344ULL);	116	EXPECT_EQ(peerInfo_.remoteEid[0], 0x5566778811223344ULL);
116	EXPECT_EQ(peerInfo_.remoteEid[1], 0xDDEEFF0099AABBCCULL);	117	EXPECT_EQ(peerInfo_.remoteEid[1], 0xDDEEFF0099AABBCCULL);
117	}	118	}
	@@ -121,8 +122,8 @@ TEST_F(HcommJettyTestSuite, Aiv_Jetty_InfoSplitsPackedHead)
121	channel_.SetPackedHead(5U, 9U);	122	channel_.SetPackedHead(5U, 9U);
122	HcommJetty jetty = MakeJetty();	123	HcommJetty jetty = MakeJetty();
123		124
124	- EXPECT_EQ(jettyInfo_.sqHead, 5U);	125	+ EXPECT_EQ(static_cast<uint32_t>(jettyInfo_.packedHead), 5U);
125	- EXPECT_EQ(jettyInfo_.expectedCqeCnt, 9U);	126	+ EXPECT_EQ(static_cast<uint32_t>(jettyInfo_.packedHead >> 32U), 9U);
126	EXPECT_EQ(jettyInfo_.sqDepth, JETTY_SQ_DEPTH);	127	EXPECT_EQ(jettyInfo_.sqDepth, JETTY_SQ_DEPTH);
127	EXPECT_EQ(jettyInfo_.numWqebbBytes, JETTY_WQE_SIZE);	128	EXPECT_EQ(jettyInfo_.numWqebbBytes, JETTY_WQE_SIZE);
128	EXPECT_EQ(jettyInfo_.cqDepth, JETTY_SQ_DEPTH);	129	EXPECT_EQ(jettyInfo_.cqDepth, JETTY_SQ_DEPTH);
	@@ -139,8 +140,8 @@ TEST_F(HcommJettyTestSuite, Aiv_Jetty_AdvanceSqAndRingDoorbell)
139	HcommJetty jetty = MakeJetty();	140	HcommJetty jetty = MakeJetty();
140		141
141	jetty.AdvanceSq(2U);	142	jetty.AdvanceSq(2U);
142	- EXPECT_EQ(jettyInfo_.sqHead, 5U);	143	+ EXPECT_EQ(static_cast<uint32_t>(jettyInfo_.packedHead), 5U);
143	- EXPECT_EQ(jettyInfo_.expectedCqeCnt, 8U);	144	+ EXPECT_EQ(static_cast<uint32_t>(jettyInfo_.packedHead >> 32U), 8U);
144	EXPECT_EQ(channel_.GetPackedHead(), (5U \| (8ULL << 32U)));	145	EXPECT_EQ(channel_.GetPackedHead(), (5U \| (8ULL << 32U)));
145		146
146	jetty.RingDoorbell();	147	jetty.RingDoorbell();
	@@ -153,8 +154,8 @@ TEST_F(HcommJettyTestSuite, Aiv_Jetty_AdvanceSqWrapsHead)
153	HcommJetty jetty = MakeJetty();	154	HcommJetty jetty = MakeJetty();
154		155
155	jetty.AdvanceSq(1U);	156	jetty.AdvanceSq(1U);
156	- EXPECT_EQ(jettyInfo_.sqHead, 0U);	157	+ EXPECT_EQ(static_cast<uint32_t>(jettyInfo_.packedHead), 0U);
157	- EXPECT_EQ(jettyInfo_.expectedCqeCnt, 2U);	158	+ EXPECT_EQ(static_cast<uint32_t>(jettyInfo_.packedHead >> 32U), 2U);
158	EXPECT_EQ(channel_.GetPackedHead(), (0U \| (2ULL << 32U)));	159	EXPECT_EQ(channel_.GetPackedHead(), (0U \| (2ULL << 32U)));
159	}	160	}
160		161