Star242
97
代码
代码
Issues49
Pull Requests79
Actions
项目讨论
Wiki
分析
Star242
97
已合并
统一样例 #173
wanghan创建于 9月12日
统一样例 #173
已合并
wanghan创建于 9月12日
从mizuki_p/asc-commfix_examples合入到cann/asc-commmaster
共 75 个文件变更+3240-3799
Mdocs/doc_contributing.md+1-1
@@ -112,7 +112,7 @@ docs/zh/guide/ ──首次引入API──→ docs/zh/api/README.md
112 ──涉及样例────→ examples/README.md112 ──涉及样例────→ examples/README.md
113 113 
114docs/zh/api/ ───使用流程──→ docs/zh/guide/hcomm_usage.md114docs/zh/api/ ───使用流程──→ docs/zh/guide/hcomm_usage.md
115- ───调用示例──→ examples/hcomm_write_read_nbi/README.md115+ ───调用示例──→ examples/aicore/hcomm/01_hcomm_write_read_nbi/README.md
116```116```
117 117 
118## 更多信息118## 更多信息
Mdocs/doc_contributing_en.md+1-1
@@ -112,7 +112,7 @@ docs/en/guide/ ──First API reference──→ docs/en/api/README.md
112 ──Sample reference──→ examples/README_en.md112 ──Sample reference──→ examples/README_en.md
113 113 
114docs/en/api/ ──Usage workflow──→ docs/en/guide/hcomm_usage.md114docs/en/api/ ──Usage workflow──→ docs/en/guide/hcomm_usage.md
115- ──Invocation sample──→ examples/hcomm_write_read_nbi/README_en.md115+ ──Invocation sample──→ examples/aicore/hcomm/01_hcomm_write_read_nbi/README_en.md
116```116```
117 117 
118## More Information118## More Information
Mdocs/en/api/README.md+2-2
@@ -69,5 +69,5 @@ AICore side:
69## Related Documents69## Related Documents
70 70 
71- [Hcomm Usage Guide](../guide/hcomm_usage.md)71- [Hcomm Usage Guide](../guide/hcomm_usage.md)
72-- [AIV Direct-driven URMA WriteNbi/ReadNbi Sample](../../../examples/hcomm_write_read_nbi/README_en.md)72+- [AIV Direct-driven URMA WriteNbi/ReadNbi Sample](../../../examples/aicore/hcomm/01_hcomm_write_read_nbi/README_en.md)
73-- [Ain Basic Ring Sample](../../../examples/ain/basic_ring/README_en.md)73+- [Ain Basic Ring Sample](../../../examples/aicore/ain/01_basic_ring/README_en.md)
Mdocs/en/api/aicore/hcomm/Hcomm.md+1-1
@@ -82,4 +82,4 @@ See [MakeMultiChannelHandle](../../host/hcomm/MakeMultiChannelHandle.md) for the
82 82 
83## Related Sample83## Related Sample
84 84 
85-[hcomm_write_read_nbi](../../../../../examples/hcomm_write_read_nbi/README_en.md) demonstrates ordinary `WriteNbi` and `ReadNbi` calls over the `COMM_PROTOCOL_UBC_CTP` path in a two-card scenario. It does not cover the RoCE or BatchHandle workflow.85+[hcomm_write_read_nbi](../../../../../examples/aicore/hcomm/01_hcomm_write_read_nbi/README_en.md) demonstrates ordinary `WriteNbi` and `ReadNbi` calls over the `COMM_PROTOCOL_UBC_CTP` path in a two-card scenario. It does not cover the RoCE or BatchHandle workflow.
Mdocs/en/api/aicore/hcomm/ReadNbi.md+1-1
@@ -73,6 +73,6 @@ __aicore__ inline int32_t ReadNbi(
73 73 
74## Related Sample74## Related Sample
75 75 
76-[hcomm_write_read_nbi](../../../../../examples/hcomm_write_read_nbi/README_en.md) demonstrates the ordinary AIV direct-driven URMA read/write workflow. See the [Hcomm Usage Guide](../../../guide/hcomm_usage.md) for the basic BatchHandle workflow.76+[hcomm_write_read_nbi](../../../../../examples/aicore/hcomm/01_hcomm_write_read_nbi/README_en.md) demonstrates the ordinary AIV direct-driven URMA read/write workflow. See the [Hcomm Usage Guide](../../../guide/hcomm_usage.md) for the basic BatchHandle workflow.
77 77 
78See [GetHandleRef](./GetHandleRef.md) for the multi-channel interface.78See [GetHandleRef](./GetHandleRef.md) for the multi-channel interface.
Mdocs/en/api/aicore/hcomm/WriteNbi.md+1-1
@@ -73,6 +73,6 @@ __aicore__ inline int32_t WriteNbi(
73 73 
74## Related Sample74## Related Sample
75 75 
76-[hcomm_write_read_nbi](../../../../../examples/hcomm_write_read_nbi/README_en.md) demonstrates the ordinary AIV direct-driven URMA read/write workflow. [hcomm_batch_write](../../../../../examples/hcomm_batch_write/README_en.md) demonstrates multi-channel batch writes. See the [Hcomm Usage Guide](../../../guide/hcomm_usage.md) for the basic BatchHandle workflow.76+[hcomm_write_read_nbi](../../../../../examples/aicore/hcomm/01_hcomm_write_read_nbi/README_en.md) demonstrates the ordinary AIV direct-driven URMA read/write workflow. [hcomm_batch_write](../../../../../examples/aicore/hcomm/02_hcomm_batch_write/README_en.md) demonstrates multi-channel batch writes. See the [Hcomm Usage Guide](../../../guide/hcomm_usage.md) for the basic BatchHandle workflow.
77 77 
78See [GetHandleRef](./GetHandleRef.md) for the multi-channel interface.78See [GetHandleRef](./GetHandleRef.md) for the multi-channel interface.
Mdocs/en/guide/build_and_test.md+2-2
@@ -201,11 +201,11 @@ bash build.sh -t --cann_3rd_lib_path=<path-to-third-party>
201 201 
202## Build and Run Samples202## Build and Run Samples
203 203 
204-`examples/hcomm_write_read_nbi` provides a point-to-point communication sample using AIV direct-driven URMA `WriteNbi` and `ReadNbi`. This sample adopts an independent CMake project for building:204+`examples/aicore/hcomm/01_hcomm_write_read_nbi` provides a point-to-point communication sample using AIV direct-driven URMA `WriteNbi` and `ReadNbi`. This sample adopts an independent CMake project for building:
205 205 
206```bash206```bash
207source /usr/local/Ascend/cann/set_env.sh207source /usr/local/Ascend/cann/set_env.sh
208-cd examples/hcomm_write_read_nbi208+cd examples/aicore/hcomm/01_hcomm_write_read_nbi
209mkdir -p build209mkdir -p build
210cd build210cd build
211cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..211cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
Mdocs/en/guide/dependencies.md+3-3
@@ -19,7 +19,7 @@ The in-repository build of asc-comm is mainly used for environment verification,
19| Scenario | Dependency | Version | Acquisition & Configuration |19| Scenario | Dependency | Version | Acquisition & Configuration |
20| --- | --- | --- | --- |20| --- | --- | --- | --- |
21| UT | googletest | 1.14.0 | System GTest is preferred. If unavailable, point `CANN_3RD_LIB_PATH` to the CANN third_party directory. |21| UT | googletest | 1.14.0 | System GTest is preferred. If unavailable, point `CANN_3RD_LIB_PATH` to the CANN third_party directory. |
22-| Samples | CANN ASC CMake utilities and hcomm library | CANN 9.1.0 or later | After running `source ${install_path}/cann/set_env.sh`, build via CMake under `examples/hcomm_write_read_nbi`. |22+| Samples | CANN ASC CMake utilities and hcomm library | CANN 9.1.0 or later | After running `source ${install_path}/cann/set_env.sh`, build via CMake under `examples/aicore/hcomm/01_hcomm_write_read_nbi`. |
23| Code Formatting | clang-format | v18.1.8 | Pulled from `pre-commit-clang/mirrors-clang-format` defined in `pre-commit-config.yaml`. |23| Code Formatting | clang-format | v18.1.8 | Pulled from `pre-commit-clang/mirrors-clang-format` defined in `pre-commit-config.yaml`. |
24| Open Source Compliance Check | oat-py | >= 1.0.1 | `scripts/oat_check.sh` attempts automatic installation. Run `pip install oat-py>=1.0.1` manually upon failure. |24| Open Source Compliance Check | oat-py | >= 1.0.1 | `scripts/oat_check.sh` attempts automatic installation. Run `pip install oat-py>=1.0.1` manually upon failure. |
25 25 
@@ -27,14 +27,14 @@ The in-repository build of asc-comm is mainly used for environment verification,
27 27 
28## Sample Runtime Dependencies28## Sample Runtime Dependencies
29 29 
30-The `examples/hcomm_write_read_nbi` sample supports Ascend 950PR / Ascend 950DT and requires at least two NPUs for runtime. Compilation verification can be completed on single-NPU environments, whereas two-card point-to-point communication runtime verification is unavailable.30+The `examples/aicore/hcomm/01_hcomm_write_read_nbi` sample supports Ascend 950PR / Ascend 950DT and requires at least two NPUs for runtime. Compilation verification can be completed on single-NPU environments, whereas two-card point-to-point communication runtime verification is unavailable.
31This sample uses `COMM_ENGINE_AIV` and `COMM_PROTOCOL_UBC_CTP` exclusively and does not cover the RoCE path.31This sample uses `COMM_ENGINE_AIV` and `COMM_PROTOCOL_UBC_CTP` exclusively and does not cover the RoCE path.
32 32 
33Sample build commands:33Sample build commands:
34 34 
35```bash35```bash
36source /usr/local/Ascend/cann/set_env.sh36source /usr/local/Ascend/cann/set_env.sh
37-cd examples/hcomm_write_read_nbi37+cd examples/aicore/hcomm/01_hcomm_write_read_nbi
38mkdir -p build38mkdir -p build
39cd build39cd build
40cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..40cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
Mdocs/en/guide/hcomm_usage.md+1-1
@@ -184,6 +184,6 @@ The lock is not reentrant. Every successful `Lock` must be followed by `Unlock`
184 184 
185## Samples185## Samples
186 186 
187-See [hcomm_write_read_nbi](../../../examples/hcomm_write_read_nbi/README_en.md) for the ordinary AIV Kernel-side API workflow and Host-side communication resource creation. See [hcomm_batch_write](../../../examples/hcomm_batch_write/README_en.md) for the multi-channel batch write workflow.187+See [hcomm_write_read_nbi](../../../examples/aicore/hcomm/01_hcomm_write_read_nbi/README_en.md) for the ordinary AIV Kernel-side API workflow and Host-side communication resource creation. See [hcomm_batch_write](../../../examples/aicore/hcomm/02_hcomm_batch_write/README_en.md) for the multi-channel batch write workflow.
188 188 
189Both samples require Ascend 950PR/Ascend 950DT and at least two NPUs for runtime verification. A single-device environment supports compilation only.189Both samples require Ascend 950PR/Ascend 950DT and at least two NPUs for runtime verification. A single-device environment supports compilation only.
Mdocs/quick_start.md+3-3
@@ -151,7 +151,7 @@ CANN包分为CANN toolkit包和CANN ops包。
151 151 
152 > [!IMPORTANT]152 > [!IMPORTANT]
153 > 安装说明153 > 安装说明
154- > 当前[Hcomm AIV直驱URMA样例](../examples/hcomm_write_read_nbi/README.md)不依赖ops包；仅在后续使用依赖算子包的功能时按需安装。154+ > 当前[Hcomm AIV直驱URMA样例](../examples/aicore/hcomm/01_hcomm_write_read_nbi/README.md)不依赖ops包；仅在后续使用依赖算子包的功能时按需安装。
155 155 
156| 参数 | 说明 |156| 参数 | 说明 |
157| :--- | :-------------------------------- |157| :--- | :-------------------------------- |
@@ -264,13 +264,13 @@ cmake --build build/ut-hcomm
264 264 
265### 🧩 样例验证<a name="sample-verify"></a>265### 🧩 样例验证<a name="sample-verify"></a>
266 266 
267-[hcomm_write_read_nbi](../examples/hcomm_write_read_nbi/README.md)提供AIV直驱URMA `WriteNbi`和`ReadNbi`点对点通信样例。样例支持Ascend 950PR/Ascend 950DT，要求CANN 9.1.0或以上版本。运行样例需要至少2张NPU；单卡环境仅支持编译验证。267+[hcomm_write_read_nbi](../examples/aicore/hcomm/01_hcomm_write_read_nbi/README.md)提供AIV直驱URMA `WriteNbi`和`ReadNbi`点对点通信样例。样例支持Ascend 950PR/Ascend 950DT，要求CANN 9.1.0或以上版本。运行样例需要至少2张NPU；单卡环境仅支持编译验证。
268 268 
269进入样例目录后执行：269进入样例目录后执行：
270 270 
271```bash271```bash
272source /usr/local/Ascend/cann/set_env.sh272source /usr/local/Ascend/cann/set_env.sh
273-cd examples/hcomm_write_read_nbi273+cd examples/aicore/hcomm/01_hcomm_write_read_nbi
274mkdir -p build274mkdir -p build
275cd build275cd build
276cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..276cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
Mdocs/quick_start_en.md+3-3
@@ -148,7 +148,7 @@ CANN packages include the CANN toolkit package and CANN ops package.
148 148 
149 > [!IMPORTANT]149 > [!IMPORTANT]
150 > Installation Note150 > Installation Note
151- > The [AIV direct-driven URMA Hcomm sample](../examples/hcomm_write_read_nbi/README_en.md) does not depend on the ops package. Install it later only when you require features relying on the operator package.151+ > The [AIV direct-driven URMA Hcomm sample](../examples/aicore/hcomm/01_hcomm_write_read_nbi/README_en.md) does not depend on the ops package. Install it later only when you require features relying on the operator package.
152 152 
153| Parameter | Description |153| Parameter | Description |
154| :--- | :-------------------------------- |154| :--- | :-------------------------------- |
@@ -261,14 +261,14 @@ Third-party open-source software used for UT execution:
261 261 
262### 🧩 Sample Verification<a name="sample-verify"></a>262### 🧩 Sample Verification<a name="sample-verify"></a>
263 263 
264-[hcomm_write_read_nbi](../examples/hcomm_write_read_nbi/README_en.md) provides a point-to-point communication sample using AIV direct-driven URMA `WriteNbi` and `ReadNbi`.264+[hcomm_write_read_nbi](../examples/aicore/hcomm/01_hcomm_write_read_nbi/README_en.md) provides a point-to-point communication sample using AIV direct-driven URMA `WriteNbi` and `ReadNbi`.
265The sample supports Ascend 950PR / Ascend 950DT and requires CANN 9.1.0 or newer. At least two NPUs are required for runtime; single-NPU environments only support compilation verification.265The sample supports Ascend 950PR / Ascend 950DT and requires CANN 9.1.0 or newer. At least two NPUs are required for runtime; single-NPU environments only support compilation verification.
266 266 
267Navigate to the sample directory and run:267Navigate to the sample directory and run:
268 268 
269```bash269```bash
270source /usr/local/Ascend/cann/set_env.sh270source /usr/local/Ascend/cann/set_env.sh
271-cd examples/hcomm_write_read_nbi271+cd examples/aicore/hcomm/01_hcomm_write_read_nbi
272mkdir -p build272mkdir -p build
273cd build273cd build
274cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..274cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
Mdocs/zh/api/README.md+2-2
@@ -69,5 +69,5 @@ AICore侧：
69## 相关文档69## 相关文档
70 70 
71- [Hcomm使用说明](../guide/hcomm_usage.md)71- [Hcomm使用说明](../guide/hcomm_usage.md)
72-- [AIV直驱URMA WriteNbi/ReadNbi样例](../../../examples/hcomm_write_read_nbi/README.md)72+- [AIV直驱URMA WriteNbi/ReadNbi样例](../../../examples/aicore/hcomm/01_hcomm_write_read_nbi/README.md)
73-- [Ain Basic Ring样例](../../../examples/ain/basic_ring/README.md)73+- [Ain Basic Ring样例](../../../examples/aicore/ain/01_basic_ring/README.md)
Mdocs/zh/api/aicore/hcomm/Hcomm.md+1-1
@@ -84,4 +84,4 @@ class Hcomm;
84 84 
85## 相关样例85## 相关样例
86 86 
87-[hcomm_write_read_nbi](../../../../../examples/hcomm_write_read_nbi/README.md)演示两卡场景下，AIV Kernel通过`COMM_PROTOCOL_UBC_CTP`路径调用普通`WriteNbi`和`ReadNbi`。该样例不覆盖RoCE路径和BatchHandle流程。87+[hcomm_write_read_nbi](../../../../../examples/aicore/hcomm/01_hcomm_write_read_nbi/README.md)演示两卡场景下，AIV Kernel通过`COMM_PROTOCOL_UBC_CTP`路径调用普通`WriteNbi`和`ReadNbi`。该样例不覆盖RoCE路径和BatchHandle流程。
Mdocs/zh/api/aicore/hcomm/ReadNbi.md+1-1
@@ -75,6 +75,6 @@ __aicore__ inline int32_t ReadNbi(
75 75 
76## 相关样例76## 相关样例
77 77 
78-[hcomm_write_read_nbi](../../../../../examples/hcomm_write_read_nbi/README.md)演示普通接口的AIV直驱URMA两卡读写流程。BatchHandle基本流程见[Hcomm使用说明](../../../guide/hcomm_usage.md)。78+[hcomm_write_read_nbi](../../../../../examples/aicore/hcomm/01_hcomm_write_read_nbi/README.md)演示普通接口的AIV直驱URMA两卡读写流程。BatchHandle基本流程见[Hcomm使用说明](../../../guide/hcomm_usage.md)。
79 79 
80多通道接口见[GetHandleRef](./GetHandleRef.md)。80多通道接口见[GetHandleRef](./GetHandleRef.md)。
Mdocs/zh/api/aicore/hcomm/WriteNbi.md+1-1
@@ -75,6 +75,6 @@ __aicore__ inline int32_t WriteNbi(
75 75 
76## 相关样例76## 相关样例
77 77 
78-[hcomm_write_read_nbi](../../../../../examples/hcomm_write_read_nbi/README.md)演示普通接口的AIV直驱URMA两卡读写流程；[hcomm_batch_write](../../../../../examples/hcomm_batch_write/README.md)演示多通道批量写流程。BatchHandle基本流程见[Hcomm使用说明](../../../guide/hcomm_usage.md)。78+[hcomm_write_read_nbi](../../../../../examples/aicore/hcomm/01_hcomm_write_read_nbi/README.md)演示普通接口的AIV直驱URMA两卡读写流程；[hcomm_batch_write](../../../../../examples/aicore/hcomm/02_hcomm_batch_write/README.md)演示多通道批量写流程。BatchHandle基本流程见[Hcomm使用说明](../../../guide/hcomm_usage.md)。
79 79 
80多通道接口见[GetHandleRef](./GetHandleRef.md)。80多通道接口见[GetHandleRef](./GetHandleRef.md)。
Mdocs/zh/api/usage_ref.md+2-2
@@ -8,5 +8,5 @@
8 8 
9## 样例代码9## 样例代码
10 10 
11-- [AIV直驱URMA WriteNbi/ReadNbi样例](../../../examples/hcomm_write_read_nbi/README.md)：展示如何在AIV Kernel中通过Hcomm的WriteNbi/ReadNbi接口实现NPU间低时延点对点通信，包含Host侧通信域创建、通道建立及Kernel侧Init → WriteNbi/ReadNbi → Drain的完整流程，支持Ascend 950PR/Ascend 950DT多卡环形拓扑。11+- [AIV直驱URMA WriteNbi/ReadNbi样例](../../../examples/aicore/hcomm/01_hcomm_write_read_nbi/README.md)：展示如何在AIV Kernel中通过Hcomm的WriteNbi/ReadNbi接口实现NPU间低时延点对点通信，包含Host侧通信域创建、通道建立及Kernel侧Init → WriteNbi/ReadNbi → Drain的完整流程，支持Ascend 950PR/Ascend 950DT多卡环形拓扑。
12-- [Ain Basic Ring样例](../../../examples/ain/basic_ring/README.md)：展示如何通过AIN接口在AIV Kernel中实现多rank环形通信，包括Put写入对端window、Get读取对端window、Flush等待通信完成及AinBarrierSession::Sync同步，支持Ascend 950PR/Ascend 950DT单机多卡场景。12+- [Ain Basic Ring样例](../../../examples/aicore/ain/01_basic_ring/README.md)：展示如何通过AIN接口在AIV Kernel中实现多rank环形通信，包括Put写入对端window、Get读取对端window、Flush等待通信完成及AinBarrierSession::Sync同步，支持Ascend 950PR/Ascend 950DT单机多卡场景。
Mdocs/zh/guide/build_and_test.md+2-2
@@ -201,11 +201,11 @@ bash build.sh -t --cann_3rd_lib_path=<path-to-third-party>
201 201 
202## 样例构建与运行202## 样例构建与运行
203 203 
204-`examples/hcomm_write_read_nbi`提供AIV直驱URMA `WriteNbi`和`ReadNbi`点对点通信样例。该样例使用独立CMake工程构建：204+`examples/aicore/hcomm/01_hcomm_write_read_nbi`提供AIV直驱URMA `WriteNbi`和`ReadNbi`点对点通信样例。该样例使用独立CMake工程构建：
205 205 
206```bash206```bash
207source /usr/local/Ascend/cann/set_env.sh207source /usr/local/Ascend/cann/set_env.sh
208-cd examples/hcomm_write_read_nbi208+cd examples/aicore/hcomm/01_hcomm_write_read_nbi
209mkdir -p build209mkdir -p build
210cd build210cd build
211cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..211cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
Mdocs/zh/guide/dependencies.md+3-3
@@ -19,7 +19,7 @@ asc-comm当前仓内构建主要用于环境检查、AICore Hcomm接口UT验证
19| 场景 | 依赖 | 版本 | 获取或配置方式 |19| 场景 | 依赖 | 版本 | 获取或配置方式 |
20| --- | --- | --- | --- |20| --- | --- | --- | --- |
21| UT | googletest | 1.14.0 | 优先使用系统GTest；没有系统GTest时，通过`CANN_3RD_LIB_PATH`指向CANN third_party目录。 |21| UT | googletest | 1.14.0 | 优先使用系统GTest；没有系统GTest时，通过`CANN_3RD_LIB_PATH`指向CANN third_party目录。 |
22-| 样例 | CANN ASC CMake能力和hcomm库 | CANN 9.1.0或以上 | 执行`source ${install_path}/cann/set_env.sh`后，在`examples/hcomm_write_read_nbi`目录下使用CMake构建。 |22+| 样例 | CANN ASC CMake能力和hcomm库 | CANN 9.1.0或以上 | 执行`source ${install_path}/cann/set_env.sh`后，在`examples/aicore/hcomm/01_hcomm_write_read_nbi`目录下使用CMake构建。 |
23| 代码格式化 | clang-format | v18.1.8 | `pre-commit-config.yaml`从`pre-commit-clang/mirrors-clang-format`拉取。 |23| 代码格式化 | clang-format | v18.1.8 | `pre-commit-config.yaml`从`pre-commit-clang/mirrors-clang-format`拉取。 |
24| 开源合规检查 | oat-py | >= 1.0.1 | `scripts/oat_check.sh`会尝试自动安装；失败时手动执行`pip install oat-py>=1.0.1`。 |24| 开源合规检查 | oat-py | >= 1.0.1 | `scripts/oat_check.sh`会尝试自动安装；失败时手动执行`pip install oat-py>=1.0.1`。 |
25 25 
@@ -27,13 +27,13 @@ asc-comm当前仓内构建主要用于环境检查、AICore Hcomm接口UT验证
27 27 
28## 样例运行依赖28## 样例运行依赖
29 29 
30-`examples/hcomm_write_read_nbi`样例支持Ascend 950PR/Ascend 950DT，运行时需要至少2张NPU。单卡环境可完成编译验证，但无法完成两卡点对点通信运行验证。该样例固定使用`COMM_ENGINE_AIV`和`COMM_PROTOCOL_UBC_CTP`，不覆盖RoCE路径。30+`examples/aicore/hcomm/01_hcomm_write_read_nbi`样例支持Ascend 950PR/Ascend 950DT，运行时需要至少2张NPU。单卡环境可完成编译验证，但无法完成两卡点对点通信运行验证。该样例固定使用`COMM_ENGINE_AIV`和`COMM_PROTOCOL_UBC_CTP`，不覆盖RoCE路径。
31 31 
32样例编译命令如下：32样例编译命令如下：
33 33 
34```bash34```bash
35source /usr/local/Ascend/cann/set_env.sh35source /usr/local/Ascend/cann/set_env.sh
36-cd examples/hcomm_write_read_nbi36+cd examples/aicore/hcomm/01_hcomm_write_read_nbi
37mkdir -p build37mkdir -p build
38cd build38cd build
39cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..39cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
Mdocs/zh/guide/hcomm_usage.md+1-1
@@ -184,6 +184,6 @@ if (lockRet == 0) {
184 184 
185## 样例185## 样例
186 186 
187-可参考[hcomm_write_read_nbi](../../../examples/hcomm_write_read_nbi/README.md)了解普通接口的AIV Kernel侧调用方式和Host侧通信资源创建流程；参考[hcomm_batch_write](../../../examples/hcomm_batch_write/README.md)了解多通道批量写流程。187+可参考[hcomm_write_read_nbi](../../../examples/aicore/hcomm/01_hcomm_write_read_nbi/README.md)了解普通接口的AIV Kernel侧调用方式和Host侧通信资源创建流程；参考[hcomm_batch_write](../../../examples/aicore/hcomm/02_hcomm_batch_write/README.md)了解多通道批量写流程。
188 188 
189两个样例均依赖Ascend 950PR/Ascend 950DT和至少2张NPU完成运行验证；单卡环境仅支持编译验证。189两个样例均依赖Ascend 950PR/Ascend 950DT和至少2张NPU完成运行验证；单卡环境仅支持编译验证。
Mdocs/zh/programming_model/memory_model.md+1-1
@@ -57,7 +57,7 @@
57- 远端地址和访问长度用于描述对端内存中的目标范围。57- 远端地址和访问长度用于描述对端内存中的目标范围。
58- 访问凭据用于表明本次操作获得了目标内存范围的访问授权。58- 访问凭据用于表明本次操作获得了目标内存范围的访问授权。
59 59 
60-具体的通信内存注册、Channel创建和远端内存获取流程，可参考[Hcomm AIV直驱URMA样例](../../../examples/hcomm_write_read_nbi/README.md)。60+具体的通信内存注册、Channel创建和远端内存获取流程，可参考[Hcomm AIV直驱URMA样例](../../../examples/aicore/hcomm/01_hcomm_write_read_nbi/README.md)。
61 61 
62### 访问凭据62### 访问凭据
63 63 
Mexamples/README.md+9-20
@@ -8,30 +8,19 @@
8 8 
9| 样例 | 说明 | 支持产品 |9| 样例 | 说明 | 支持产品 |
10| --- | --- | --- |10| --- | --- | --- |
11-| [ain/basic_ring](./ain/basic_ring/README.md) | 演示多卡环形场景下AIV Kernel通过AIN接口调用`Put`/`Get`进行点对点单边通信，并通过`AinBarrierSession`完成同步与结果校验。 | Ascend 950PR/Ascend 950DT |11+| [aicore/ain/01_basic_ring](./aicore/ain/01_basic_ring/README.md) | 演示多卡环形场景下AIV Kernel通过AIN接口调用`Put`/`Get`进行点对点单边通信，并通过`AinBarrierSession`完成同步与结果校验。 | Ascend 950PR/Ascend 950DT |
12-| [hcomm_batch_write](./hcomm_batch_write/README.md) | 演示多个URMA Channel共享Jetty，并通过`MakeBatchHandle`、`GetHandleRef`、`BatchCommit`和`Drain`批量提交跨peer写任务。 | Ascend 950PR/Ascend 950DT |12+| [aicore/hcomm/01_hcomm_write_read_nbi](./aicore/hcomm/01_hcomm_write_read_nbi/README.md) | 演示多卡场景下AIV Kernel通过URMA路径调用`Hcomm::WriteNbi`和`Hcomm::ReadNbi`，并校验通信结果。 | Ascend 950PR/Ascend 950DT |
13+| [aicore/hcomm/02_hcomm_batch_write](./aicore/hcomm/02_hcomm_batch_write/README.md) | 演示多个URMA Channel共享Jetty，并通过`MakeBatchHandle`、`GetHandleRef`、`BatchCommit`和`Drain`批量提交跨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 |
16+| [aicore/hcomm/05_simt_urma_perftest](./aicore/hcomm/05_simt_urma_perftest/README.md) | 测量上述五个 SIMT URMA 接口的下发时延与完成带宽，覆盖立即提交/延迟提交两种发布策略。 | Ascend 950PR / Ascend 950DT |
13| [hcomm_jetty_write](./hcomm_jetty_write/README.md) | 演示多卡场景下AIV Kernel通过`HcommJetty::Write`和`HcommJetty::WriteValue`直接向Jetty SQ提交WQE，并通过`Drain`等待完成与校验结果。 | 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 |
14-| [hcomm_write_read_nbi](./hcomm_write_read_nbi/README.md) | 演示多卡场景下AIV Kernel通过URMA路径调用`Hcomm::WriteNbi`和`Hcomm::ReadNbi`，并校验通信结果。 | Ascend 950PR/Ascend 950DT |18+| [ccu/ccu_direct](./ccu/ccu_direct/01_allgather/README.md) | 演示基于HCCL通信域和CCU数据面接口，以直调`<<<>>>`方式实现AllGather等集合通信操作。 | Ascend 950PR/Ascend 950DT |
15-| [one_multi_path](./one_multi_path/README.md) | 查询`UB_MEM`链路，为每个peer创建one path/multi path Channel，逐个处理peer，并通过双Stream并发搬运和校验该peer的远端数据。 | Ascend 950PR/Ascend 950DT |
16-| [simt_urma](./simt_urma/README.md) | 演示并验证全部 SIMT URMA 接口：`WriteNbi`、`WriteValueNbi`、`WriteWithNotifyNbi`、`AtomicFAA` 和 `AtomicCAS`，覆盖立即提交与批量提交。 | Ascend 950PR / Ascend 950DT |
17-| [simt_urma_perftest](./simt_urma_perftest/README.md) | 测量上述五个 SIMT URMA 接口的下发时延与完成带宽，覆盖立即提交/延迟提交两种发布策略。 | Ascend 950PR / Ascend 950DT |
18 19 
19-## simt_urma20+各样例的编译与运行方式参见对应目录下的README。
20- 
21-`simt_urma`验证SIMT Kernel通过URMA路径调用`WriteNbi`、`WriteValueNbi`、`WriteWithNotifyNbi`、
22-`AtomicFAA`和`AtomicCAS`。样例包含单接口用例，以及串行提交和batch-last提交场景，共十种模式。
23- 
24-编译运行方式参见[simt_urma/README.md](./simt_urma/README.md)。
25- 
26-## simt_urma_perftest
27- 
28-`simt_urma_perftest`测量上述五个接口的性能，分别统计下发时延与完成带宽，并对`write`和`notify`
29-支持payload扫描。
30- 
31-性能测试方式参见[simt_urma_perftest/README.md](./simt_urma_perftest/README.md)。
32 21 
33## 运行约束22## 运行约束
34 23 
35- 样例支持Ascend 950PR/Ascend 950DT，CANN软件版本要求为9.1.0或以上。24- 样例支持Ascend 950PR/Ascend 950DT，CANN软件版本要求为9.1.0或以上。
36- 样例运行需要至少2张NPU；单卡环境仅支持编译验证。25- 样例运行需要至少2张NPU；单卡环境仅支持编译验证。
37-- 样例编译依赖CANN ASC CMake能力，并在链接阶段依赖CANN `hcomm`库。26+- 样例编译依赖CANN ASC CMake能力，并在链接阶段依赖CANN `hcomm`库。
Mexamples/README_en.md+10-18
@@ -8,27 +8,19 @@ This directory contains usage samples for asc-comm APIs.
8 8 
9| Sample | Description | Supported Products |9| Sample | Description | Supported Products |
10| --- | --- | --- |10| --- | --- | --- |
11-| [ain/basic_ring](./ain/basic_ring/README_en.md) | Demonstrates AIV Kernel invoking `Put`/`Get` over the AIN interface for point-to-point one-sided communication in a multi-card ring scenario, with synchronization and result verification through `AinBarrierSession`. | Ascend 950PR / Ascend 950DT |11+| [aicore/ain/01_basic_ring](./aicore/ain/01_basic_ring/README_en.md) | Demonstrates AIV Kernel invoking `Put`/`Get` over the AIN interface for point-to-point one-sided communication in a multi-card ring scenario, with synchronization and result verification through `AinBarrierSession`. | Ascend 950PR / Ascend 950DT |
12-| [hcomm_batch_write](./hcomm_batch_write/README_en.md) | Demonstrates multiple URMA channels sharing a Jetty and submitting cross-peer writes through `MakeBatchHandle`, `GetHandleRef`, `BatchCommit`, and `Drain`. | Ascend 950PR / Ascend 950DT |12+| [aicore/hcomm/01_hcomm_write_read_nbi](./aicore/hcomm/01_hcomm_write_read_nbi/README_en.md) | Demonstrates AIV Kernel invoking `Hcomm::WriteNbi` and `Hcomm::ReadNbi` over the URMA path in a multi-card scenario, with communication result verification. | Ascend 950PR / Ascend 950DT |
13+| [aicore/hcomm/02_hcomm_batch_write](./aicore/hcomm/02_hcomm_batch_write/README_en.md) | Demonstrates multiple URMA channels sharing a Jetty and submitting cross-peer writes through `MakeBatchHandle`, `GetHandleRef`, `BatchCommit`, and `Drain`. | 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 |
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 |
13| [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 |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 |
14-| [hcomm_write_read_nbi](./hcomm_write_read_nbi/README_en.md) | Demonstrates AIV Kernel invoking `Hcomm::WriteNbi` and `Hcomm::ReadNbi` over the URMA path in a multi-card scenario, with communication result 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 |
15-| [one_multi_path](./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 |
16-| [simt_urma](./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 |
17-| [simt_urma_perftest](./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 |
18 19 
19-## simt_urma20+See the README in each sample directory for build and run instructions.
20- 
21-`simt_urma` validates SIMT URMA `WriteNbi`, `WriteValueNbi`, `WriteWithNotifyNbi`, `AtomicFAA`, and `AtomicCAS`, covering isolated interface cases plus serialized and batch-last submission across ten modes.
22- 
23-See [simt_urma/README_en.md](./simt_urma/README_en.md) for build and run instructions.
24- 
25-## simt_urma_perftest
26- 
27-`simt_urma_perftest` measures those five interfaces, reporting issue latency and completion bandwidth, with payload sweeps for `write` and `notify`.
28- 
29-See [simt_urma_perftest/README_en.md](./simt_urma_perftest/README_en.md) for performance testing instructions.
30 21 
31## Runtime Constraints22## Runtime Constraints
23+ 
32- Supported chips: Ascend 950PR / Ascend 950DT. CANN version 9.1.0 or later is required.24- Supported chips: Ascend 950PR / Ascend 950DT. CANN version 9.1.0 or later is required.
33- At least two NPUs are required for runtime execution; single-NPU environments only support compilation verification.25- At least two NPUs are required for runtime execution; single-NPU environments only support compilation verification.
34-- Sample build relies on CANN ASC CMake utilities and links against the CANN `hcomm` library at link time.26+- Sample build relies on CANN ASC CMake utilities and links against the CANN `hcomm` library.
Rexamples/ain/basic_ring/CMakeLists.txt→examples/aicore/ain/01_basic_ring/CMakeLists.txt+16-29
@@ -10,50 +10,37 @@
10 10 
11cmake_minimum_required(VERSION 3.16)11cmake_minimum_required(VERSION 3.16)
12 12 
13-set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture")13+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
14- 
15-if(NOT CMAKE_ASC_ARCHITECTURES STREQUAL "dav-3510")
16- message(FATAL_ERROR "ain_basic_ring_sample only supports dav-3510, but got ${CMAKE_ASC_ARCHITECTURES}")
17-endif()
18 14 
19find_package(ASC REQUIRED)15find_package(ASC REQUIRED)
20- 16+project(basic_ring_demo LANGUAGES ASC CXX)
21-project(ain_basic_ring_sample LANGUAGES ASC CXX)
22- 
23-set(CMAKE_CXX_STANDARD 17)
24-set(CMAKE_CXX_STANDARD_REQUIRED ON)
25- 
26-set_source_files_properties(
27- basic_ring.asc
28- PROPERTIES LANGUAGE ASC
29-)
30 17 
31add_executable(basic_ring_demo18add_executable(basic_ring_demo
32 basic_ring.asc19 basic_ring.asc
33)20)
34 21 
22+set_target_properties(basic_ring_demo PROPERTIES
23+ LINKER_LANGUAGE ASC
24+)
25+ 
35target_include_directories(basic_ring_demo PRIVATE26target_include_directories(basic_ring_demo PRIVATE
36- ${ASCEND_CANN_PACKAGE_PATH}/pkg_inc27+ ${CMAKE_CURRENT_SOURCE_DIR}
37- ${ASCEND_CANN_PACKAGE_PATH}/include28+ ${CMAKE_CURRENT_SOURCE_DIR}/../..
38- ${ASCEND_CANN_PACKAGE_PATH}/include/hcomm29+ $ENV{ASCEND_HOME_PATH}/include
39- ${ASCEND_CANN_PACKAGE_PATH}/asc/include/comm_api/aicore30+ $ENV{ASCEND_HOME_PATH}/asc/include/comm_api/aicore
31+ $ENV{ASCEND_HOME_PATH}/asc/include/comm_api/hcomm
40)32)
41 33 
42target_link_directories(basic_ring_demo PRIVATE34target_link_directories(basic_ring_demo PRIVATE
43- ${ASCEND_CANN_PACKAGE_PATH}/lib6435+ $ENV{ASCEND_HOME_PATH}/lib64
36+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/lib64
44)37)
45 38 
46target_link_libraries(basic_ring_demo PRIVATE39target_link_libraries(basic_ring_demo PRIVATE
47- -Wl,--no-as-needed
48- hccl
49 hcomm40 hcomm
50- -Wl,--as-needed41+ hccl
51- ascendcl
52- runtime
53- stdc++
54- dl
55)42)
56 43 
57-set_target_properties(basic_ring_demo PROPERTIES44+target_compile_options(basic_ring_demo PRIVATE
58- BUILD_RPATH "${ASCEND_CANN_PACKAGE_PATH}/lib64;${ASCEND_CANN_PACKAGE_LINUX_PATH}/lib64"45+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
59)46)
Aexamples/aicore/ain/01_basic_ring/README.md+121-0
@@ -0,0 +1,121 @@
1+# AIN Basic Ring样例
2+ 
3+## 概述
4+ 
5+本样例展示如何在Ascend C AIV Kernel中通过`ain/ain.h`接口使用AIN基础通信能力。样例采用多rank
6+对称执行方式：每个rank向下一个rank执行`Put`，并对上一个rank执行`Get`，最后通过
7+`AinBarrierSession`完成同步。
8+ 
9+Host侧负责内存申请、初始化HCCL通信域、创建HCCL Team、注册对称window并创建AIV通信channel；
10+Kernel侧只通过AIN接口提交数据面通信任务。
11+ 
12+## 本样例支持的产品及CANN软件版本
13+ 
14+| 产品 | CANN软件版本 |
15+| --- | --- |
16+| Ascend 950PR/Ascend 950DT | >= CANN 9.1.0 |
17+ 
18+## 目录结构介绍
19+ 
20+```text
21+basic_ring
22+├── CMakeLists.txt // 编译工程文件
23+├── basic_ring.asc // Host侧资源准备、AICore侧AIN调用和共享定义
24+├── README.md // 中文样例说明
25+└── README_en.md // 英文样例说明
26+```
27+ 
28+## 样例描述
29+ 
30+### 功能说明
31+ 
32+`nranks`个rank组成环形拓扑，当前rank `r`执行两次单边通信：
33+ 
34+- `Put`：将本rank标识写入下一个rank（`(r + 1) % nranks`）的`recvWindow[rankId]`
35+- `Get`：从上一个rank（`(r - 1 + nranks) % nranks`）的`sendWindow[0]`读取数据到本rank的
36+ `recvWindow[rankNum + rankId]`
37+ 
38+通信完成后每个rank回读`recvBuf`并校验`Put`和`Get`的结果。
39+ 
40+### 样例规格
41+ 
42+| 项目 | 说明 |
43+| --- | --- |
44+| 通信模式 | 环形拓扑AIN点对点单边通信 |
45+| 通信引擎 | `COMM_ENGINE_AIV`，协议`COMM_PROTOCOL_UB_CTP` |
46+| 部署形态 | 单机多卡 |
47+| 支持rank数 | >= 2，且不超过当前可用NPU数量 |
48+ 
49+### 实现流程
50+ 
51+1. 可执行文件直接拉起全部rank进程（进程内fork，参见`examples/aicore/utils/process_manager.h`）；
52+ rank间通过命令行指定的TCP控制通道（`examples/aicore/utils/rank_sync.h`）交换root info并
53+ 创建多rank HCCL通信域。
54+2. 每个rank初始化ACL运行环境，通过`HcclTeamCreateDescInit`配置team描述（包括`barrierCount`、
55+ `rankIds`、通信引擎`COMM_ENGINE_AIV`、协议`COMM_PROTOCOL_UB_CTP`和`channelCnt`），
56+ 调用`HcclTeamCreate`创建team并自动建立AIV + UB_CTP通信channel。
57+3. 分别将`sendBuf`和`recvBuf`通过`HcclCommSymWinRegister`注册为对称window。
58+4. 下发`team`、`sendWin`、`recvWin`到Kernel。
59+5. Kernel侧构造`AscendC::Ain`对象并初始化Hcomm UB临时工作区，执行`Put`和`Get`后调用`Flush`
60+ 等待通信任务完成，最后通过`AinBarrierSession::Sync`完成rank间同步。
61+ 
62+## 编译运行
63+ 
64+在本样例根目录下执行如下步骤，编译并运行样例。本样例仅支持NPU运行模式。
65+ 
66+- 配置环境变量
67+ 
68+ 请根据当前环境上CANN开发套件包的安装方式配置环境变量。
69+ 
70+ ```bash
71+ source ${install_path}/cann/set_env.sh
72+ ```
73+ 
74+ > **说明：** `${install_path}`为CANN包安装目录，未指定安装目录时默认安装至`/usr/local/Ascend`。
75+ 
76+- 样例执行
77+ 
78+ 在本样例目录下执行如下命令。可执行文件内部fork出全部rank进程，直接运行即可：
79+ 
80+ ```bash
81+ mkdir -p build
82+ cd build
83+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
suqwe
suqwesuqwe9月17日

为什么统一成最麻烦的方式呢

likedislike
wanghan
wanghan
9月17日 评论：
84+ make -j
85+ cd .. && ./build/basic_ring_demo tcp://127.0.0.1:29622 <nranks>
86+ ```
87+ 
88+- 启动参数说明
89+ 
90+ | 参数 | 默认值 | 说明 |
91+ | --- | --- | --- |
92+ | `<ip:port>` | 无默认 | rank 0监听的控制通道地址（`tcp://ip:port`格式），用于root info交换；并行运行多个实例时错开端口 |
93+ | `<nranks>` | 无默认 | 启动的rank数量，必须为不小于2的整数，且不超过当前可用NPU数量 |
94+ 
95+- 编译选项说明
96+ 
97+ | 选项 | 可选值 | 说明 |
98+ | --- | --- | --- |
99+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认） | 运行模式，本样例仅支持NPU运行 |
100+ | `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU架构，对应Ascend 950PR/Ascend 950DT |
101+ 
102+- 执行结果
103+ 
104+ 每个rank会回读`recvBuf`并打印`Put`和`Get`结果，例如：
105+ 
106+ ```text
107+ [rank 0] basic_ring | PUT recvBuf[1]=1, GET recvBuf[2]=1 | PASS
108+ [rank 1] basic_ring | PUT recvBuf[0]=0, GET recvBuf[3]=0 | PASS
109+ RESULT | Example=basic_ring | Status=PASS
110+ ```
111+ 
112+ 所有rank均输出`PASS`且最终输出`Status=PASS`表示样例执行成功。
113+ 
114+## 注意事项
115+ 
116+- 运行样例需要至少2张NPU；单卡环境仅支持编译验证。
117+- 编译和运行前必须先加载CANN环境变量，确保`ASCEND_CANN_PACKAGE_PATH`、ASC CMake模块和运行时
118+ 动态库可用。
119+- 样例依赖CANN中的`hccl`、`hcomm`、`ascendcl`和`runtime`库。
120+- 运行时rank数量不能超过当前可用NPU数量。
121+- 样例仅支持单机多卡运行，不支持跨节点多机运行。
Rexamples/ain/basic_ring/README_en.md→examples/aicore/ain/01_basic_ring/README_en.md+12-14
@@ -26,8 +26,7 @@ basic_ring
26├── CMakeLists.txt // Build project file26├── CMakeLists.txt // Build project file
27├── README.md // Chinese README27├── README.md // Chinese README
28├── README_en.md // English README28├── README_en.md // English README
29-├── basic_ring.asc // Host-side setup, AICore-side AIN calls, and shared definitions29+└── basic_ring.asc // Host-side setup, AICore-side AIN calls, and shared definitions
30-└── run.sh // Multi-rank launch script
31```30```
32 31 
33## Sample Flow32## Sample Flow
@@ -35,11 +34,10 @@ basic_ring
35### Host-Side Resource Setup34### Host-Side Resource Setup
36 35 
371. Call `aclInit`, `aclrtSetDevice`, and `aclrtCreateStream` to initialize the runtime environment.361. Call `aclInit`, `aclrtSetDevice`, and `aclrtCreateStream` to initialize the runtime environment.
38-2. Create the multi-rank HCCL communicator through `HcclGetRootInfo` and `HcclCommInitRootInfo`.37+2. Create the multi-rank HCCL communicator through the `RankSyncContext` helper (which exchanges root info over a TCP control channel) and `examples::InitCommByRootInfo`.
39-3. Configure `barrierCount` and call `HcclWorldTeamCreate` to create the world team.38+3. Initialize a team descriptor via `HcclTeamCreateDescInit` (including `barrierCount`, `rankIds`, engine `COMM_ENGINE_AIV`, protocol `COMM_PROTOCOL_UB_CTP`, and `channelCnt`), then call `HcclTeamCreate` to create the team with AIV + UB_CTP communication channels.
40-4. Register `sendBuf` and `recvBuf` as symmetric windows through `HcclTeamWindowRegister`.39+4. Register `sendBuf` and `recvBuf` as symmetric windows through `HcclCommSymWinRegister`.
41-5. Call `HcclTeamChannelsCreate` to create AIV + UB_CTP communication channels.40+5. Pass `team`, `sendWin`, and `recvWin` to the Kernel.
42-6. Pass `worldTeam`, `sendWin`, and `recvWin` to the Kernel.
43 41 
44### Kernel-Side AIN Communication42### Kernel-Side AIN Communication
45 43 
@@ -74,28 +72,28 @@ make -j
74After the build completes, run the following command in the sample directory:72After the build completes, run the following command in the sample directory:
75 73 
76```bash74```bash
77-./run.sh <nranks>75+./build/basic_ring_demo tcp://127.0.0.1:29622 <nranks>
78```76```
79 77 
80For example, run with two cards:78For example, run with two cards:
81 79 
82```bash80```bash
83-./run.sh 281+./build/basic_ring_demo tcp://127.0.0.1:29622 2
84```82```
85 83 
86-`run.sh` starts one `basic_ring_demo` process for each rank and uses a temporary file to exchange the HCCL root info. The default script timeout is 300 seconds.84+The executable forks all rank processes internally and uses a TCP control channel to exchange the HCCL root info.
87 85 
88## Expected Result86## Expected Result
89 87 
90Each rank reads back `recvBuf` and prints the `Put` and `Get` results, for example:88Each rank reads back `recvBuf` and prints the `Put` and `Get` results, for example:
91 89 
92```text90```text
93-PUT RES rank 0 recvBuf[1]=1 expect=191+[rank 0] basic_ring | PUT recvBuf[1]=1, GET recvBuf[2]=1 | PASS
94-GET RES rank 0 recvBuf[2]=1 expect=192+[rank 1] basic_ring | PUT recvBuf[0]=0, GET recvBuf[3]=0 | PASS
95-rank 0 test pass93+RESULT | Example=basic_ring | Status=PASS
96```94```
97 95 
98-The sample succeeds when all ranks print `test pass`.96+The sample succeeds when all ranks print `PASS` and the final status is `PASS`.
99 97 
100## Build Option98## Build Option
101 99 
Rexamples/ain/basic_ring/basic_ring.asc→examples/aicore/ain/01_basic_ring/basic_ring.asc+76-163
@@ -13,19 +13,13 @@
13 * \brief AIN basic ring 样例.13 * \brief AIN basic ring 样例.
14 */14 */
15 15 
16-#include <chrono>
17#include <cstdint>16#include <cstdint>
18-#include <cstdio>
19#include <cstdlib>17#include <cstdlib>
20-#include <fstream>
21#include <iostream>18#include <iostream>
22#include <string>19#include <string>
23-#include <thread>
24#include <vector>20#include <vector>
25-#include <sys/stat.h>
26 21 
27#include "acl/acl.h"22#include "acl/acl.h"
28-#include "acl/acl_rt.h"
29#include "hccl/hccl.h"23#include "hccl/hccl.h"
30#include "hccl/hccl_comm.h"24#include "hccl/hccl_comm.h"
31#include "hccl/hccl_types.h"25#include "hccl/hccl_types.h"
@@ -34,123 +28,15 @@
34#include "kernel_operator.h"28#include "kernel_operator.h"
35#include "ain/ain.h"29#include "ain/ain.h"
36 30 
37-#define ACL_CHECK(expr) \31+#include "utils/process_manager.h"
38- do { \32+#include "utils/check.h"
39- aclError ret = (expr); \33+#include "utils/hccl_comm_init.h"
40- if (ret != ACL_SUCCESS) { \34+#include "utils/rank_sync.h"
41- std::cout << "acl error " << __FILE__ << ":" << __LINE__ << " ret=" << static_cast<int32_t>(ret) \35+#include "utils/arg_parser.h"
42- << std::endl; \
43- return ret; \
44- } \
45- } while (0)
46- 
47-#define HCCL_CHECK(expr) \
48- do { \
49- HcclResult ret = (expr); \
50- if (ret != HCCL_SUCCESS) { \
51- std::cout << "hccl error " << __FILE__ << ":" << __LINE__ << " ret=" << static_cast<int32_t>(ret) \
52- << std::endl; \
53- return ret; \
54- } \
55- } while (0)
56- 
57-#define HCOMM_CHECK(expr) \
58- do { \
59- HcommResult ret = (expr); \
60- if (ret != HCCL_SUCCESS) { \
61- std::cout << "hcomm error " << __FILE__ << ":" << __LINE__ << " ret=" << static_cast<int32_t>(ret) \
62- << std::endl; \
63- return ret; \
64- } \
65- } while (0)
66 36 
67namespace {37namespace {
68 38 
69constexpr size_t COMM_BUF_SIZE = 2U * 1024U * 1024U;39constexpr size_t COMM_BUF_SIZE = 2U * 1024U * 1024U;
70-constexpr uint32_t ROOT_INFO_WAIT_SECONDS = 60U;
71-constexpr uint32_t ROOT_INFO_POLL_MS = 100U;
72- 
73-bool IsRootInfoFileReady(const std::string& path)
74-{
75- struct stat st {};
76- return stat(path.c_str(), &st) == 0 && st.st_size >= static_cast<off_t>(sizeof(HcclRootInfo));
77-}
78- 
79-HcclResult WriteRootInfoFile(const std::string& path, const HcclRootInfo& rootInfo)
80-{
81- const std::string tmpPath = path + ".tmp";
82- std::ofstream ofs(tmpPath, std::ios::binary | std::ios::trunc);
83- if (!ofs.is_open()) {
84- std::cout << "open root info tmp file failed: " << tmpPath << std::endl;
85- return HCCL_E_OPEN_FILE_FAILURE;
86- }
87- ofs.write(reinterpret_cast<const char*>(&rootInfo), sizeof(rootInfo));
88- ofs.close();
89- if (!ofs) {
90- std::cout << "write root info file failed: " << tmpPath << std::endl;
91- return HCCL_E_SYSCALL;
92- }
93- if (std::rename(tmpPath.c_str(), path.c_str()) != 0) {
94- std::cout << "rename root info file failed: " << tmpPath << " -> " << path << std::endl;
95- return HCCL_E_SYSCALL;
96- }
97- return ::HCCL_SUCCESS;
98-}
99- 
100-HcclResult ReadRootInfoFile(const std::string& path, HcclRootInfo* rootInfo)
101-{
102- const uint32_t maxPollTimes = ROOT_INFO_WAIT_SECONDS * 1000U / ROOT_INFO_POLL_MS;
103- for (uint32_t i = 0; i < maxPollTimes; ++i) {
104- if (!IsRootInfoFileReady(path)) {
105- std::this_thread::sleep_for(std::chrono::milliseconds(ROOT_INFO_POLL_MS));
106- continue;
107- }
108- std::ifstream ifs(path, std::ios::binary);
109- if (!ifs.is_open()) {
110- return HCCL_E_OPEN_FILE_FAILURE;
111- }
112- ifs.read(reinterpret_cast<char*>(rootInfo), sizeof(*rootInfo));
113- if (ifs.gcount() == static_cast<std::streamsize>(sizeof(*rootInfo))) {
114- return ::HCCL_SUCCESS;
115- }
116- std::this_thread::sleep_for(std::chrono::milliseconds(ROOT_INFO_POLL_MS));
117- }
118- std::cout << "wait root info file timeout: " << path << std::endl;
119- return HCCL_E_TIMEOUT;
120-}
121- 
122-HcclResult InitCommByRootInfo(uint32_t rank, uint32_t nranks, const std::string& rootInfoFile, HcclComm* comm)
123-{
124- HcclRootInfo rootInfo{};
125- HcclResult ret = HCCL_SUCCESS;
126- if (rank == 0U) {
127- HCCL_CHECK(HcclGetRootInfo(&rootInfo));
128- ret = WriteRootInfoFile(rootInfoFile, rootInfo);
129- if (ret != HCCL_SUCCESS) {
130- return ret;
131- }
132- } else {
133- ret = ReadRootInfoFile(rootInfoFile, &rootInfo);
134- if (ret != HCCL_SUCCESS) {
135- return ret;
136- }
137- }
138- HCCL_CHECK(HcclCommInitRootInfo(nranks, &rootInfo, rank, comm));
139- return HCCL_SUCCESS;
140-}
141- 
142-int32_t ParseInt32(const char* arg)
143-{
144- if (arg == nullptr) {
145- return -1;
146- }
147- char* end = nullptr;
148- long value = std::strtol(arg, &end, 10);
149- if (end == arg || *end != '\0') {
150- return -1;
151- }
152- return static_cast<int32_t>(value);
153-}
154 40 
155} // namespace41} // namespace
156 42 
@@ -181,53 +67,83 @@ extern "C" [[bisheng::core_ratio(0, 1)]] __global__ __aicore__ void basic_ring_k
181 ainBarrier.Sync(AscendC::AinMemoryOrder::AIN_MEMORY_ORDER_RELAX, hcommUbuf);67 ainBarrier.Sync(AscendC::AinMemoryOrder::AIN_MEMORY_ORDER_RELAX, hcommUbuf);
182}68}
183 69 
70+int32_t RunBasicRingRank(uint32_t rank, uint32_t nranks, const std::string& ctrlEndpoint);
71+ 
184int32_t main(int32_t argc, char* argv[])72int32_t main(int32_t argc, char* argv[])
185{73{
186- if (argc < 4) {74+ if (argc != 3) {
187- std::cout << "Usage: " << argv[0] << " <rank> <nranks> <root_info_file>" << std::endl;75+ std::fprintf(stderr, "Usage: %s <tcp://ip:port> <nranks>\n", argv[0]);
76+ return 1;
77+ }
78+ uint32_t nranks = 0;
79+ if (!examples::ParseUint32(argv[2], nranks, 2)) {
80+ std::fprintf(stderr, "invalid nranks: %s, expected an integer >= 2\n", argv[2]);
81+ return 1;
82+ }
83+ const std::string endpoint = argv[1];
84+ 
85+ std::vector<examples::ArgsOf<decltype(RunBasicRingRank)>> perRankArgs;
86+ perRankArgs.reserve(nranks);
87+ for (uint32_t rank = 0; rank < nranks; ++rank) {
88+ perRankArgs.emplace_back(rank, nranks, endpoint);
89+ }
90+ 
91+ examples::ProcessGroup group;
92+ if (!group.Launch(nranks, RunBasicRingRank, std::move(perRankArgs))) {
93+ return 1;
94+ }
95+ const bool pass = group.WaitAll();
96+ 
97+ if (group.Interrupted()) {
98+ std::cout << "RESULT | Example=basic_ring | Status=INTERRUPTED" << std::endl;
99+ return 1;
100+ }
101+ std::cout << "RESULT | Example=basic_ring | Status=" << (pass ? "PASS" : "FAIL") << std::endl;
102+ return pass ? 0 : 1;
103+}
104+ 
105+// 单个rank的完整流程。由ProcessGroup在每个子进程中调用。
106+// 本样例host侧无barrier（rank间同步由设备侧AinBarrierSession完成）。
107+int32_t RunBasicRingRank(uint32_t rank, uint32_t nranks, const std::string& ctrlEndpoint)
108+{
109+ examples::RankSyncContext sync = rank == 0U ? examples::RankSyncContext::Listen(ctrlEndpoint, nranks) :
110+ examples::RankSyncContext::Connect(rank, nranks, ctrlEndpoint);
111+ if (!sync.Ok()) {
112+ std::cout << "[rank " << rank << "] control channel setup failed" << std::endl;
188 return 1;113 return 1;
189 }114 }
190 115 
191- int32_t rank = ParseInt32(argv[1]);116+ EXAMPLES_ACLCHECK(aclInit(nullptr));
192- int32_t nranks = ParseInt32(argv[2]);
193- std::string rootInfoFile = argv[3];
194- if (rank < 0 || nranks < 2 || rank >= nranks) {
195- std::cout << "Invalid rank arguments: rank=" << argv[1] << " nranks=" << argv[2] << std::endl;
196- return 1;
197- }
198- 
199- ACL_CHECK(aclInit(nullptr));
200 uint32_t devCount = 0;117 uint32_t devCount = 0;
201- ACL_CHECK(aclrtGetDeviceCount(&devCount));118+ EXAMPLES_ACLCHECK(aclrtGetDeviceCount(&devCount));
202- std::cout << "Found " << devCount << " NPU device(s), rank=" << rank << "/" << nranks << std::endl;119+ EXAMPLES_ACLCHECK(aclrtSetDevice(static_cast<uint32_t>(rank)));
203- ACL_CHECK(aclrtSetDevice(static_cast<uint32_t>(rank)));
204 aclrtStream stream = nullptr;120 aclrtStream stream = nullptr;
205- ACL_CHECK(aclrtCreateStream(&stream));121+ EXAMPLES_ACLCHECK(aclrtCreateStream(&stream));
206 122 
207 HcclComm hcclComm = nullptr;123 HcclComm hcclComm = nullptr;
208 // 创建通信域124 // 创建通信域
209- HCCL_CHECK(InitCommByRootInfo(static_cast<uint32_t>(rank), static_cast<uint32_t>(nranks), rootInfoFile, &hcclComm));125+ EXAMPLES_HCCLCHECK(examples::InitCommByRootInfo(sync, &hcclComm));
210 126 
211 // 初始化本地数据127 // 初始化本地数据
212 void* sendBuf = nullptr;128 void* sendBuf = nullptr;
213- HCOMM_CHECK(HcommMemAlloc(&sendBuf, COMM_BUF_SIZE));129+ EXAMPLES_HCOMMCHECK(HcommMemAlloc(&sendBuf, COMM_BUF_SIZE));
214 std::vector<int8_t> sendHost(COMM_BUF_SIZE, static_cast<int8_t>(rank));130 std::vector<int8_t> sendHost(COMM_BUF_SIZE, static_cast<int8_t>(rank));
215- ACL_CHECK(aclrtMemcpy(sendBuf, COMM_BUF_SIZE, sendHost.data(), COMM_BUF_SIZE, ACL_MEMCPY_HOST_TO_DEVICE));131+ EXAMPLES_ACLCHECK(aclrtMemcpy(sendBuf, COMM_BUF_SIZE, sendHost.data(), COMM_BUF_SIZE, ACL_MEMCPY_HOST_TO_DEVICE));
216 132 
217 void* recvBuf = nullptr;133 void* recvBuf = nullptr;
218- HCOMM_CHECK(HcommMemAlloc(&recvBuf, COMM_BUF_SIZE));134+ EXAMPLES_HCOMMCHECK(HcommMemAlloc(&recvBuf, COMM_BUF_SIZE));
219- ACL_CHECK(aclrtMemset(recvBuf, COMM_BUF_SIZE, 0, COMM_BUF_SIZE));135+ EXAMPLES_ACLCHECK(aclrtMemset(recvBuf, COMM_BUF_SIZE, 0, COMM_BUF_SIZE));
220 136 
221 // 将本地内存注册为window137 // 将本地内存注册为window
222 HcclCommSymWindow sendWin{nullptr};138 HcclCommSymWindow sendWin{nullptr};
223- HCCL_CHECK(HcclCommSymWinRegister(hcclComm, sendBuf, COMM_BUF_SIZE, &sendWin, 1));139+ EXAMPLES_HCCLCHECK(HcclCommSymWinRegister(hcclComm, sendBuf, COMM_BUF_SIZE, &sendWin, 1));
224 140 
225 HcclCommSymWindow recvWin{nullptr};141 HcclCommSymWindow recvWin{nullptr};
226- HCCL_CHECK(HcclCommSymWinRegister(hcclComm, recvBuf, COMM_BUF_SIZE, &recvWin, 1));142+ EXAMPLES_HCCLCHECK(HcclCommSymWinRegister(hcclComm, recvBuf, COMM_BUF_SIZE, &recvWin, 1));
227 143 
228 // 创建team144 // 创建team
229 HcclTeamCreateDesc desc;145 HcclTeamCreateDesc desc;
230- HCCL_CHECK(HcclTeamCreateDescInit(&desc));146+ EXAMPLES_HCCLCHECK(HcclTeamCreateDescInit(&desc));
231 std::vector<uint32_t> worldRankIds;147 std::vector<uint32_t> worldRankIds;
232 for (uint32_t idx = 0; idx < nranks; idx++) {148 for (uint32_t idx = 0; idx < nranks; idx++) {
233 worldRankIds.emplace_back(idx);149 worldRankIds.emplace_back(idx);
@@ -241,35 +157,32 @@ int32_t main(int32_t argc, char* argv[])
241 desc.protocol = COMM_PROTOCOL_UB_CTP;157 desc.protocol = COMM_PROTOCOL_UB_CTP;
242 desc.channelCnt = 1;158 desc.channelCnt = 1;
243 HcommTeamHandle team{nullptr};159 HcommTeamHandle team{nullptr};
244- HCCL_CHECK(HcclTeamCreate(hcclComm, &desc, &team));160+ EXAMPLES_HCCLCHECK(HcclTeamCreate(hcclComm, &desc, &team));
245 161 
246 basic_ring_kernel<<<1U, nullptr, stream>>>(team, sendWin, recvWin, nranks, rank);162 basic_ring_kernel<<<1U, nullptr, stream>>>(team, sendWin, recvWin, nranks, rank);
247- ACL_CHECK(aclrtSynchronizeStream(stream));163+ EXAMPLES_ACLCHECK(aclrtSynchronizeStream(stream));
248 164 
249 // 结果校验165 // 结果校验
250 int32_t prev = (rank - 1 + nranks) % nranks;166 int32_t prev = (rank - 1 + nranks) % nranks;
251 int32_t getOffset = nranks + rank;167 int32_t getOffset = nranks + rank;
252 int8_t* testBuf = nullptr;168 int8_t* testBuf = nullptr;
253- ACL_CHECK(aclrtMallocHost(reinterpret_cast<void**>(&testBuf), COMM_BUF_SIZE));169+ EXAMPLES_ACLCHECK(aclrtMallocHost(reinterpret_cast<void**>(&testBuf), COMM_BUF_SIZE));
254- ACL_CHECK(aclrtMemcpy(testBuf, COMM_BUF_SIZE, recvBuf, COMM_BUF_SIZE, ACL_MEMCPY_DEVICE_TO_HOST));170+ EXAMPLES_ACLCHECK(aclrtMemcpy(testBuf, COMM_BUF_SIZE, recvBuf, COMM_BUF_SIZE, ACL_MEMCPY_DEVICE_TO_HOST));
255- std::cout << "PUT RES rank " << rank << " recvBuf[" << prev << "]=" << static_cast<int32_t>(testBuf[prev])
256- << " expect=" << prev << std::endl;
257- std::cout << "GET RES rank " << rank << " recvBuf[" << getOffset << "]=" << static_cast<int32_t>(testBuf[getOffset])
258- << " expect=" << prev << std::endl;
259- 
260 bool pass = testBuf[prev] == static_cast<int8_t>(prev) && testBuf[getOffset] == static_cast<int8_t>(prev);171 bool pass = testBuf[prev] == static_cast<int8_t>(prev) && testBuf[getOffset] == static_cast<int8_t>(prev);
261- std::cout << "rank " << rank << " test " << (pass ? "pass" : "failed") << std::endl;172+ std::cout << "[rank " << rank << "] basic_ring | PUT recvBuf[" << prev
173+ << "]=" << static_cast<int32_t>(testBuf[prev]) << ", GET recvBuf[" << getOffset
174+ << "]=" << static_cast<int32_t>(testBuf[getOffset]) << " | " << (pass ? "PASS" : "FAIL") << std::endl;
262 175 
263- HCCL_CHECK(HcclTeamDestroy(team));176+ EXAMPLES_HCCLCHECK(HcclTeamDestroy(team));
264- HCCL_CHECK(HcclCommSymWinDeregister(recvWin));177+ EXAMPLES_HCCLCHECK(HcclCommSymWinDeregister(recvWin));
265- HCCL_CHECK(HcclCommSymWinDeregister(sendWin));178+ EXAMPLES_HCCLCHECK(HcclCommSymWinDeregister(sendWin));
266- ACL_CHECK(aclrtFreeHost(testBuf));179+ EXAMPLES_ACLCHECK(aclrtFreeHost(testBuf));
267- HCOMM_CHECK(HcommMemFree(sendBuf));180+ EXAMPLES_HCOMMCHECK(HcommMemFree(sendBuf));
268- HCOMM_CHECK(HcommMemFree(recvBuf));181+ EXAMPLES_HCOMMCHECK(HcommMemFree(recvBuf));
269- HCCL_CHECK(HcclCommDestroy(hcclComm));182+ EXAMPLES_HCCLCHECK(HcclCommDestroy(hcclComm));
270- ACL_CHECK(aclrtDestroyStream(stream));183+ EXAMPLES_ACLCHECK(aclrtDestroyStream(stream));
271- ACL_CHECK(aclrtResetDevice(static_cast<uint32_t>(rank)));184+ EXAMPLES_ACLCHECK(aclrtResetDevice(static_cast<uint32_t>(rank)));
272- ACL_CHECK(aclFinalize());185+ EXAMPLES_ACLCHECK(aclFinalize());
273 186 
274- return 0;187+ return pass ? 0 : 1;
275}188}
Rexamples/hcomm_write_read_nbi/CMakeLists.txt→examples/aicore/hcomm/01_hcomm_write_read_nbi/CMakeLists.txt+54-49
@@ -1,49 +1,54 @@
1-# ----------------------------------------------------------------------------------------------------------1+# ----------------------------------------------------------------------------------------------------------
2-# Copyright (c) 2026 Huawei Technologies Co., Ltd.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 of3+# 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").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.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,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.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.8+# See LICENSE in the root of the software repository for the full text of the License.
9-# ----------------------------------------------------------------------------------------------------------9+# ----------------------------------------------------------------------------------------------------------
10- 10+ 
11-cmake_minimum_required(VERSION 3.16)11+cmake_minimum_required(VERSION 3.16)
12- 12+ 
13-set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")13+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
14- 14+ 
15-if(NOT CMAKE_ASC_ARCHITECTURES STREQUAL "dav-3510")15+find_package(ASC REQUIRED)
16- message(FATAL_ERROR "hcomm_write_read_nbi only supports dav-3510, but got ${CMAKE_ASC_ARCHITECTURES}")16+project(hcomm_write_read_nbi LANGUAGES ASC CXX)
17-endif()17+ 
18- 18+option(ASCENDC_DEBUG "Enable AscendC/Hcomm kernel debug logs" OFF)
19-find_package(ASC REQUIRED)19+ 
20- 20+set_source_files_properties(hcomm_write_read_nbi_kernel.cpp PROPERTIES LANGUAGE ASC)
21-project(kernel_samples LANGUAGES ASC CXX)21+ 
22- 22+add_executable(hcomm_write_read_nbi
23-option(ASCENDC_DEBUG "Enable AscendC/Hcomm kernel debug logs" OFF)23+ hcomm_write_read_nbi.asc
24- 24+ hcomm_write_read_nbi_kernel.cpp
25-set_source_files_properties(25+)
26- hcomm_write_read_nbi_kernel.cpp26+ 
27- utils.cpp27+set_target_properties(hcomm_write_read_nbi PROPERTIES
28- PROPERTIES LANGUAGE ASC28+ LINKER_LANGUAGE ASC
29-)29+)
30- 30+ 
31-add_executable(demo31+target_include_directories(hcomm_write_read_nbi PRIVATE
32- hcomm_write_read_nbi.asc32+ ${CMAKE_CURRENT_SOURCE_DIR}
33- hcomm_write_read_nbi_kernel.cpp33+ ${CMAKE_CURRENT_SOURCE_DIR}/../..
34- utils.cpp34+)
35-)35+ 
36- 36+target_link_directories(hcomm_write_read_nbi PRIVATE
37-target_link_libraries(demo PRIVATE37+ $ENV{ASCEND_HOME_PATH}/lib64
38- hcomm38+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/lib64
39-)39+)
40- 40+ 
41-target_compile_options(demo PRIVATE41+target_link_libraries(hcomm_write_read_nbi PRIVATE
42- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>42+ hcomm
43-)43+ hccl
44- 44+)
45-if(ASCENDC_DEBUG)45+ 
46- target_compile_definitions(demo PRIVATE46+target_compile_options(hcomm_write_read_nbi PRIVATE
47- ASCENDC_DEBUG47+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
48- )48+)
49-endif()49+ 
50+if(ASCENDC_DEBUG)
51+ target_compile_definitions(hcomm_write_read_nbi PRIVATE
52+ ASCENDC_DEBUG
53+ )
54+endif()
Aexamples/aicore/hcomm/01_hcomm_write_read_nbi/README.md+143-0
@@ -0,0 +1,143 @@
1+# Hcomm AIV直驱URMA WriteNbi/ReadNbi点对点通信样例
2+ 
3+## 概述
4+ 
5+本样例展示如何在Ascend C AIV Kernel中，基于**AIV直驱URMA**架构，使用`Hcomm`类的`WriteNbi`和
6+`ReadNbi`接口实现NPU间低时延的点对点（P2P）通信。样例采用多rank环形拓扑：每个rank与前后相邻rank
7+建立P2P通道，通过地址偏移区分数据段，Kernel侧直接提交通信任务，最后在Host侧校验通信结果。
8+ 
9+Host侧只负责创建通信域、注册通信内存和建立通道；通信任务由AIV Kernel直接提交，数据面执行期间
10+不需要Host逐次参与。该模式适用于MoE Dispatch/Combine、Pipeline并行等对通信时延敏感的场景。
11+ 
12+## 本样例支持的产品及CANN软件版本
13+ 
14+| 产品 | CANN软件版本 |
15+|------|-------------|
16+| Ascend 950PR / Ascend 950DT | >= CANN 9.1.0 |
17+ 
18+## 目录结构介绍
19+ 
20+```text
21+hcomm_write_read_nbi
22+├── CMakeLists.txt // 编译工程文件
23+├── hcomm_write_read_nbi.asc // Host侧资源准备与Kernel调用
24+├── hcomm_write_read_nbi_kernel.cpp // AIV Kernel侧Hcomm调用
25+├── hcomm_rw_def.h // Host与Kernel共享定义
26+├── README.md // 中文样例说明
27+└── README_en.md // 英文样例说明
28+```
29+ 
30+## 样例描述
31+ 
32+### 功能说明
33+ 
34+| API | 数据流向 | 语义说明（AIV直驱URMA） |
35+|------|---------|------|
36+| `WriteNbi` | 本地GM → 远端GM | AIV直驱URMA写接口：将本地Global Memory数据直接写入远端NPU指定地址，无需远端CPU介入。 |
37+| `ReadNbi` | 远端GM → 本地GM | AIV直驱URMA读接口：直接从远端NPU指定地址读取数据到本地Global Memory，无需远端CPU介入。 |
38+ 
39+每个rank的通信缓冲区划分为三段，通过地址偏移区分：
40+ 
41+- **seg0** `[0, DATA_SIZE)`：本卡pattern（Host侧以`rank + index`模式预初始化，首字节为rankId，确保不同rank的数据来源可区分）
42+- **seg1** `[DATA_SIZE, 2*DATA_SIZE)`：接收`prev`通过`WriteNbi`写入的数据
43+- **seg2** `[2*DATA_SIZE, 3*DATA_SIZE)`：接收本卡通过`ReadNbi`从`next`读回的数据
44+ 
45+Kernel执行完毕后，Host侧回读结果并校验seg1和seg2是否与对端pattern完全一致。
46+ 
47+### 样例规格
48+ 
49+| 项目 | 说明 |
50+|------|------|
51+| 通信模式 | 环形拓扑AIV直驱URMA点对点通信 |
52+| 调用方式 | AIV Kernel内`Hcomm::WriteNbi`/`ReadNbi` |
53+| 进程启动 | 可执行文件直接拉起，内部fork多rank |
54+| 支持rank数 | >= 2，且不超过当前可用NPU数量 |
55+ 
56+### 实现流程
57+ 
58+1. 可执行文件直接拉起全部rank进程（进程内fork，参见`examples/aicore/utils/process_manager.h`）；
59+ rank间通过命令行指定的TCP控制通道（`examples/aicore/utils/rank_sync.h`）交换root info并
60+ 创建通信域。**注意：AIV直驱模式下，无需配置`hcclOpExpansionMode`。**
61+2. 调用`HcclCommMemReg`向通信域注册本卡通信buffer，Channel创建时内存信息自动交换给对端。
62+3. 通过`HcclRankGraphGetLayers`和`HcclRankGraphGetLinks`获取到`prev`和`next`的链路Endpoint。
63+4. 调用`HcclChannelAcquire`创建P2P通道，指定引擎为`COMM_ENGINE_AIV`、协议为`COMM_PROTOCOL_UB_CTP`（URMA协议）。
64+5. 调用`HcclChannelGetRemoteMems`获取对端注册内存地址，作为Kernel侧`WriteNbi`/`ReadNbi`的远端目标地址。
65+6. Host预初始化seg0，将`ChannelHandle`和buffer地址封装到`CommContext`下发到各卡GM。
66+7. Kernel侧执行`Init()`分配UB工作空间后提交通信任务：将本卡seg0通过`WriteNbi`写入`next`的
67+ seg1，并通过`ReadNbi`从`next`的seg0读到本卡seg2，最后`Drain()`等待通信任务完成。
68+8. Host侧通过控制通道barrier同步后，回读结果并校验seg1和seg2的pattern。
69+ 
70+```cpp
71+// 单个rank上的AIV直驱URMA通信操作示例：
72+// 1. 本卡seg0 → 对端seg1 (WriteNbi)
73+hcomm_.WriteNbi(channel, remoteBuf + DATA_SIZE, localBuf, DATA_SIZE);
74+ 
75+// 2. 对端seg0 → 本卡seg2 (ReadNbi)
76+hcomm_.ReadNbi(channel, localBuf + 2 * DATA_SIZE, remoteBuf, DATA_SIZE);
77+ 
78+// 3. 等待AIV硬件通信任务完成
79+hcomm_.Drain(channel);
80+```
81+ 
82+## 编译运行
83+ 
84+在本样例根目录下执行如下步骤，编译并运行样例。本样例仅支持NPU运行模式。
85+ 
86+- 配置环境变量
87+ 
88+ 请根据当前环境上CANN开发套件包的安装方式配置环境变量。
89+ 
90+ ```bash
91+ source ${install_path}/cann/set_env.sh
92+ ```
93+ 
94+ > **说明：** `${install_path}`为CANN包安装目录，未指定安装目录时默认安装至`/usr/local/Ascend`。
95+ 
96+- 样例执行
97+ 
98+ 在本样例目录下执行如下命令。可执行文件内部fork出全部rank进程，直接运行即可：
99+ 
100+ ```bash
101+ mkdir -p build
102+ cd build
103+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
104+ make -j
105+ cd .. && ./build/hcomm_write_read_nbi tcp://127.0.0.1:29623 <nranks>
106+ ```
107+ 
108+- 启动参数说明
109+ 
110+ | 参数 | 默认值 | 说明 |
111+ |------|--------|------|
112+ | `<ip:port>` | 无默认 | rank 0监听的控制通道地址（`tcp://ip:port`格式），用于root info交换与Host侧barrier；并行运行多个实例时错开端口 |
113+ | `<nranks>` | 无默认 | 启动的rank数量，必须 >= 2，且不超过当前可用NPU数量 |
114+ 
115+- 编译选项说明
116+ 
117+ | 选项 | 可选值 | 说明 |
118+ |------|--------|------|
119+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认） | 运行模式，本样例仅支持NPU运行 |
120+ | `CMAKE_ASC_ARCHITECTURES` | `dav-3510`（默认） | NPU架构：`dav-3510`对应Ascend 950PR/Ascend 950DT |
121+ 
122+- 执行结果
123+ 
124+ 多卡执行成功后，终端将输出如下信息，说明AIV直驱URMA通信成功（写入读出结果一致）：
125+ 
126+ ```text
127+ [rank 0] hcomm_write_read_nbi | sent to rank 1, received from rank 1 | PASS
128+ [rank 1] hcomm_write_read_nbi | sent to rank 0, received from rank 0 | PASS
129+ RESULT | Example=hcomm_write_read_nbi | Status=PASS
130+ ```
131+ 
132+## 注意事项
133+ 
134+- 运行样例需要至少2张NPU；单卡环境仅支持编译验证。
135+- Kernel侧`Init()`分配的UB工作空间需 >= 512字节，用于存放Hcomm内部WQE/CQE等状态。
136+- `WriteNbi`/`ReadNbi`默认`commit=true`，入队后立即触发doorbell；如需批量优化，可改用
137+ `commit=false`多次入队后统一调用`Commit()`。`Drain()`返回后数据可见性才有保证。
138+- 当`nranks == 2`时，`prev`和`next`指向同一个对端rank，seg1和seg2会校验同一个对端pattern。
139+- 当前样例已建立`prev`和`next`两个方向的P2P通道，但默认只执行`next`方向的数据面操作。若要
140+ 扩展为完整的all-to-neighbor通信，可复用已建立的`prev`通道：执行两次Kernel，第一次使用`prev`
141+ 通道执行`ReadNbi`从`prev`的seg0读到本卡seg2，第二次使用`next`通道执行`WriteNbi`将本卡seg0
142+ 写入`next`的seg1，使每个rank的seg1和seg2都匹配`prev`的pattern。扩展时需保留Kernel前后的
143+ Host侧barrier；若调整数据段偏移，需同步更新Host侧pattern校验的期望rank和offset。
Aexamples/aicore/hcomm/01_hcomm_write_read_nbi/README_en.md+132-0
@@ -0,0 +1,132 @@
1+# Hcomm AIV Direct-Drive URMA WriteNbi/ReadNbi Point-to-Point Communication Sample
2+ 
3+## Overview
4+ 
5+This sample demonstrates how to use the `WriteNbi` and `ReadNbi` APIs of the `Hcomm` class to implement low-latency point-to-point (P2P) communication between NPUs from an Ascend C AIV Kernel, based on the **AIV direct-drive URMA** architecture. The sample adopts a multi-rank ring topology: each rank establishes P2P channels with its neighboring ranks, separates data segments via address offsets, submits communication tasks directly from the Kernel side, and validates the communication results on the Host side.
6+ 
7+The Host side is only responsible for creating the communication domain, registering communication memory, and acquiring channels; communication tasks are submitted directly by the AIV Kernel, and the Host does not need to participate in each data-plane transfer. This mode is suitable for latency-sensitive scenarios such as MoE Dispatch/Combine and pipeline parallelism.
8+ 
9+## Supported Products and CANN Software Versions
10+ 
11+| Product | CANN Software Version |
12+|---------|-----------------------|
13+| Ascend 950PR / Ascend 950DT | >= CANN 9.1.0 |
14+ 
15+## Directory Structure
16+ 
17+```text
18+hcomm_write_read_nbi
19+├── CMakeLists.txt // CMake build file
20+├── hcomm_write_read_nbi.asc // Host resource setup and Kernel invocation
21+├── hcomm_write_read_nbi_kernel.cpp // Hcomm calls from the AIV Kernel
22+├── hcomm_rw_def.h // Definitions shared by Host and Kernel
23+├── README.md // Chinese sample documentation
24+└── README_en.md // English sample documentation
25+```
26+ 
27+## Sample Description
28+ 
29+### Functionality Description
30+ 
31+| API | Data Flow | Semantic Description (AIV Direct-Drive URMA) |
32+|-----|-----------|----------------------------------------------|
33+| `WriteNbi` | Local GM → Remote GM | AIV direct-drive URMA write API: directly writes local Global Memory data to a specified address on the remote NPU, without requiring remote CPU intervention. |
34+| `ReadNbi` | Remote GM → Local GM | AIV direct-drive URMA read API: directly reads data from a specified address on the remote NPU into local Global Memory, without requiring remote CPU intervention. |
35+ 
36+Each rank's communication buffer is divided into three segments, distinguished by address offsets:
37+ 
38+- **seg0** `[0, DATA_SIZE)`: Local pattern (pre-initialized by the Host using a `rank + index` pattern, with the first byte set to rankId, ensuring that data sources from different ranks are distinguishable)
39+- **seg1** `[DATA_SIZE, 2*DATA_SIZE)`: Receives data written by `prev` through `WriteNbi`
40+- **seg2** `[2*DATA_SIZE, 3*DATA_SIZE)`: Receives data read back from `next` by the local `ReadNbi`
41+ 
42+After Kernel execution, the Host reads back the results and validates whether seg1 and seg2 match the peer patterns exactly.
43+ 
44+### Sample Specifications
45+ 
46+| Item | Description |
47+|------|-------------|
48+| Communication Mode | Ring topology AIV direct-drive URMA point-to-point communication |
49+| Invocation Method | `Hcomm::WriteNbi`/`ReadNbi` within AIV Kernel |
50+| Process Launch | Executable directly launches, internally forks multiple ranks |
51+| Supported Rank Count | >= 2, not exceeding the number of available NPUs |
52+ 
53+### Implementation Flow
54+ 
55+1. The executable directly launches all rank processes (in-process fork, see `examples/aicore/utils/process_manager.h`); ranks exchange root info and create the communication domain through the TCP control channel specified on the command line (`examples/aicore/utils/rank_sync.h`). **Note: In AIV direct-drive mode, there is no need to configure `hcclOpExpansionMode`.**
56+2. Call `HcclCommMemReg` to register the local communication buffer with the communication domain; memory information is automatically exchanged with the peer during channel creation.
57+3. Use `HcclRankGraphGetLayers` and `HcclRankGraphGetLinks` to obtain link endpoints for `prev` and `next`.
58+4. Call `HcclChannelAcquire` to create P2P channels, specifying `COMM_ENGINE_AIV` as the engine and `COMM_PROTOCOL_UB_CTP` as the protocol (URMA protocol).
59+5. Call `HcclChannelGetRemoteMems` to retrieve the memory addresses registered by peer ranks, which serve as remote target addresses for `WriteNbi`/`ReadNbi` in the Kernel.
60+6. The Host pre-initializes seg0, encapsulates the `ChannelHandle` and buffer addresses into `CommContext`, and downloads it to the GM of each card.
61+7. The Kernel side executes `Init()` to allocate UB workspace, then submits communication tasks: writes local seg0 to `next`'s seg1 through `WriteNbi`, and reads `next`'s seg0 into local seg2 through `ReadNbi`, finally calling `Drain()` to wait for communication tasks to complete.
62+8. The Host synchronizes via control channel barrier, then reads back results and validates seg1 and seg2 patterns.
63+ 
64+```cpp
65+// Example AIV direct-drive URMA operations on one rank:
66+// 1. Local seg0 → Peer seg1 (WriteNbi)
67+hcomm_.WriteNbi(channel, remoteBuf + DATA_SIZE, localBuf, DATA_SIZE);
68+ 
69+// 2. Peer seg0 → Local seg2 (ReadNbi)
70+hcomm_.ReadNbi(channel, localBuf + 2 * DATA_SIZE, remoteBuf, DATA_SIZE);
71+ 
72+// 3. Wait for AIV hardware communication tasks to complete
73+hcomm_.Drain(channel);
74+```
75+ 
76+## Compilation and Execution
77+ 
78+Follow the steps below in the root directory of this sample to compile and run it. This sample only supports NPU execution mode.
79+ 
80+- Configure environment variables
81+ 
82+ Configure the environment variables according to the installation method of the CANN development kit on your current environment.
83+ 
84+ ```bash
85+ source ${install_path}/cann/set_env.sh
86+ ```
87+ 
88+ > **Note:** `${install_path}` is the CANN package installation directory. If not specified, it defaults to `/usr/local/Ascend`.
89+ 
90+- Run the sample
91+ 
92+ Execute the following commands in the sample directory. The executable internally forks all rank processes, so it can be run directly:
93+ 
94+ ```bash
95+ mkdir -p build
96+ cd build
97+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
98+ make -j
99+ cd .. && ./build/hcomm_write_read_nbi tcp://127.0.0.1:29623 <nranks>
100+ ```
101+ 
102+- Launch parameter description
103+ 
104+ | Parameter | Default | Description |
105+ |-----------|---------|-------------|
106+ | `<ip:port>` | None | Control channel address (`tcp://ip:port` format) that rank 0 listens on, used for root info exchange and Host-side barriers; use different ports when running multiple instances in parallel |
107+ | `<nranks>` | None | Number of ranks to launch; must be >= 2 and not exceed the number of available NPUs |
108+ 
109+- Build options description
110+ 
111+ | Option | Allowed Values | Description |
112+ |--------|----------------|-------------|
113+ | `CMAKE_ASC_RUN_MODE` | `npu` (default) | Execution mode; this sample only supports NPU execution |
114+ | `CMAKE_ASC_ARCHITECTURES` | `dav-3510` (default) | NPU architecture: `dav-3510` corresponds to Ascend 950PR / Ascend 950DT |
115+ 
116+- Expected output
117+ 
118+ After a successful multi-card run, the terminal will output the following, indicating that AIV direct-drive URMA communication completed successfully (write/read results are consistent):
119+ 
120+ ```text
121+ [rank 0] hcomm_write_read_nbi | sent to rank 1, received from rank 1 | PASS
122+ [rank 1] hcomm_write_read_nbi | sent to rank 0, received from rank 0 | PASS
123+ RESULT | Example=hcomm_write_read_nbi | Status=PASS
124+ ```
125+ 
126+## Notes
127+ 
128+- Running the sample requires at least 2 NPUs; a single-card environment only supports compilation verification.
129+- The UB workspace allocated by `Init()` on the Kernel side must be >= 512 bytes, used to store internal Hcomm states such as WQE/CQE.
130+- `WriteNbi`/`ReadNbi` defaults to `commit=true`, triggering the doorbell immediately after enqueueing; for batch optimization, use `commit=false` to enqueue multiple tasks and call `Commit()` uniformly. Data visibility is only guaranteed after `Drain()` returns.
131+- When `nranks == 2`, `prev` and `next` refer to the same peer rank, so seg1 and seg2 are validated against the same peer pattern.
132+- The current sample has established P2P channels in both the `prev` and `next` directions, but only executes data-plane operations in the `next` direction by default. To extend to full all-to-neighbor communication, reuse the established `prev` channel: launch the Kernel twice, first using the `prev` channel to execute `ReadNbi` reading `prev`'s seg0 into local seg2, then using the `next` channel to execute `WriteNbi` writing local seg0 into `next`'s seg1, so that each rank's seg1 and seg2 both match `prev`'s pattern. When extending, keep the Host-side barriers before and after Kernel execution; if data segment offsets are adjusted, update the expected rank and offset used by Host-side pattern validation accordingly.
Rexamples/hcomm_write_read_nbi/hcomm_rw_def.h→examples/aicore/hcomm/01_hcomm_write_read_nbi/hcomm_rw_def.h+67-66
@@ -1,66 +1,67 @@
1-/**1+/**
2- * Copyright (c) 2026 Huawei Technologies Co., Ltd.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 of3+ * 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").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.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,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.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.8+ * See LICENSE in the root of the software repository for the full text of the License.
9- */9+ */
10- 10+ 
11-#ifndef HCOMM_RW_DEF_H11+#ifndef HCOMM_RW_DEF_H
12-#define HCOMM_RW_DEF_H12+#define HCOMM_RW_DEF_H
13- 13+ 
14-#include <cstdint>14+#include <cstdint>
15- 15+ 
16-#include "adv_api/hcomm/hcomm.h"16+#include "adv_api/hcomm/hcomm.h"
17-#include "hccl/hccl.h"17+#include "hccl/hccl.h"
18-#include "hccl/hccl_comm.h"18+ 
19- 19+// Host侧流程的返回值与错误检查宏（原TCP环工具utils.h中仅剩的被使用部分）。
20-constexpr uint32_t DATA_SIZE = 256U;20+constexpr int32_t SUCCESS = 0;
21-constexpr uint64_t COMM_BUF_SIZE = 4096U;21+constexpr int32_t FAIL = -1;
22-constexpr uint64_t SEND_DATA_OFFSET = 0U;22+ 
23-constexpr uint64_t WRITE_RESULT_OFFSET = DATA_SIZE;23+constexpr uint32_t DATA_SIZE = 256U;
24-constexpr uint64_t READ_RESULT_OFFSET = 2U * DATA_SIZE;24+constexpr uint64_t COMM_BUF_SIZE = 4096U;
25- 25+constexpr uint64_t SEND_DATA_OFFSET = 0U;
26-constexpr uint32_t NRANKS = 2U;26+constexpr uint64_t WRITE_RESULT_OFFSET = DATA_SIZE;
27-constexpr uint16_t BASE_PORT = 29620U;27+constexpr uint64_t READ_RESULT_OFFSET = 2U * DATA_SIZE;
28-constexpr uint32_t ROOT_SERVER_RANK = 0U;28+ 
29-constexpr CommProtocol HOST_COMM_PROTOCOL = COMM_PROTOCOL_UB_CTP;29+constexpr uint32_t NRANKS = 2U;
30-constexpr AscendC::CommProtocol KERNEL_COMM_PROTOCOL = AscendC::COMM_PROTOCOL_UB_CTP;30+constexpr CommProtocol HOST_COMM_PROTOCOL = COMM_PROTOCOL_UB_CTP;
31-static_assert(31+constexpr AscendC::CommProtocol KERNEL_COMM_PROTOCOL = AscendC::COMM_PROTOCOL_UB_CTP;
32- static_cast<int32_t>(HOST_COMM_PROTOCOL) == static_cast<int32_t>(KERNEL_COMM_PROTOCOL),32+static_assert(
33- "Host and Kernel communication protocols must match");33+ static_cast<int32_t>(HOST_COMM_PROTOCOL) == static_cast<int32_t>(KERNEL_COMM_PROTOCOL),
34-constexpr uint32_t HCOMM_WORKSPACE_SIZE = 512U;34+ "Host and Kernel communication protocols must match");
35- 35+constexpr uint32_t HCOMM_WORKSPACE_SIZE = 512U;
36-enum TestResult : uint32_t {36+ 
37- TEST_SUCCESS = 0U,37+enum TestResult : uint32_t {
38- TEST_HCOMM_INIT_FAILED = 1U,38+ TEST_SUCCESS = 0U,
39- TEST_WRITE_FAILED = 2U,39+ TEST_HCOMM_INIT_FAILED = 1U,
40- TEST_READ_FAILED = 3U,40+ TEST_WRITE_FAILED = 2U,
41- TEST_DRAIN_FAILED = 4U,41+ TEST_READ_FAILED = 3U,
42-};42+ TEST_DRAIN_FAILED = 4U,
43-static_assert(sizeof(TestResult) == sizeof(uint32_t), "TestResult must remain uint32_t-sized");43+};
44- 44+static_assert(sizeof(TestResult) == sizeof(uint32_t), "TestResult must remain uint32_t-sized");
45-constexpr uint32_t HCOMM_READ_ONLY = 0x01U;45+ 
46-constexpr uint32_t HCOMM_WRITE_ONLY = 0x02U;46+constexpr uint32_t HCOMM_READ_ONLY = 0x01U;
47-constexpr uint32_t HCOMM_READ_WRITE = HCOMM_READ_ONLY | HCOMM_WRITE_ONLY;47+constexpr uint32_t HCOMM_WRITE_ONLY = 0x02U;
48- 48+constexpr uint32_t HCOMM_READ_WRITE = HCOMM_READ_ONLY | HCOMM_WRITE_ONLY;
49-namespace HcommExample {49+ 
50- 50+namespace HcommExample {
51-struct CommContext {51+ 
52- uint64_t channelHandle;52+struct CommContext {
53- uint64_t localBufferAddr;53+ uint64_t channelHandle;
54- uint64_t remoteBufferAddr;54+ uint64_t localBufferAddr;
55- TestResult testResult;55+ uint64_t remoteBufferAddr;
56-};56+ TestResult testResult;
57- 57+};
58-} // namespace HcommExample58+ 
59- 59+} // namespace HcommExample
60-extern "C" __vector__ __global__ __aicore__ void kernel_hcomm_read_nbi(GM_ADDR context);60+ 
61- 61+extern "C" __vector__ __global__ __aicore__ void kernel_hcomm_read_nbi(GM_ADDR context);
62-extern "C" __vector__ __global__ __aicore__ void kernel_hcomm_write_nbi(GM_ADDR context);62+ 
63- 63+extern "C" __vector__ __global__ __aicore__ void kernel_hcomm_write_nbi(GM_ADDR context);
64-extern "C" __vector__ __global__ __aicore__ void kernel_hcomm_write_read_nbi(GM_ADDR context);64+ 
65- 65+extern "C" __vector__ __global__ __aicore__ void kernel_hcomm_write_read_nbi(GM_ADDR context);
66-#endif // HCOMM_RW_DEF_H66+ 
67+#endif // HCOMM_RW_DEF_H
Rexamples/hcomm_write_read_nbi/hcomm_write_read_nbi.asc→examples/aicore/hcomm/01_hcomm_write_read_nbi/hcomm_write_read_nbi.asc+445-504
Rexamples/hcomm_write_read_nbi/hcomm_write_read_nbi_kernel.cpp→examples/aicore/hcomm/01_hcomm_write_read_nbi/hcomm_write_read_nbi_kernel.cpp+0-0
文件重命名但无更改。
Rexamples/hcomm_batch_write/CMakeLists.txt→examples/aicore/hcomm/02_hcomm_batch_write/CMakeLists.txt+14-29
@@ -11,52 +11,37 @@
11cmake_minimum_required(VERSION 3.16)11cmake_minimum_required(VERSION 3.16)
12 12 
13set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")13set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
14-if(NOT CMAKE_ASC_ARCHITECTURES STREQUAL "dav-3510")
15- message(FATAL_ERROR "hcomm_batch_write only supports dav-3510, but got ${CMAKE_ASC_ARCHITECTURES}")
16-endif()
17 14 
18find_package(ASC REQUIRED)15find_package(ASC REQUIRED)
19project(hcomm_batch_write LANGUAGES ASC CXX)16project(hcomm_batch_write LANGUAGES ASC CXX)
20 17 
21-set(CMAKE_CXX_STANDARD 17)18+set_source_files_properties(hcomm_batch_write_kernel.cpp PROPERTIES LANGUAGE ASC)
22-set(CMAKE_CXX_STANDARD_REQUIRED ON)
23- 
24-if(NOT ASCEND_CANN_PACKAGE_LINUX_PATH)
25- if(CMAKE_HOST_SYSTEM_PROCESSOR STREQUAL "x86_64")
26- set(ASCEND_CANN_PACKAGE_LINUX_PATH "${ASCEND_CANN_PACKAGE_PATH}/x86_64-linux")
27- else()
28- set(ASCEND_CANN_PACKAGE_LINUX_PATH "${ASCEND_CANN_PACKAGE_PATH}/aarch64-linux")
29- endif()
30-endif()
31- 
32-set_source_files_properties(
33- hcomm_batch_write_kernel.cpp
34- PROPERTIES LANGUAGE ASC
35-)
36 19 
37add_executable(hcomm_batch_write20add_executable(hcomm_batch_write
38 hcomm_batch_write.asc21 hcomm_batch_write.asc
39 hcomm_batch_write_kernel.cpp22 hcomm_batch_write_kernel.cpp
40)23)
41 24 
25+set_target_properties(hcomm_batch_write PROPERTIES
26+ LINKER_LANGUAGE ASC
27+)
28+ 
42target_include_directories(hcomm_batch_write PRIVATE29target_include_directories(hcomm_batch_write PRIVATE
43- "${ASCEND_CANN_PACKAGE_LINUX_PATH}/pkg_inc"30+ ${CMAKE_CURRENT_SOURCE_DIR}
44- "${ASCEND_CANN_PACKAGE_LINUX_PATH}/include"31+ ${CMAKE_CURRENT_SOURCE_DIR}/../..
45- "${ASCEND_CANN_PACKAGE_LINUX_PATH}/include/hcomm"32+ $ENV{ASCEND_HOME_PATH}/asc/include/comm_api
46- "${ASCEND_CANN_PACKAGE_LINUX_PATH}/asc/include/adv_api"
47)33)
48 34 
49target_link_directories(hcomm_batch_write PRIVATE35target_link_directories(hcomm_batch_write PRIVATE
50- "${ASCEND_CANN_PACKAGE_PATH}/lib64"36+ $ENV{ASCEND_HOME_PATH}/lib64
51- "${ASCEND_CANN_PACKAGE_LINUX_PATH}/lib64"37+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/lib64
52)38)
53 39 
54-target_link_libraries(hcomm_batch_write PRIVATE hcomm)40+target_link_libraries(hcomm_batch_write PRIVATE
41+ hcomm
42+ hccl
43+)
55 44 
56target_compile_options(hcomm_batch_write PRIVATE45target_compile_options(hcomm_batch_write PRIVATE
57 $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>46 $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
58)47)
59- 
60-set_target_properties(hcomm_batch_write PROPERTIES
61- BUILD_RPATH "${ASCEND_CANN_PACKAGE_PATH}/lib64;${ASCEND_CANN_PACKAGE_LINUX_PATH}/lib64"
62-)
Aexamples/aicore/hcomm/02_hcomm_batch_write/README.md+112-0
@@ -0,0 +1,112 @@
1+# Hcomm Batch Write样例
2+ 
3+## 概述
4+ 
5+本样例演示如何使用Hcomm批量写接口，在单机多卡环境中由rank 0向其他rank批量写入数据。多个URMA
6+Channel共享同一个Jetty，通过`MakeBatchHandle`、`GetHandleRef`、`BatchCommit`和`Drain`批量提交
7+跨peer写任务。程序会准备待发送数据，提交一批写操作并等待完成；接收rank随后校验收到的数据，主进程
8+最后汇总所有rank的执行结果。
9+ 
10+## 本样例支持的产品及CANN软件版本
11+ 
12+| 产品 | CANN软件版本 |
13+| --- | --- |
14+| Ascend 950PR/Ascend 950DT | >= CANN 9.1.0 |
15+ 
16+## 目录结构介绍
17+ 
18+```text
19+hcomm_batch_write
20+├── CMakeLists.txt // 编译工程文件
21+├── hcomm_batch_write.asc // Host侧资源准备与Kernel调用
22+├── hcomm_batch_write_kernel.cpp // AIV Kernel侧批量写调用
23+├── hcomm_batch_write_def.h // Host与Kernel共享定义
24+├── README.md // 中文样例说明
25+└── README_en.md // 英文样例说明
26+```
27+ 
28+## 样例描述
29+ 
30+### 功能说明
31+ 
32+rank 0将一段本端数据通过批量写接口同时写入其他所有rank的接收缓冲区。批量提交将多个写任务打包为
33+一个batch handle，通过一次`BatchCommit`统一下发，减少提交次数；`Drain`等待整批任务完成后，各
34+接收rank校验落地的数据是否与本端发送数据一致。
35+ 
36+### 样例规格
37+ 
38+| 项目 | 说明 |
39+| --- | --- |
40+| 通信模式 | rank 0到其他所有rank的一对多批量写 |
41+| 调用方式 | AIV Kernel内`MakeBatchHandle`/`GetHandleRef`/`BatchCommit`/`Drain` |
W
Wwuzhf99月16日

当前样例仅覆盖 MultiChannelHandle + GetHandleRef 的多通道批量写流程。即使配置为 2 个 rank，也只是 MultiChannelHandle 中包含一个逻辑通道，尚未覆盖基于 ChannelHandle 创建 BatchHandle 的单通道接口用法。建议增加单通道调用分支或独立样例，演示MakeBatchHandle(ChannelHandle, ...)、WriteNbi、BatchCommit 和 Drain 的完整流程。

likedislike
wanghan
wanghan
9月16日 评论：
42+| 数据校验 | 每个接收rank校验落地数据 |
43+| 支持rank数 | 2至16，且不超过本机可用NPU数量 |
44+ 
45+### 实现流程
46+ 
47+1. 可执行文件直接拉起全部rank进程（进程内fork，参见`examples/aicore/utils/process_manager.h`）；
48+ rank间通过命令行指定的TCP控制通道（`examples/aicore/utils/rank_sync.h`）交换root info并
49+ 创建通信域，rank 0准备待发送数据。
50+2. 各rank注册通信内存并通过控制通道交换缓冲区地址，rank 0与每个peer建立共享Jetty的URMA Channel。
51+3. Kernel侧构造batch handle，为每个peer填充一个写任务引用后统一`BatchCommit`。
52+4. `Drain`等待整批任务完成。
53+5. 各接收rank校验落地数据，主进程汇总所有rank的执行结果。
54+ 
55+## 编译运行
56+ 
57+在本样例根目录下执行如下步骤，编译并运行样例。本样例仅支持NPU运行模式。
58+ 
59+- 配置环境变量
60+ 
61+ 请根据当前环境上CANN开发套件包的安装方式配置环境变量。
62+ 
63+ ```bash
64+ source ${install_path}/cann/set_env.sh
65+ ```
66+ 
67+ > **说明：** `${install_path}`为CANN包安装目录，未指定安装目录时默认安装至`/usr/local/Ascend`。
68+ 
69+- 样例执行
70+ 
71+ 在本样例目录下执行如下命令。可执行文件内部fork出全部rank进程，直接运行即可。
72+ 
73+ ```bash
74+ mkdir -p build
75+ cd build
76+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
77+ make -j
78+ cd .. && ./build/hcomm_batch_write tcp://127.0.0.1:29624 <rank_num>
79+ ```
80+ 
81+- 启动参数说明
82+ 
83+ | 参数 | 默认值 | 说明 |
84+ | --- | --- | --- |
85+ | `<ip:port>` | 无默认 | rank 0监听的控制通道地址（`tcp://ip:port`格式），用于root info交换与Host侧barrier；并行运行多个实例时错开端口 |
86+ | `<rank_num>` | 无默认 | 启动的rank数量，取值范围`2-16`，不能超过本机可用NPU数量 |
87+ 
88+- 编译选项说明
89+ 
90+ | 选项 | 可选值 | 说明 |
91+ | --- | --- | --- |
92+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认） | 运行模式，本样例仅支持NPU运行 |
93+ | `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU架构，对应Ascend 950PR/Ascend 950DT |
94+ 
95+- 执行结果
96+ 
97+ 4卡运行成功时，每个接收rank输出校验通过信息，主进程最终输出：
98+ 
99+ ```text
100+ [rank 1] hcomm_batch_write | received and verified data from rank 0 | PASS
101+ [rank 2] hcomm_batch_write | received and verified data from rank 0 | PASS
102+ [rank 3] hcomm_batch_write | received and verified data from rank 0 | PASS
103+ RESULT | Example=hcomm_batch_write | Status=PASS
104+ ```
105+ 
106+ 主进程输出`Status=FAIL`或任一rank报错时，表示本次批量写入或数据校验未通过。
107+ 
108+## 注意事项
109+ 
110+- 运行样例需要至少2张NPU；单卡环境仅支持编译验证。
111+- rank数量不能超过本机可用NPU数量，且运行用户需具备设备访问权限。
112+- 编译依赖CANN ASC CMake能力，并在链接阶段依赖CANN `hcomm`库。
Rexamples/hcomm_batch_write/README_en.md→examples/aicore/hcomm/02_hcomm_batch_write/README_en.md+10-13
@@ -23,17 +23,11 @@ source ${install_path}/cann/set_env.sh
23Build and run with two NPUs by default, or pass a rank count from `2` to `16`:23Build and run with two NPUs by default, or pass a rank count from `2` to `16`:
24 24 
25```bash25```bash
26-bash examples/hcomm_batch_write/run.sh26+mkdir -p build
27-bash examples/hcomm_batch_write/run.sh 427+cd build
28-```28+cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
29- 29+make -j
30-To build and run manually:30+cd .. && ./build/hcomm_batch_write tcp://127.0.0.1:29624 4
31- 
32-```bash
33-cmake -S examples/hcomm_batch_write -B build/examples/hcomm_batch_write \
34- -DCMAKE_ASC_ARCHITECTURES=dav-3510
35-cmake --build build/examples/hcomm_batch_write -j
36-./build/examples/hcomm_batch_write/hcomm_batch_write 4
37```31```
38 32 
39The final command starts four ranks. Before running, make sure the requested number of NPUs is available and that the current user can access them.33The final command starts four ranks. Before running, make sure the requested number of NPUs is available and that the current user can access them.
@@ -43,7 +37,10 @@ The final command starts four ranks. Before running, make sure the requested num
43Each receiving rank prints a verification message when its data is correct. A successful four-rank run ends with:37Each receiving rank prints a verification message when its data is correct. A successful four-rank run ends with:
44 38 
45```text39```text
46-hcomm 4-rank grouped batch write test passed40+[rank 1] hcomm_batch_write | received and verified data from rank 0 | PASS
41+[rank 2] hcomm_batch_write | received and verified data from rank 0 | PASS
42+[rank 3] hcomm_batch_write | received and verified data from rank 0 | PASS
43+RESULT | Example=hcomm_batch_write | Status=PASS
47```44```
48 45 
49-`failed` in the parent-process summary or an error from any rank means that the batch write or data verification did not complete successfully.46+`Status=FAIL` in the parent-process summary or an error from any rank means that the batch write or data verification did not complete successfully.
Rexamples/hcomm_batch_write/hcomm_batch_write.asc→examples/aicore/hcomm/02_hcomm_batch_write/hcomm_batch_write.asc+82-258
@@ -14,14 +14,9 @@
14#include <cstring>14#include <cstring>
15#include <signal.h>15#include <signal.h>
16#include <string>16#include <string>
17-#include <sys/socket.h>
18-#include <sys/time.h>
19-#include <sys/wait.h>
20-#include <unistd.h>
21#include <vector>17#include <vector>
22 18 
23#include "acl/acl.h"19#include "acl/acl.h"
24-#include "acl/acl_rt.h"
25#include "hccl/hccl_res.h"20#include "hccl/hccl_res.h"
26#include "hccl/hccl.h"21#include "hccl/hccl.h"
27#include "hccl/hccl_comm.h"22#include "hccl/hccl_comm.h"
Rexamples/hcomm_batch_write/hcomm_batch_write_def.h→examples/aicore/hcomm/02_hcomm_batch_write/hcomm_batch_write_def.h+0-0
文件重命名但无更改。
Rexamples/hcomm_batch_write/hcomm_batch_write_kernel.cpp→examples/aicore/hcomm/02_hcomm_batch_write/hcomm_batch_write_kernel.cpp+0-0
文件重命名但无更改。
Rexamples/one_multi_path/CMakeLists.txt→examples/aicore/hcomm/03_one_multi_path/CMakeLists.txt+16-17
@@ -10,24 +10,11 @@
10 10 
11cmake_minimum_required(VERSION 3.16)11cmake_minimum_required(VERSION 3.16)
12 12 
13-set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, sim")
14set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")13set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
15 14 
16find_package(ASC REQUIRED)15find_package(ASC REQUIRED)
17- 
18project(one_multi_path LANGUAGES ASC CXX)16project(one_multi_path LANGUAGES ASC CXX)
19 17 
20-set(CMAKE_CXX_STANDARD 17)
21-set(CMAKE_CXX_STANDARD_REQUIRED ON)
22- 
23-if(NOT DEFINED ASCEND_HOME_PATH)
24- if(DEFINED ENV{ASCEND_HOME_PATH})
25- set(ASCEND_HOME_PATH $ENV{ASCEND_HOME_PATH})
26- else()
27- set(ASCEND_HOME_PATH /usr/local/Ascend/ascend-toolkit/latest)
28- endif()
29-endif()
30- 
31add_executable(one_multi_path18add_executable(one_multi_path
32 main.cpp19 main.cpp
33 ubmem_data_copy.asc20 ubmem_data_copy.asc
@@ -38,15 +25,27 @@ set_target_properties(one_multi_path PROPERTIES
38)25)
39 26 
40target_include_directories(one_multi_path PRIVATE27target_include_directories(one_multi_path PRIVATE
41- ${ASCEND_HOME_PATH}/include28+ ${CMAKE_CURRENT_SOURCE_DIR}
29+ ${CMAKE_CURRENT_SOURCE_DIR}/../..
30+ $ENV{ASCEND_HOME_PATH}/include
42)31)
43 32 
44target_link_directories(one_multi_path PRIVATE33target_link_directories(one_multi_path PRIVATE
45- ${ASCEND_HOME_PATH}/lib6434+ $ENV{ASCEND_HOME_PATH}/lib64
46- ${ASCEND_HOME_PATH}/runtime/lib6435+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/lib64
47)36)
48 37 
49-target_link_libraries(one_multi_path PRIVATE hccl hcomm ascendcl c_sec pthread)38+target_link_libraries(one_multi_path PRIVATE
39+ hcomm
40+ hccl
41+ ascendcl
42+ acl_rt
43+ runtime
44+ c_sec
45+ stdc++
46+ dl
47+ pthread
48+)
50 49 
51target_compile_options(one_multi_path PRIVATE50target_compile_options(one_multi_path PRIVATE
52 $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>51 $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
Rexamples/one_multi_path/README.md→examples/aicore/hcomm/03_one_multi_path/README.md+71-42
@@ -2,11 +2,13 @@
2 2 
3## 概述3## 概述
4 4 
5-本样例演示如何通过HCCL Resource API查询`COMM_PROTOCOL_UB_MEM`链路，在同一个通信域中为每个peer创建`pathMode=1`和`pathMode=2`两条Channel。各peer按顺序处理；处理单个peer时，通过两个Stream并发搬运5+本样例演示如何通过HCCL Resource API查询`COMM_PROTOCOL_UB_MEM`链路，在同一个通信域中为每个peer
6-不同的数据分片。6+创建`pathMode=1`和`pathMode=2`两条Channel。各peer按顺序处理；处理单个peer时，通过两个Stream
7+并发搬运不同的数据分片。
7 8 
8-样例采用一进程一Device的运行方式。每个rank使用自己的rank ID填充本端HCCL Buffer，再从其他rank读取9+样例采用一进程一Device的运行方式。每个rank使用自己的rank ID填充本端HCCL Buffer，再从其他rank
9-4 KB数据并校验。`run.sh`负责启动本机rank进程，rank 0进程负责分发`HcclRootInfo`。10+读取4 KB数据并校验。可执行文件直接拉起本机全部rank进程；多机部署时每台机器各起一个进程，
11+rank 0进程负责分发`HcclRootInfo`。
10 12 
11## 本样例支持的产品及CANN软件版本13## 本样例支持的产品及CANN软件版本
12 14 
@@ -18,20 +20,20 @@
18 20 
19```text21```text
20one_multi_path22one_multi_path
21-├── CMakeLists.txt // 编译工程文件23+├── CMakeLists.txt // 编译工程文件
22-├── main.cpp // Host侧通信域、Channel和校验流程24+├── main.cpp // Host侧通信域、Channel和校验流程
23-├── README.md // 中文样例说明25+├── ubmem_data_copy.asc // Ascend C数据搬运Kernel
24-├── README_en.md // 英文样例说明26+├── README.md // 中文样例说明
25-├── run.sh // 多rank启动脚本27+└── README_en.md // 英文样例说明
26-└── ubmem_data_copy.asc // Ascend C数据搬运Kernel
27```28```
28 29 
29## 样例描述30## 样例描述
30 31 
31### 功能说明32### 功能说明
32 33 
33-每个rank先获取本端HCCL Buffer，并使用自己的rank ID填充前4 KB。Channel创建完成后，每个rank通过每个34+每个rank先获取本端HCCL Buffer，并使用自己的rank ID填充前4 KB。Channel创建完成后，每个rank
34-peer对应的两条Channel读取对端HCCL Buffer，将读取结果写入本端校验Buffer，最后拷贝到Host侧逐字节校验。35+通过每个peer对应的两条Channel读取对端HCCL Buffer，将读取结果写入本端校验Buffer，最后拷贝到
36+Host侧逐字节校验。
35 37 
36```text38```text
37对端HCCL Buffer -> UB_MEM链路 -> Ascend C DataCopy -> 本端校验Buffer -> Host校验39对端HCCL Buffer -> UB_MEM链路 -> Ascend C DataCopy -> 本端校验Buffer -> Host校验
@@ -44,19 +46,28 @@ peer对应的两条Channel读取对端HCCL Buffer，将读取结果写入本端
44| 1 | `COMM_PROTOCOL_UB_MEM` | `1` | one path | 前2 KB |46| 1 | `COMM_PROTOCOL_UB_MEM` | `1` | one path | 前2 KB |
45| 2 | `COMM_PROTOCOL_UB_MEM` | `2` | multi path | 后2 KB |47| 2 | `COMM_PROTOCOL_UB_MEM` | `2` | multi path | 后2 KB |
46 48 
49+### 样例规格
50+ 
51+| 项目 | 说明 |
52+| --- | --- |
53+| 通信模式 | 每peer两条`COMM_PROTOCOL_UB_MEM`路径（one path/multi path）并发读取 |
54+| 调用方式 | Host侧Resource API创建Channel + Ascend C DataCopy Kernel |
55+| 数据规模 | 每peer读取4 KB（两条path各2 KB） |
56+| 部署形态 | 单机多卡/多机多卡 |
57+ 
47### 实现流程58### 实现流程
48 59 
49-1. rank 0生成并分发一份`HcclRootInfo`，各rank初始化同一个通信域。60+1. 可执行文件直接拉起本机rank进程（进程内fork，参见`examples/aicore/utils/process_manager.h`）；
61+ rank间通过命令行指定的TCP控制通道（`examples/aicore/utils/rank_sync.h`）交换`HcclRootInfo`，
62+ 各rank初始化同一个通信域。
502. 每个rank获取本端HCCL Buffer并写入rank ID，再通过`HcclCommMemReg`注册Channel所需的内存资源。632. 每个rank获取本端HCCL Buffer并写入rank ID，再通过`HcclCommMemReg`注册Channel所需的内存资源。
513. 每个rank遍历RankGraph各层，为每个peer筛选一条`COMM_PROTOCOL_UB_MEM`链路。643. 每个rank遍历RankGraph各层，为每个peer筛选一条`COMM_PROTOCOL_UB_MEM`链路。
524. 每个peer基于同一组Endpoint构造`pathMode=1`和`pathMode=2`两个描述。654. 每个peer基于同一组Endpoint构造`pathMode=1`和`pathMode=2`两个描述。
535. 将所有peer的Channel描述一次性传入`HcclChannelAcquire`，批量创建Channel。665. 将所有peer的Channel描述一次性传入`HcclChannelAcquire`，批量创建Channel。
54-6. 各peer按顺序处理；单个peer的两条Channel分别在两个Stream上读取前后两个2 KB分片，两个Kernel均下发后67+6. 各peer按顺序处理；单个peer的两条Channel分别在两个Stream上读取前后两个2 KB分片，两个Kernel
55- 再同步Stream，因此并发范围仅限同一peer的两条path。68+ 均下发后再同步Stream，因此并发范围仅限同一peer的两条path。
567. 将本端校验Buffer拷贝到Host，检查每个字节是否等于对端rank ID。697. 将本端校验Buffer拷贝到Host，检查每个字节是否等于对端rank ID。
57-8. 所有peer的数据校验完成后执行`HcclBarrier`，确保全部rank完成远端读取后再销毁通信资源。70+8. 所有peer的数据校验完成后通过控制通道barrier同步，确保全部rank完成远端读取后再销毁通信资源。
58- 
59-`rank_graph_topo`日志只显示所选UB_MEM链路所在的RankGraph层，不参与`pathMode`选择。
60 71 
61## 编译运行72## 编译运行
62 73 
@@ -64,8 +75,7 @@ peer对应的两条Channel读取对端HCCL Buffer，将读取结果写入本端
64 75 
65- 配置环境变量76- 配置环境变量
66 77 
67- 请根据当前环境上CANN开发套件包的[安装方式](../../docs/quick_start.md)配置78+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../docs/quick_start.md)配置环境变量。
68- 环境变量。
69 79 
70 ```bash80 ```bash
71 source ${install_path}/cann/set_env.sh81 source ${install_path}/cann/set_env.sh
@@ -75,38 +85,55 @@ peer对应的两条Channel读取对端HCCL Buffer，将读取结果写入本端
75 85 
76- 单机两卡运行86- 单机两卡运行
77 87 
78- 默认启动两个进程，分别使用Device 0和Device 1。88+ 可执行文件内部fork出全部rank进程，Device与rank一一对应：
79 89 
80 ```bash90 ```bash
81- cmake -S . -B build -DCMAKE_ASC_ARCHITECTURES=dav-351091+ mkdir -p build
82- cmake --build build -j92+ cd build
83- bash run.sh93+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
94+ make -j
95+ cd .. && ./build/one_multi_path tcp://127.0.0.1:29625 <nranks>
84 ```96 ```
85 97 
86- 启动参数说明98- 启动参数说明
87 99 
100+ 单机模式：
101+ 
102+ ```bash
103+ ./build/one_multi_path <ip:port> <nranks>
104+ ```
105+ 
88 | 参数 | 默认值 | 说明 |106 | 参数 | 默认值 | 说明 |
89 | --- | --- | --- |107 | --- | --- | --- |
90- | `-pes` | `2` | 通信域中的总rank数，必须为不小于2的整数 |108+ | `<ip:port>` | 无默认 | rank 0监听的控制通道地址（`tcp://ip:port`格式），用于root info交换与Host侧barrier；并行运行多个实例时错开端口 |
91- | `-ipport` | `tcp://127.0.0.1:8899` | rank 0分发`HcclRootInfo`使用的地址；IPv6地址使用`tcp://[地址]:端口`格式 |109+ | `<nranks>` | 无默认 | 通信域中的总rank数，必须为不小于2的整数，且不超过本机可用NPU数量 |
92- | `-gnpus` | `2` | 当前机器启动的rank进程数 |
93- | `-fpe` | `0` | 当前机器的起始rank ID |
94- | `-fnpu` | `0` | 当前机器的起始Device ID |
95 110 
96- 跨机运行111- 跨机运行
97 112 
98- 两机各4卡、总共8卡时，先在rank 0所在机器执行第一条命令，再在另一台机器执行第二条命令：113+ 多机部署时每台机器直接运行可执行文件，逐rank指定全局rank与Device：
114+ 
115+ ```bash
116+ ./build/one_multi_path <ip:port> <rank_size> <rank> <device>
117+ ```
118+ 
119+ 两机各4卡、总共8卡时，先在rank 0所在机器执行第一条命令，再在另一台机器执行：
99 120 
100 ```bash121 ```bash
101 # 机器A：rank 0-3，Device 0-3122 # 机器A：rank 0-3，Device 0-3
102- bash run.sh -pes 8 -ipport tcp://<机器A_IP>:8899 -gnpus 4 -fpe 0 -fnpu 0123+ ./build/one_multi_path tcp://<机器A_IP>:8899 8 0 0
124+ ./build/one_multi_path tcp://<机器A_IP>:8899 8 1 1
125+ ./build/one_multi_path tcp://<机器A_IP>:8899 8 2 2
126+ ./build/one_multi_path tcp://<机器A_IP>:8899 8 3 3
103 127 
104 # 机器B：rank 4-7，Device 0-3128 # 机器B：rank 4-7，Device 0-3
105- bash run.sh -pes 8 -ipport tcp://<机器A_IP>:8899 -gnpus 4 -fpe 4 -fnpu 0129+ ./build/one_multi_path tcp://<机器A_IP>:8899 8 4 0
130+ ./build/one_multi_path tcp://<机器A_IP>:8899 8 5 1
131+ ./build/one_multi_path tcp://<机器A_IP>:8899 8 6 2
132+ ./build/one_multi_path tcp://<机器A_IP>:8899 8 7 3
106 ```133 ```
107 134 
108- 两机各8卡、总共16卡时，将`-pes`和`-gnpus`分别设置为`16`和`8`，两台机器的`-fpe`分别设置为135+ 其中`<ip:port>`统一指向rank 0所在机器；`<rank_size>`为全局rank总数，`<rank>`为本进程的全局
109- `0`和`8`。136+ rank，`<device>`为本机Device ID。
110 137 
111- 编译选项说明138- 编译选项说明
112 139 
@@ -120,19 +147,21 @@ peer对应的两条Channel读取对端HCCL Buffer，将读取结果写入本端
120 日志顺序可能因多进程并发而不同。每个rank都应完成到其他rank的数据校验，例如：147 日志顺序可能因多进程并发而不同。每个rank都应完成到其他rank的数据校验，例如：
121 148 
122 ```text149 ```text
123- rank 0/2: device=0, channels_per_peer=2, channel_descs=2150+ [rank 0] one_multi_path | channels_per_peer=2, channel_descs=2
124 rank 0: peer=1, rank_graph_topo=5, protocol=UB_MEM, path_modes=1/2151 rank 0: peer=1, rank_graph_topo=5, protocol=UB_MEM, path_modes=1/2
125 rank 0: dual-path data copy from remote rank 1 passed, bytes=4096152 rank 0: dual-path data copy from remote rank 1 passed, bytes=4096
126- rank 0: dual-path validation passed153+ [rank 0] one_multi_path | dual-path validation | PASS
127 ```154 ```
128 155 
129- 所有rank均为每个peer取得两条不同的Channel Handle、完成数据校验且脚本返回0，表示样例执行成功。156+ 所有rank均输出`PASS`且最终输出`RESULT | Example=one_multi_path | Status=PASS`，表示样例执行成功。
130 157 
131## 注意事项158## 注意事项
132 159 
133-- 默认场景需要Device 0和Device 1可用，且两个Device之间存在`COMM_PROTOCOL_UB_MEM`链路。160+- 运行样例需要至少2张NPU；单卡环境仅支持编译验证。
134-- 多机运行时，所有机器必须使用相同的`-pes`，rank范围必须完整且不能重复。161+- 默认场景需要Device 0和Device 1可用，且两个Device之间存在`COMM_PROTOCOL_UB_MEM`链路；出现
135-- `-ipport`只用于分发`HcclRootInfo`，不参与HCCL数据链路选择。162+ `no UB_MEM link found`时，表示当前rank与目标rank之间不存在可用的`COMM_PROTOCOL_UB_MEM`链路。
163+- `rank_graph_topo`日志只显示所选UB_MEM链路所在的RankGraph层，不参与`pathMode`选择。
164+- 多机运行时，所有机器必须使用相同的`<rank_size>`与`<ip:port>`（指向rank 0），rank范围必须
165+ 完整且不能重复。
166+- `<ip:port>`控制通道只用于分发`HcclRootInfo`与Host侧同步，不参与HCCL数据链路选择。
136- `HcclChannelAcquire`失败时，可根据前序`peer`链路筛选日志定位失败范围。167- `HcclChannelAcquire`失败时，可根据前序`peer`链路筛选日志定位失败范围。
137-- 出现`no UB_MEM link found`时，表示当前rank与目标rank之间不存在可用的
138- `COMM_PROTOCOL_UB_MEM`链路。
Rexamples/one_multi_path/README_en.md→examples/aicore/hcomm/03_one_multi_path/README_en.md+36-28
@@ -4,7 +4,7 @@
4 4 
5This sample demonstrates how to query a `COMM_PROTOCOL_UB_MEM` link through HCCL Resource APIs and create two channels with `pathMode=1` and `pathMode=2` for every peer in one communication domain. Peers are processed serially; for each peer, separate data shards are copied concurrently on two streams.5This sample demonstrates how to query a `COMM_PROTOCOL_UB_MEM` link through HCCL Resource APIs and create two channels with `pathMode=1` and `pathMode=2` for every peer in one communication domain. Peers are processed serially; for each peer, separate data shards are copied concurrently on two streams.
6 6 
7-The sample uses one process per device. Each rank fills its local HCCL Buffer with its rank ID, reads 4 KB from every peer, and verifies the data. `run.sh` starts the local rank processes, and rank 0 distributes7+The sample uses one process per device. Each rank fills its local HCCL Buffer with its rank ID, reads 4 KB from every peer, and verifies the data. The executable forks all local rank processes internally, and rank 0 distributes
8`HcclRootInfo`.8`HcclRootInfo`.
9 9 
10## Supported Products and CANN Software Versions10## Supported Products and CANN Software Versions
@@ -21,7 +21,6 @@ one_multi_path
21├── main.cpp // Host-side communication domain, channel, and verification flow21├── main.cpp // Host-side communication domain, channel, and verification flow
22├── README.md // Chinese documentation22├── README.md // Chinese documentation
23├── README_en.md // English documentation23├── README_en.md // English documentation
24-├── run.sh // Multi-rank launch script
25└── ubmem_data_copy.asc // Ascend C data copy kernel24└── ubmem_data_copy.asc // Ascend C data copy kernel
26```25```
27 26 
@@ -45,14 +44,14 @@ For the same peer, the two channels use the same communication endpoints and are
45 44 
46### Implementation Flow45### Implementation Flow
47 46 
48-1. Rank 0 creates and distributes one `HcclRootInfo`, and every rank initializes the same communication domain.47+1. The executable forks all rank processes (via `examples/aicore/utils/process_manager.h`); ranks exchange `HcclRootInfo` through a TCP control channel (`examples/aicore/utils/rank_sync.h`) and initialize the same communication domain.
492. Each rank obtains its local HCCL Buffer, fills it with the rank ID, and registers the memory resource required by the channels through `HcclCommMemReg`.482. Each rank obtains its local HCCL Buffer, fills it with the rank ID, and registers the memory resource required by the channels through `HcclCommMemReg`.
503. Each rank scans the RankGraph layers and selects one `COMM_PROTOCOL_UB_MEM` link for every peer.493. Each rank scans the RankGraph layers and selects one `COMM_PROTOCOL_UB_MEM` link for every peer.
514. For every peer, the sample builds two descriptors with the same endpoints and `pathMode=1` and `pathMode=2`.504. For every peer, the sample builds two descriptors with the same endpoints and `pathMode=1` and `pathMode=2`.
525. All peer channel descriptors are passed to one `HcclChannelAcquire` call to create the channels in one batch.515. All peer channel descriptors are passed to one `HcclChannelAcquire` call to create the channels in one batch.
536. Peers are processed serially. For each peer, the two channels read the first and second 2 KB shards on separate streams, and both kernels are submitted before either stream is synchronized. Concurrency is limited to the two paths of that peer.526. Peers are processed serially. For each peer, the two channels read the first and second 2 KB shards on separate streams, and both kernels are submitted before either stream is synchronized. Concurrency is limited to the two paths of that peer.
547. The local verification buffer is copied to the host, where every byte is checked against the peer rank ID.537. The local verification buffer is copied to the host, where every byte is checked against the peer rank ID.
55-8. After all peer data is verified, `HcclBarrier` ensures that every rank has completed its remote reads before communication resources are destroyed.54+8. After all peer data is verified, a barrier (via `RankSyncContext`) ensures that every rank has completed its remote reads before communication resources are destroyed.
56 55 
57The `rank_graph_topo` log reports only the RankGraph layer containing the selected UB_MEM link. It does not56The `rank_graph_topo` log reports only the RankGraph layer containing the selected UB_MEM link. It does not
58select `pathMode`.57select `pathMode`.
@@ -64,7 +63,7 @@ Perform the following steps in the sample root directory. This sample supports N
64- Set Environment Variables63- Set Environment Variables
65 64 
66 Configure the environment variables according to the CANN development toolkit65 Configure the environment variables according to the CANN development toolkit
67- [installation method](../../docs/quick_start_en.md).66+ [installation method](../../../../docs/quick_start_en.md).
68 67 
69 ```bash68 ```bash
70 source ${install_path}/cann/set_env.sh69 source ${install_path}/cann/set_env.sh
@@ -73,39 +72,48 @@ Perform the following steps in the sample root directory. This sample supports N
73 > **Note:** `${install_path}` is the CANN installation directory. The default installation directory is72 > **Note:** `${install_path}` is the CANN installation directory. The default installation directory is
74 > `/usr/local/Ascend`.73 > `/usr/local/Ascend`.
75 74 
76-- Run on Two Devices75+- Single-machine run
77 76 
78- By default, the script starts two processes on devices 0 and 1.77+ The executable forks all rank processes internally; each rank is assigned to the corresponding NPU device:
79 78 
80 ```bash79 ```bash
81- cmake -S . -B build -DCMAKE_ASC_ARCHITECTURES=dav-351080+ mkdir -p build
82- cmake --build build -j81+ cd build
83- bash run.sh82+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
83+ make -j
84+ cd .. && ./build/one_multi_path tcp://127.0.0.1:29625 <nranks>
84 ```85 ```
85 86 
86-- Launch Options87+ | Parameter | Default | Description |
87- 
88- | Option | Default | Description |
89 | --- | --- | --- |88 | --- | --- | --- |
90- | `-pes` | `2` | Total number of ranks; it must be an integer greater than or equal to 2 |89+ | `<ip:port>` | Required | Control channel address that rank 0 listens on (`tcp://ip:port`); use different ports for parallel instances |
91- | `-ipport` | `tcp://127.0.0.1:8899` | Address used by rank 0 to distribute `HcclRootInfo`; use `tcp://[address]:port` for IPv6 |90+ | `<nranks>` | Required | Total number of ranks; must be >= 2 and not exceed available NPU count |
92- | `-gnpus` | `2` | Number of rank processes started on the current host |
93- | `-fpe` | `0` | First rank ID on the current host |
94- | `-fnpu` | `0` | First device ID on the current host |
95 91 
96-- Run on Multiple Hosts92+- Multi-machine run
97 93 
98- For eight devices across two hosts, start the first command on the rank 0 host, followed by the second command on the other host:94+ Each host runs the executable with per-rank arguments:
95+ 
96+ ```bash
97+ ./build/one_multi_path <tcp://ip:port> <rank_size> <rank> <device>
98+ ```
99+ 
100+ For eight devices across two hosts:
99 101 
100 ```bash102 ```bash
101 # Host A: ranks 0-3, devices 0-3103 # Host A: ranks 0-3, devices 0-3
102- bash run.sh -pes 8 -ipport tcp://<host_A_IP>:8899 -gnpus 4 -fpe 0 -fnpu 0104+ ./build/one_multi_path tcp://<host_A_IP>:8899 8 0 0
105+ ./build/one_multi_path tcp://<host_A_IP>:8899 8 1 1
106+ ./build/one_multi_path tcp://<host_A_IP>:8899 8 2 2
107+ ./build/one_multi_path tcp://<host_A_IP>:8899 8 3 3
103 108 
104 # Host B: ranks 4-7, devices 0-3109 # Host B: ranks 4-7, devices 0-3
105- bash run.sh -pes 8 -ipport tcp://<host_A_IP>:8899 -gnpus 4 -fpe 4 -fnpu 0110+ ./build/one_multi_path tcp://<host_A_IP>:8899 8 4 0
111+ ./build/one_multi_path tcp://<host_A_IP>:8899 8 5 1
112+ ./build/one_multi_path tcp://<host_A_IP>:8899 8 6 2
113+ ./build/one_multi_path tcp://<host_A_IP>:8899 8 7 3
106 ```114 ```
107 115 
108- For 16 devices with eight devices per host, set `-pes` to `16`, `-gnpus` to `8`, and use `-fpe 0` and `-fpe 8` on the two hosts.116+ `<ip:port>` must point to the rank 0 host on every machine; `<rank_size>` is the global rank count; `<rank>` is this process's global rank ID; `<device>` is the local device ID.
109 117 
110- Build Options118- Build Options
111 119 
@@ -119,20 +127,20 @@ Perform the following steps in the sample root directory. This sample supports N
119 Log order can vary because the rank processes run concurrently. Every rank should validate every peer:127 Log order can vary because the rank processes run concurrently. Every rank should validate every peer:
120 128 
121 ```text129 ```text
122- rank 0/2: device=0, channels_per_peer=2, channel_descs=2130+ [rank 0] one_multi_path | channels_per_peer=2, channel_descs=2
123 rank 0: peer=1, rank_graph_topo=5, protocol=UB_MEM, path_modes=1/2131 rank 0: peer=1, rank_graph_topo=5, protocol=UB_MEM, path_modes=1/2
124 rank 0: dual-path data copy from remote rank 1 passed, bytes=4096132 rank 0: dual-path data copy from remote rank 1 passed, bytes=4096
125- rank 0: dual-path validation passed133+ [rank 0] one_multi_path | dual-path validation | PASS
126 ```134 ```
127 135 
128 The sample succeeds when every rank obtains two different channel handles for every peer, verifies the data,136 The sample succeeds when every rank obtains two different channel handles for every peer, verifies the data,
129- and the script exits with 0.137+ and the final output is `RESULT | Example=one_multi_path | Status=PASS`.
130 138 
131## Notes139## Notes
132 140 
133- The default run requires devices 0 and 1 and a `COMM_PROTOCOL_UB_MEM` link between them.141- The default run requires devices 0 and 1 and a `COMM_PROTOCOL_UB_MEM` link between them.
134-- For a multi-host run, every host must use the same `-pes`, and the rank ranges must be complete and unique.142+- For a multi-host run, every host must use the same `<rank_size>` and `<ip:port>` (pointing to rank 0), and the rank ranges must be complete and unique.
135-- `-ipport` distributes `HcclRootInfo` only. It does not select the HCCL data link.143+- `<ip:port>` is used only for distributing `HcclRootInfo` and host-side synchronization. It does not select the HCCL data link.
136- If `HcclChannelAcquire` fails, use the preceding peer link-selection logs to narrow down the failing range.144- If `HcclChannelAcquire` fails, use the preceding peer link-selection logs to narrow down the failing range.
137- `no UB_MEM link found` means that no `COMM_PROTOCOL_UB_MEM` link is available between the local and target145- `no UB_MEM link found` means that no `COMM_PROTOCOL_UB_MEM` link is available between the local and target
138 ranks.146 ranks.
Rexamples/one_multi_path/main.cpp→examples/aicore/hcomm/03_one_multi_path/main.cpp+85-317
@@ -8,7 +8,6 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11-#include <chrono>
12#include <cerrno>11#include <cerrno>
13#include <cstdio>12#include <cstdio>
14#include <cstdint>13#include <cstdint>
@@ -17,37 +16,19 @@
17#include <iostream>16#include <iostream>
18#include <limits>17#include <limits>
19#include <string>18#include <string>
20-#include <thread>
21#include <vector>19#include <vector>
22 20 
23-#include <netdb.h>
24-#include <sys/socket.h>
25-#include <sys/time.h>
26-#include <unistd.h>
27- 
28#include "acl/acl.h"21#include "acl/acl.h"
29#include "hccl/hccl.h"22#include "hccl/hccl.h"
30#include "hccl/hccl_rank_graph.h"23#include "hccl/hccl_rank_graph.h"
31#include "hccl/hccl_res.h"24#include "hccl/hccl_res.h"
32#include "hccl/hccl_types.h"25#include "hccl/hccl_types.h"
33 26 
34-#define HCCLCHECK(ret) \27+#include "utils/check.h"
35- do { \28+#include "utils/hccl_comm_init.h"
36- const HcclResult checkRet = (ret); \29+#include "utils/process_manager.h"
37- if (checkRet != HCCL_SUCCESS) { \30+#include "utils/rank_sync.h"
38- std::fprintf(stderr, "%s:%d: %s return %d.\n", __FILE__, __LINE__, #ret, checkRet); \31+#include "utils/arg_parser.h"
39- return checkRet; \
40- } \
41- } while (0)
42- 
43-#define ACLCHECK(ret) \
44- do { \
45- const aclError checkRet = (ret); \
46- if (checkRet != ACL_SUCCESS) { \
47- std::fprintf(stderr, "%s:%d: %s return %d.\n", __FILE__, __LINE__, #ret, checkRet); \
48- return HCCL_E_RUNTIME; \
49- } \
50- } while (0)
51 32 
52extern "C" aclError LaunchUbmemDataCopy(void* dst, void* src, uint32_t byteSize, aclrtStream stream);33extern "C" aclError LaunchUbmemDataCopy(void* dst, void* src, uint32_t byteSize, aclrtStream stream);
53 34 
@@ -60,9 +41,10 @@ constexpr uint8_t PATH_MODE_ONE = 1;
60constexpr uint8_t PATH_MODE_MULTI = 2;41constexpr uint8_t PATH_MODE_MULTI = 2;
61constexpr CommProtocol TARGET_PROTOCOL = COMM_PROTOCOL_UB_MEM;42constexpr CommProtocol TARGET_PROTOCOL = COMM_PROTOCOL_UB_MEM;
62constexpr uint32_t PATH_COUNT = 2;43constexpr uint32_t PATH_COUNT = 2;
63-constexpr time_t SOCKET_IO_TIMEOUT_SECONDS = 30;
64constexpr uint8_t TARGET_PATH_MODES[PATH_COUNT] = {PATH_MODE_ONE, PATH_MODE_MULTI};44constexpr uint8_t TARGET_PATH_MODES[PATH_COUNT] = {PATH_MODE_ONE, PATH_MODE_MULTI};
65static_assert(VERIFY_BYTES % PATH_COUNT == 0, "VERIFY_BYTES must be divisible by PATH_COUNT");45static_assert(VERIFY_BYTES % PATH_COUNT == 0, "VERIFY_BYTES must be divisible by PATH_COUNT");
46+// 收尾barrier的tag：本样例只有这一个host侧同步点。
47+constexpr uint32_t kTagFinalRelease = examples::kTagUserBase;
66 48 
67struct Options {49struct Options {
68 uint32_t rankSize = 0;50 uint32_t rankSize = 0;
@@ -78,7 +60,6 @@ struct CommResources {
78 void* verifyBuffer = nullptr;60 void* verifyBuffer = nullptr;
79};61};
80 62 
81-int ExchangeRootInfo(uint32_t rank, uint32_t rankSize, const std::string& ipPort, HcclRootInfo& rootInfo);
82int ReportError(uint32_t rank, const char* operation, int result);63int ReportError(uint32_t rank, const char* operation, int result);
83int BuildChannelDescs(64int BuildChannelDescs(
84 HcclComm comm, uint32_t rank, uint32_t peerRank, HcclMemHandle& memHandle, HcclChannelDesc* descs);65 HcclComm comm, uint32_t rank, uint32_t peerRank, HcclMemHandle& memHandle, HcclChannelDesc* descs);
@@ -86,14 +67,14 @@ int TransferAndVerifyShards(
86 uint32_t rank, uint32_t peerRank, HcclComm comm, const ChannelHandle* channels, CommResources& resources);67 uint32_t rank, uint32_t peerRank, HcclComm comm, const ChannelHandle* channels, CommResources& resources);
87int CleanupCommRes(CommResources& resources, int status);68int CleanupCommRes(CommResources& resources, int status);
88 69 
89-int TestDualPath(const Options& options, const HcclRootInfo& rootInfo)70+int TestDualPath(const Options& options, examples::RankSyncContext& sync)
90{71{
91 CommResources resources;72 CommResources resources;
92 HcclMemHandle memHandle = nullptr;73 HcclMemHandle memHandle = nullptr;
93 std::vector<uint32_t> peerRanks;74 std::vector<uint32_t> peerRanks;
94 75 
95 // Initialize the communication domain and prepare data in the local HCCL Buffer.76 // Initialize the communication domain and prepare data in the local HCCL Buffer.
96- int result = HcclCommInitRootInfo(options.rankSize, &rootInfo, options.rank, &resources.comm);77+ int result = examples::InitCommByRootInfo(sync, &resources.comm);
97 if (result != HCCL_SUCCESS) {78 if (result != HCCL_SUCCESS) {
98 return CleanupCommRes(resources, ReportError(options.rank, "HcclCommInitRootInfo", result));79 return CleanupCommRes(resources, ReportError(options.rank, "HcclCommInitRootInfo", result));
99 }80 }
@@ -188,14 +169,11 @@ int TestDualPath(const Options& options, const HcclRootInfo& rootInfo)
188 }169 }
189 }170 }
190 171 
191- // Keep every rank's HCCL Buffer valid until all remote reads have completed.172+ // Keep every rank's HCCL Buffer valid until all remote reads have completed. Each rank has
192- result = HcclBarrier(resources.comm, resources.streams[0]);173+ // already synchronized its own streams above, so the host-side barrier suffices: once every
193- if (result != HCCL_SUCCESS) {174+ // rank passes it, all remote reads in the group are done and freeing is safe.
194- return CleanupCommRes(resources, ReportError(options.rank, "HcclBarrier", result));175+ if (!sync.Barrier(kTagFinalRelease)) {
195- }176+ return CleanupCommRes(resources, ReportError(options.rank, "Barrier", HCCL_E_INTERNAL));
196- result = aclrtSynchronizeStream(resources.streams[0]);
197- if (result != ACL_SUCCESS) {
198- return CleanupCommRes(resources, ReportError(options.rank, "aclrtSynchronizeStream after HcclBarrier", result));
199 }177 }
200 return CleanupCommRes(resources, HCCL_SUCCESS);178 return CleanupCommRes(resources, HCCL_SUCCESS);
201}179}
@@ -368,301 +346,37 @@ int CleanupCommRes(CommResources& resources, int status)
368 return status != HCCL_SUCCESS ? status : cleanupStatus;346 return status != HCCL_SUCCESS ? status : cleanupStatus;
369}347}
370 348 
371-bool ParseNonNegativeInt(const char* text, int& value)
372-{
373- if (text == nullptr || *text == '\0') {
374- return false;
375- }
376- 
377- char* parseEnd = nullptr;
378- errno = 0;
379- const long parsed = std::strtol(text, &parseEnd, 10);
380- if (errno == ERANGE || parseEnd == text || *parseEnd != '\0' || parsed < 0 ||
381- parsed > std::numeric_limits<int>::max()) {
382- return false;
383- }
384- value = static_cast<int>(parsed);
385- return true;
386-}
387- 
388-bool ParseIpPort(const std::string& ipPort, std::string& host, uint16_t& port)
389-{
390- constexpr const char* TCP_PREFIX = "tcp://";
391- std::string value = ipPort;
392- if (value.compare(0, std::strlen(TCP_PREFIX), TCP_PREFIX) == 0) {
393- value = value.substr(std::strlen(TCP_PREFIX));
394- }
395- 
396- size_t separator = value.rfind(':');
397- if (separator == std::string::npos || separator == 0 || separator + 1 >= value.size()) {
398- return false;
399- }
400- 
401- const std::string portText = value.substr(separator + 1);
402- char* parseEnd = nullptr;
403- errno = 0;
404- const long parsedPort = std::strtol(portText.c_str(), &parseEnd, 10);
405- if (errno == ERANGE || parseEnd == portText.c_str() || *parseEnd != '\0' || parsedPort <= 0 ||
406- parsedPort > UINT16_MAX) {
407- return false;
408- }
409- host = value.substr(0, separator);
410- if (host.front() == '[') {
411- if (host.size() < 3 || host.back() != ']') {
412- return false;
413- }
414- host = host.substr(1, host.size() - 2);
415- }
416- port = static_cast<uint16_t>(parsedPort);
417- return true;
418-}
419- 
420-bool SendAll(int fd, const void* data, size_t size)
421-{
422- const char* buffer = static_cast<const char*>(data);
423- size_t sent = 0;
424- while (sent < size) {
425- const ssize_t count = send(fd, buffer + sent, size - sent, MSG_NOSIGNAL);
426- if (count < 0 && errno == EINTR) {
427- continue;
428- }
429- if (count < 0) {
430- std::cerr << "send failed, errno=" << errno << std::endl;
431- return false;
432- }
433- if (count == 0) {
434- std::cerr << "send returned 0 before all data was sent" << std::endl;
435- return false;
436- }
437- sent += static_cast<size_t>(count);
438- }
439- return true;
440-}
441- 
442-bool RecvAll(int fd, void* data, size_t size)
443-{
444- char* buffer = static_cast<char*>(data);
445- size_t received = 0;
446- while (received < size) {
447- const ssize_t count = recv(fd, buffer + received, size - received, 0);
448- if (count < 0 && errno == EINTR) {
449- continue;
450- }
451- if (count < 0) {
452- std::cerr << "recv failed, errno=" << errno << std::endl;
453- return false;
454- }
455- if (count == 0) {
456- std::cerr << "peer closed the connection before all data was received" << std::endl;
457- return false;
458- }
459- received += static_cast<size_t>(count);
460- }
461- return true;
462-}
463- 
464-bool SetSocketTimeouts(int fd)
465-{
466- timeval timeout{};
467- timeout.tv_sec = SOCKET_IO_TIMEOUT_SECONDS;
468- if (setsockopt(fd, SOL_SOCKET, SO_RCVTIMEO, &timeout, sizeof(timeout)) != 0) {
469- std::cerr << "setsockopt SO_RCVTIMEO failed, errno=" << errno << std::endl;
470- return false;
471- }
472- if (setsockopt(fd, SOL_SOCKET, SO_SNDTIMEO, &timeout, sizeof(timeout)) != 0) {
473- std::cerr << "setsockopt SO_SNDTIMEO failed, errno=" << errno << std::endl;
474- return false;
475- }
476- return true;
477-}
478- 
479-int CreateListenSocket(const std::string& host, uint16_t port)
480-{
481- addrinfo hints{};
482- hints.ai_family = AF_UNSPEC;
483- hints.ai_socktype = SOCK_STREAM;
484- hints.ai_flags = AI_PASSIVE;
485- const std::string portText = std::to_string(port);
486- addrinfo* addresses = nullptr;
487- const int resolveResult = getaddrinfo(host.c_str(), portText.c_str(), &hints, &addresses);
488- if (resolveResult != 0) {
489- std::cerr << "getaddrinfo for listen address failed: " << gai_strerror(resolveResult) << std::endl;
490- return -1;
491- }
492- 
493- int listenFd = -1;
494- int lastError = 0;
495- for (const addrinfo* address = addresses; address != nullptr; address = address->ai_next) {
496- int fd = socket(address->ai_family, address->ai_socktype, address->ai_protocol);
497- if (fd < 0) {
498- lastError = errno;
499- continue;
500- }
501- if (!SetSocketTimeouts(fd)) {
502- lastError = errno;
503- close(fd);
504- continue;
505- }
506- int reuseAddress = 1;
507- if (setsockopt(fd, SOL_SOCKET, SO_REUSEADDR, &reuseAddress, sizeof(reuseAddress)) != 0) {
508- lastError = errno;
509- close(fd);
510- continue;
511- }
512- if (bind(fd, address->ai_addr, address->ai_addrlen) != 0 || listen(fd, 64) != 0) {
513- lastError = errno;
514- close(fd);
515- continue;
516- }
517- listenFd = fd;
518- break;
519- }
520- freeaddrinfo(addresses);
521- if (listenFd < 0) {
522- std::cerr << "listen on " << host << ":" << port << " failed, errno=" << lastError << std::endl;
523- }
524- return listenFd;
525-}
526- 
527-int ConnectToRoot(const std::string& host, uint16_t port)
528-{
529- constexpr int RETRY_TIMES = 300;
530- constexpr int RETRY_INTERVAL_MS = 100;
531- 
532- addrinfo hints{};
533- hints.ai_family = AF_UNSPEC;
534- hints.ai_socktype = SOCK_STREAM;
535- const std::string portText = std::to_string(port);
536- addrinfo* addresses = nullptr;
537- const int resolveResult = getaddrinfo(host.c_str(), portText.c_str(), &hints, &addresses);
538- if (resolveResult != 0) {
539- std::cerr << "getaddrinfo for root address failed: " << gai_strerror(resolveResult) << std::endl;
540- return -1;
541- }
542- 
543- for (int retry = 0; retry < RETRY_TIMES; ++retry) {
544- for (const addrinfo* address = addresses; address != nullptr; address = address->ai_next) {
545- int fd = socket(address->ai_family, address->ai_socktype, address->ai_protocol);
546- if (fd < 0) {
547- continue;
548- }
549- if (!SetSocketTimeouts(fd)) {
550- close(fd);
551- continue;
552- }
553- if (connect(fd, address->ai_addr, address->ai_addrlen) == 0) {
554- freeaddrinfo(addresses);
555- return fd;
556- }
557- close(fd);
558- }
559- if (retry + 1 < RETRY_TIMES) {
560- std::this_thread::sleep_for(std::chrono::milliseconds(RETRY_INTERVAL_MS));
561- }
562- }
563- freeaddrinfo(addresses);
564- return -1;
565-}
566- 
567-int ExchangeRootInfo(uint32_t rank, uint32_t rankSize, const std::string& ipPort, HcclRootInfo& rootInfo)
568-{
569- std::string host;
570- uint16_t port = 0;
571- if (!ParseIpPort(ipPort, host, port)) {
572- std::cerr << "invalid ipport: " << ipPort << std::endl;
573- return -1;
574- }
575- 
576- if (rank == 0) {
577- int listenFd = CreateListenSocket(host, port);
578- if (listenFd < 0) {
579- return -1;
580- }
581- for (uint32_t peer = 1; peer < rankSize; ++peer) {
582- int connectionFd = -1;
583- do {
584- connectionFd = accept(listenFd, nullptr, nullptr);
585- } while (connectionFd < 0 && errno == EINTR);
586- if (connectionFd < 0) {
587- std::cerr << "accept failed, errno=" << errno << std::endl;
588- close(listenFd);
589- return -1;
590- }
591- if (!SetSocketTimeouts(connectionFd)) {
592- close(connectionFd);
593- close(listenFd);
594- return -1;
595- }
596- if (!SendAll(connectionFd, &rootInfo, sizeof(rootInfo))) {
597- std::cerr << "send root info failed" << std::endl;
598- close(connectionFd);
599- close(listenFd);
600- return -1;
601- }
602- close(connectionFd);
603- }
604- close(listenFd);
605- return 0;
606- }
607- 
608- int fd = ConnectToRoot(host, port);
609- if (fd < 0) {
610- std::cerr << "connect ipport failed: " << ipPort << std::endl;
611- return -1;
612- }
613- if (!RecvAll(fd, &rootInfo, sizeof(rootInfo))) {
614- std::cerr << "recv root info failed" << std::endl;
615- close(fd);
616- return -1;
617- }
618- close(fd);
619- return 0;
620-}
621- 
622} // namespace349} // namespace
623 350 
624-int main(int argc, char* argv[])351+// 单个rank的完整流程，单机与多机模式共用。由ProcessGroup（单机）或main的5参数分支
352+// （多机）调用。
353+int RunOneMultiPathRank(uint32_t rankSize, uint32_t rank, const std::string& ipPort, uint32_t device)
625{354{
626- if (argc != 5) {355+ Options options{rankSize, rank, device, ipPort};
627- std::cerr << "invalid arguments" << std::endl;356+ 
357+ examples::RankSyncContext sync = rank == 0U ? examples::RankSyncContext::Listen(ipPort, rankSize) :
358+ examples::RankSyncContext::Connect(rank, rankSize, ipPort);
359+ if (!sync.Ok()) {
360+ std::cerr << "[rank " << rank << "] control channel setup failed" << std::endl;
628 return 1;361 return 1;
629 }362 }
630 363 
631- int rankSize = 0;364+ EXAMPLES_ACLCHECK(aclInit(nullptr));
632- int rank = 0;
633- int device = 0;
634- if (!ParseNonNegativeInt(argv[1], rankSize) || !ParseNonNegativeInt(argv[2], rank) ||
635- !ParseNonNegativeInt(argv[4], device) || rankSize < 2 || rank >= rankSize) {
636- std::cerr << "invalid arguments: rank_size must be at least 2, rank must be in [0, rank_size), and device "
637- "must be a non-negative integer"
638- << std::endl;
639- return 1;
640- }
641- Options options{
642- static_cast<uint32_t>(rankSize), static_cast<uint32_t>(rank), static_cast<uint32_t>(device), argv[3]};
643- 
644- ACLCHECK(aclInit(nullptr));
645 uint32_t deviceCount = 0;365 uint32_t deviceCount = 0;
646- ACLCHECK(aclrtGetDeviceCount(&deviceCount));366+ EXAMPLES_ACLCHECK(aclrtGetDeviceCount(&deviceCount));
647 if (options.device >= deviceCount) {367 if (options.device >= deviceCount) {
648 std::cerr << "device " << options.device << " is out of range, device_count=" << deviceCount << std::endl;368 std::cerr << "device " << options.device << " is out of range, device_count=" << deviceCount << std::endl;
649 return 1;369 return 1;
650 }370 }
651- ACLCHECK(aclrtSetDevice(static_cast<int32_t>(options.device)));371+ EXAMPLES_ACLCHECK(aclrtSetDevice(static_cast<int32_t>(options.device)));
652- 
653- HcclRootInfo rootInfo{};
654- if (options.rank == 0) {
655- HCCLCHECK(HcclGetRootInfo(&rootInfo));
656- }
657 372 
658 int status = HCCL_E_INTERNAL;373 int status = HCCL_E_INTERNAL;
659- if (ExchangeRootInfo(options.rank, options.rankSize, options.ipPort, rootInfo) == 0) {374+ {
660- std::cout << "rank " << options.rank << "/" << options.rankSize << ": device=" << options.device375+ std::cout << "[rank " << options.rank << "] one_multi_path | channels_per_peer=" << PATH_COUNT
661- << ", channels_per_peer=" << PATH_COUNT << ", channel_descs=" << (options.rankSize - 1) * PATH_COUNT376+ << ", channel_descs=" << (options.rankSize - 1) * PATH_COUNT << std::endl;
662- << std::endl;377+ status = TestDualPath(options, sync);
663- status = TestDualPath(options, rootInfo);
664 if (status == HCCL_SUCCESS) {378 if (status == HCCL_SUCCESS) {
665- std::cout << "rank " << options.rank << ": dual-path validation passed" << std::endl;379+ std::cout << "[rank " << options.rank << "] one_multi_path | dual-path validation | PASS" << std::endl;
666 }380 }
667 }381 }
668 382 
@@ -674,3 +388,57 @@ int main(int argc, char* argv[])
674 }388 }
675 return status == HCCL_SUCCESS ? 0 : 1;389 return status == HCCL_SUCCESS ? 0 : 1;
676}390}
391+ 
392+// 单机模式：one_multi_path <tcp://ip:port> <nranks>，父进程fork全部rank进程，
393+// device固定为rank对应的NPU。多机模式：one_multi_path <tcp://ip:port> <rank_size> <rank> <device>，
394+// 每台机器各自启动本机的rank进程，endpoint统一指向rank 0所在机器。
395+int main(int argc, char* argv[])
396+{
397+ // 单机：endpoint + nranks。
398+ if (argc == 3) {
399+ uint32_t uRankSize = 0;
400+ if (!examples::ParseUint32(argv[2], uRankSize, 2)) {
401+ std::cerr << "invalid nranks: " << argv[2] << ", expected an integer >= 2" << std::endl;
402+ return 1;
403+ }
404+ const std::string endpoint = argv[1];
405+ 
406+ std::vector<examples::ArgsOf<decltype(RunOneMultiPathRank)>> perRankArgs;
407+ perRankArgs.reserve(uRankSize);
408+ for (uint32_t rank = 0; rank < uRankSize; ++rank) {
409+ perRankArgs.emplace_back(uRankSize, rank, endpoint, rank);
410+ }
411+ 
412+ examples::ProcessGroup group;
413+ if (!group.Launch(uRankSize, RunOneMultiPathRank, std::move(perRankArgs))) {
414+ return 1;
415+ }
416+ const bool pass = group.WaitAll();
417+ if (group.Interrupted()) {
418+ std::cerr << "RESULT | Example=one_multi_path | Status=INTERRUPTED" << std::endl;
419+ return 1;
420+ }
421+ std::cout << "RESULT | Example=one_multi_path | Status=" << (pass ? "PASS" : "FAIL") << std::endl;
422+ return pass ? 0 : 1;
423+ }
424+ 
425+ // 多机：每rank一进程，参数为 <endpoint> <rank_size> <rank> <device>。
426+ if (argc != 5) {
427+ std::cerr << "Usage: " << argv[0] << " <tcp://ip:port> <nranks> (single machine)" << std::endl;
428+ std::cerr << " " << argv[0] << " <tcp://ip:port> <rank_size> <rank> <device> (per-rank, multi machine)"
429+ << std::endl;
430+ return 1;
431+ }
432+ 
433+ uint32_t uRankSize = 0;
434+ uint32_t uRank = 0;
435+ uint32_t uDevice = 0;
436+ if (!examples::ParseUint32(argv[2], uRankSize, 2) || !examples::ParseUint32(argv[3], uRank) ||
437+ !examples::ParseUint32(argv[4], uDevice) || uRank >= uRankSize) {
438+ std::cerr << "invalid arguments: rank_size must be at least 2, rank must be in [0, rank_size), and device "
439+ "must be a non-negative integer"
440+ << std::endl;
441+ return 1;
442+ }
443+ return RunOneMultiPathRank(uRankSize, uRank, argv[1], uDevice);
444+}
Rexamples/one_multi_path/ubmem_data_copy.asc→examples/aicore/hcomm/03_one_multi_path/ubmem_data_copy.asc+0-0
文件重命名但无更改。
Rexamples/simt_urma/CMakeLists.txt→examples/aicore/hcomm/04_simt_urma/CMakeLists.txt+3-15
@@ -12,23 +12,12 @@ cmake_minimum_required(VERSION 3.16)
12 12 
13set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")13set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
14 14 
15-if(NOT CMAKE_ASC_ARCHITECTURES STREQUAL "dav-3510")
16- message(FATAL_ERROR "simt_urma only supports dav-3510, but got ${CMAKE_ASC_ARCHITECTURES}")
17-endif()
18- 
19find_package(ASC REQUIRED)15find_package(ASC REQUIRED)
20- 
21project(simt_urma LANGUAGES ASC CXX)16project(simt_urma LANGUAGES ASC CXX)
22 17 
23-set(CMAKE_CXX_STANDARD 17)
24-set(CMAKE_CXX_STANDARD_REQUIRED ON)
25set(CMAKE_POSITION_INDEPENDENT_CODE ON)18set(CMAKE_POSITION_INDEPENDENT_CODE ON)
26 19 
27-# Only the kernel is compiled as ASC. main.cpp remains on the host C++ toolchain.20+set_source_files_properties(op_kernel.cpp PROPERTIES LANGUAGE ASC)
28-set_source_files_properties(
29- op_kernel.cpp
30- PROPERTIES LANGUAGE ASC
31-)
32 21 
33add_executable(simt_urma22add_executable(simt_urma
34 main.cpp23 main.cpp
@@ -43,13 +32,12 @@ set_target_properties(simt_urma PROPERTIES
43 LINKER_LANGUAGE ASC32 LINKER_LANGUAGE ASC
44)33)
45 34 
46-target_include_directories(simt_urma BEFORE PRIVATE35+target_include_directories(simt_urma PRIVATE
47 ${CMAKE_CURRENT_SOURCE_DIR}36 ${CMAKE_CURRENT_SOURCE_DIR}
37+ ${CMAKE_CURRENT_SOURCE_DIR}/../..
48 $ENV{ASCEND_HOME_PATH}/include38 $ENV{ASCEND_HOME_PATH}/include
49- $<$<COMPILE_LANGUAGE:ASC>:$ENV{ASCEND_HOME_PATH}/asc/include/adv_api>
50)39)
51 40 
52-# Host-side communication-domain, memory-registration and channel APIs are provided by libhcomm.so.
53target_link_directories(simt_urma PRIVATE41target_link_directories(simt_urma PRIVATE
54 $ENV{ASCEND_HOME_PATH}/lib6442 $ENV{ASCEND_HOME_PATH}/lib64
55 $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/lib6443 $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/lib64
Aexamples/aicore/hcomm/04_simt_urma/README.md+163-0
@@ -0,0 +1,163 @@
1+# Hcomm SIMT URMA 功能样例
2+ 
3+## 概述
4+ 
5+本样例验证`hcomm/hcomm_simt.h`中全部SIMT URMA点对点通信接口：`WriteNbi`、`WriteValueNbi`、
6+`WriteWithNotifyNbi`、`AtomicFAA`和`AtomicCAS`，并覆盖立即提交和延迟提交两种发布方式。
7+ 
8+拓扑为环形：每个rank把操作发送到`next = (rank + 1) % nranks`的注册内存，并校验
9+`prev = (rank - 1 + nranks) % nranks`写入的结果。每个rank既是发送方也是接收方，两侧都校验：
10+`sendBuf`保存本rank原子操作取回的旧值，`recvBuf`保存prev写入的数据，环上每个`recvBuf`恰好只有
11+一个写者，因此期望值都是确定的。
12+ 
13+所有操作均由单个SIMT线程发起。一个channel必须由一个lane驱动：延迟提交的WQE停留在发送队列中，
14+由后续`commit=true`的提交发布覆盖整批PI，因此同一批任务必须由同一个lane提交。
15+ 
16+## 本样例支持的产品及CANN软件版本
17+ 
18+| 产品 | CANN软件版本 |
19+| --- | --- |
20+| Ascend 950PR/Ascend 950DT | >= CANN 9.1.0 |
21+ 
22+## 目录结构介绍
23+ 
24+```text
25+simt_urma
26+├── CMakeLists.txt // 编译工程文件
27+├── main.cpp // Host侧通信域、内存注册和校验流程
28+├── op_kernel.cpp // SIMT Kernel侧URMA接口调用
29+├── simt_urma_common.h // Host与Kernel共享定义
30+├── README.md // 中文样例说明
31+└── README_en.md // 英文样例说明
32+```
33+ 
34+## 样例描述
35+ 
36+### 功能说明
37+ 
38+样例提供十种测试模式，覆盖五个接口各自的立即提交与延迟提交：
39+ 
40+**Write系列**
41+ 
42+`WriteNbi`从本地缓冲区取数据，WQE中是指向该缓冲区的SGE。第`i`个slot携带值`i + 1`，因此接收侧
43+任何错值都能直接定位到它本该来自哪个slot。
44+ 
45+`WriteValueNbi`把待写值直接内联在WQE中，不读发送缓冲区。要求`config`开启inline（默认即
46+`URMA_INLINE_CFG`），且`sizeof(T)`不超过WQE内联负载区，两者均有编译期检查。
47+ 
48+| `mode` | 提交方式 |
49+| --- | --- |
50+| `write_single` | 32次独立的`WriteNbi<commit=true>` |
51+| `write_batch_last` | 31次串行`WriteNbi<false>`，最后1次`WriteNbi<true>` |
52+| `write_value_single` | 32次独立的`WriteValueNbi<commit=true>` |
53+| `write_value_batch_last` | 31次串行`WriteValueNbi<false>`，最后1次`WriteValueNbi<true>` |
54+ 
55+**Notify/Atomic系列**
56+ 
57+| `mode` | 说明 |
58+| --- | --- |
59+| `notify` | 提交一次Notify操作并调用`Drain` |
60+| `faa` | 提交一次FAA操作，校验远端累加结果和本地取回的旧值 |
61+| `cas` | 提交一次CAS操作，校验远端交换结果和本地取回的旧值 |
62+| `single`（默认） | 串行提交Notify、FAA和CAS，每次操作后分别调用`Drain` |
63+| `batch_last` | Notify和FAA使用`commit=false`延迟提交，最后由`commit=true`的CAS提交整批 |
64+| `notify_immediate_repeat` | 单lane连续提交`kNotifyImmediateRepeatCount`个`WriteWithNotifyNbi<true>`（当前为100），最后统一调用一次`Drain` |
65+ 
66+两个系列使用互不重叠的slot布局，缓冲区按两者中较大的布局分配。SIMT没有`Commit`接口，延迟提交
67+的任务只能由后续一次`commit=true`的提交带出；所有模式最后都调用`Drain`，确保kernel返回前CQ已被
68+消费。
69+ 
70+### 样例规格
71+ 
72+| 项目 | 说明 |
73+| --- | --- |
74+| 通信模式 | 环形拓扑SIMT URMA点对点通信 |
75+| 调用方式 | SIMT Kernel内直接调用URMA接口 |
76+| 测试模式 | `write_single`、`write_batch_last`、`write_value_single`、`write_value_batch_last`、`notify`、`faa`、`cas`、`single`、`batch_last`、`notify_immediate_repeat` |
77+| 支持rank数 | 任意 >= 2 |
78+ 
79+### 实现流程
80+ 
81+1. 可执行文件直接拉起全部rank进程（进程内fork，参见`examples/aicore/utils/process_manager.h`）；
82+ rank间通过命令行指定的TCP控制通道（`examples/aicore/utils/rank_sync.h`）交换root info
83+ 并创建通信域。
84+2. 每个rank通过`HcclCommMemReg`注册`sendBuf`和`recvBuf`。
85+3. 通过`HcclRankGraphGetLayers`/`HcclRankGraphGetLinks`获取到`next`的链路Endpoint，调用
86+ `HcclChannelAcquire`创建Channel，并通过`HcclChannelGetRemoteMems`获取对端内存地址。
87+4. Kernel侧按选定模式提交URMA操作，最后调用`Drain`等待完成。
88+5. Host侧通过控制通道的两次barrier（`ready`、`sent`）保证所有rank完成注册和写入后再回读校验。
89+6. 校验`sendBuf`取回值和`recvBuf`落地数据，输出PASS/FAIL。
90+ 
91+## 编译运行
92+ 
93+在本样例根目录下执行如下步骤，编译并运行样例。本样例仅支持NPU运行模式。
94+ 
95+- 配置环境变量
96+ 
97+ 样例编译时使用已安装到CANN目录的asc-comm头文件。若尚未安装，先在仓库根目录生成并安装开发验证
98+ run包，参数说明参见[`docs/zh/guide/build_and_test.md`](../../../../docs/zh/guide/build_and_test.md)：
99+ 
100+ ```bash
101+ bash build.sh --pkg
102+ ./build_out/cann-asc-comm_1.0.0_linux-<arch>.run --full
103+ ```
104+ 
105+ 然后加载CANN环境变量：
106+ 
107+ ```bash
108+ source ${install_path}/cann/set_env.sh
109+ ```
110+ 
111+ > **说明：** `${install_path}`为CANN包安装目录，未指定安装目录时默认安装至`/usr/local/Ascend`。
112+ 
113+- 样例执行
114+ 
115+ 在本样例目录下执行如下命令。可执行文件内部fork出全部rank进程，直接运行即可：
116+ 
117+ ```bash
118+ mkdir -p build
119+ cd build
120+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
121+ make -j
122+ cd .. && ./build/simt_urma tcp://127.0.0.1:29620 <nranks> [mode]
123+ ```
124+ 
125+- 启动参数说明
126+ 
127+ | 参数 | 默认值 | 说明 |
128+ | --- | --- | --- |
129+ | `<ip:port>` | 无默认 | rank 0监听的控制通道地址（`tcp://ip:port`格式），用于root info交换与Host侧barrier；并行运行多个实例时错开端口 |
130+ | `<nranks>` | 无默认 | 通信域rank数，必须为不小于2的整数 |
131+ | `[mode]` | `single` | 测试模式，取值见样例描述；无法识别的mode会直接报错，不会静默回退 |
132+ 
133+- 编译选项说明
134+ 
135+ | 选项 | 可选值 | 说明 |
136+ | --- | --- | --- |
137+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认） | 运行模式，本样例仅支持NPU运行 |
138+ | `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU架构，对应Ascend 950PR/Ascend 950DT |
139+ 
140+- 执行结果
141+ 
142+ 以2卡运行`write_single`为例，每个rank打印一组结果：
143+ 
144+ ```text
145+ [rank 0] simt_urma write_single | sent to rank 1, received from rank 1 | PASS
146+ [rank 1] simt_urma write_single | sent to rank 0, received from rank 0 | PASS
147+ RESULT | Example=simt_urma Mode=write_single | Status=PASS
148+ ```
149+ 
150+ 所有rank均输出PASS且最终输出`Status=PASS`，表示样例执行成功。
151+ 
152+## 注意事项
153+ 
154+- 运行样例需要至少2张NPU；单卡环境仅支持编译验证。
155+- Kernel侧通过`AscendC::simt::LocalBufferAddr`和`RemoteBufferAddr`解析缓冲区基址时，传入的是
156+ 通道缓冲区表的下标，而不是`memHandles`数组下标：本地表下标0由CCL buffer预留，第一个显式注册的
157+ 缓冲区使用下标1；远端缓冲区下标通过`HcclChannelGetRemoteMems`按tag查找。直接传`memHandles`
158+ 下标会解析到CCL buffer，此时CQE正常返回但接收方看到的是CCL buffer内容而非payload。
159+- Host侧`sent` barrier是校验结果可信的关键：没有它，某rank会在写入方kernel尚未完成时读到清零的
160+ 缓冲区。
161+- 可执行文件捕获`SIGINT`/`SIGTERM`并终止全部rank进程，Ctrl+C不会留下阻塞在Host barrier上、
162+ 仍占用设备的进程。
163+- 编译依赖CANN ASC CMake能力，并在链接阶段依赖CANN `hcomm`库。
Rexamples/simt_urma/README_en.md→examples/aicore/hcomm/04_simt_urma/README_en.md+47-58
@@ -72,71 +72,64 @@ sized to whichever layout is larger.
72SIMT has no `Commit` interface, so a deferred task can only be carried out by a later `commit=true`72SIMT has no `Commit` interface, so a deferred task can only be carried out by a later `commit=true`
73post. Every mode ends with `Drain` to make sure the CQ has been consumed before the kernel returns.73post. Every mode ends with `Drain` to make sure the CQ has been consumed before the kernel returns.
74 74 
75-## Preparing the Development Verification Environment75+## Build and Run
76 76 
77-The sample builds against the asc-comm headers installed into the CANN directory. Generate and77+Perform the following steps in the sample root directory. This sample supports NPU run mode only.
78-install the development verification run package from the repository root first:
79 78 
80-```bash79+- Set Environment Variables
81-source /usr/local/Ascend/cann/set_env.sh
82-bash build.sh --pkg
83-./build_out/cann-asc-comm_1.0.0_linux-<arch>.run --full
84-```
85 80 
86-Pass `--install-path=/path/to/cann` to select a different CANN directory. See81+ The sample builds against the asc-comm headers installed into the CANN directory. If not yet
87-[`docs/en/guide/build_and_test.md`](../../docs/en/guide/build_and_test.md) for the full option list.82+ installed, generate and install the development verification run package from the repository root
83+ (see [`docs/en/guide/build_and_test.md`](../../../../docs/en/guide/build_and_test.md)):
88 84 
89-## Build85+ ```bash
86+ bash build.sh --pkg
87+ ./build_out/cann-asc-comm_1.0.0_linux-<arch>.run --full
88+ ```
90 89 
91-```bash90+ Then load the CANN environment:
92-cd examples/simt_urma
93-bash build.sh
94-```
95 91 
96-`build.sh` removes the existing `build` directory before reconfiguring and rebuilding, so stale92+ ```bash
97-headers and object files are never reused.93+ source ${install_path}/cann/set_env.sh
94+ ```
98 95 
99-The output is `examples/simt_urma/build/simt_urma`.96+ > **Note:** `${install_path}` is the CANN installation directory. If not specified, the default is
97+ > `/usr/local/Ascend`.
100 98 
101-## Run99+- Run the Sample
102 100 
103-The communicator is created with `HcclGetRootInfo` and `HcclCommInitRootInfo`, so no rank table is101+ The executable forks all rank processes internally:
104-required. `run.sh` creates a temporary root info file for each run, written by rank 0 and read by the
105-other ranks. Rank `i` runs on device `i`.
106 102 
107-```bash103+ ```bash
108-examples/simt_urma/run.sh <nranks> [mode]104+ mkdir -p build
109-```105+ cd build
106+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..
107+ make -j
108+ cd .. && ./build/simt_urma tcp://127.0.0.1:29620 <nranks> [mode]
109+ ```
110 110 
111-`nranks` may be any value >= 2: every rank sends and receives on the ring.111+ | Parameter | Default | Description |
112+ | --- | --- | --- |
113+ | `<ip:port>` | Required | Control channel address that rank 0 listens on (`tcp://ip:port`); use different ports for parallel instances |
114+ | `<nranks>` | Required | Number of ranks; must be >= 2 |
115+ | `[mode]` | `single` | Test mode; see Test Modes above; unrecognized modes are rejected outright |
112 116 
113-An unrecognized `mode` is rejected outright rather than silently falling back to the default.117+- Build Options
114 118 
115-```bash119+ | Option | Values | Description |
116-examples/simt_urma/run.sh 2 write_single120+ | --- | --- | --- |
117-examples/simt_urma/run.sh 2 write_batch_last121+ | `CMAKE_ASC_RUN_MODE` | `npu` (default) | Run mode; this sample supports NPU execution only |
118-examples/simt_urma/run.sh 2 write_value_single122+ | `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU architecture for Ascend 950PR/Ascend 950DT |
119-examples/simt_urma/run.sh 2 write_value_batch_last
120-examples/simt_urma/run.sh 2 notify
121-examples/simt_urma/run.sh 2 faa
122-examples/simt_urma/run.sh 2 cas
123-examples/simt_urma/run.sh 2 single
124-examples/simt_urma/run.sh 2 batch_last
125-examples/simt_urma/run.sh 2 notify_immediate_repeat
126-```
127 123 
128-Expected output (two cards, one group printed per rank):124+- Expected Output
129 125 
130-```text126+ ```text
131-[rank 0] simt_urma mode=<mode> sending to rank 1127+ [rank 0] simt_urma write_single | sent to rank 1, received from rank 1 | PASS
132-[rank 1] simt_urma mode=<mode> sending to rank 0128+ [rank 1] simt_urma write_single | sent to rank 0, received from rank 0 | PASS
133-[rank 0] simt_urma <mode> PASS | sent to rank 1, received from rank 1129+ RESULT | Example=simt_urma Mode=write_single | Status=PASS
134-[rank 1] simt_urma <mode> PASS | sent to rank 0, received from rank 0130+ ```
135-RESULT | Mode=<mode> | Status=PASS
136-```
137 131 
138-`run.sh` traps `SIGINT`/`SIGTERM` and kills the rank processes it started, so Ctrl+C does not leave132+ The sample succeeds when all ranks print `PASS` and the final status is `PASS`.
139-processes blocked on a host barrier still holding devices.
140 133 
141## Buffer Index Constraints134## Buffer Index Constraints
142 135 
@@ -153,8 +146,8 @@ receiver observes the CCL buffer's contents rather than the payload.
153 146 
154## Host-side Synchronization147## Host-side Synchronization
155 148 
156-Two host barriers make the verification trustworthy. They use marker files in a temporary directory149+Two host barriers make the verification trustworthy. They use the `RankSyncContext` barrier
157-created fresh for each run.150+primitives over the TCP control channel.
158 151 
159| Barrier | Purpose |152| Barrier | Purpose |
160|---|---|153|---|---|
@@ -165,14 +158,10 @@ The `sent` barrier is what makes the result trustworthy. Without it a rank runs
165to `aclrtMemcpy` with no ordering constraint against the kernel of the rank writing into it, reads158to `aclrtMemcpy` with no ordering constraint against the kernel of the rank writing into it, reads
166the zeroed buffer, and reports every slot as never landed even though all CQEs returned cleanly.159the zeroed buffer, and reports every slot as never landed even though all CQEs returned cleanly.
167 160 
168-The directory path is passed as a command-line argument rather than an environment variable: an
169-exported variable is easy to lose across a shell wrapper or a scheduler, and once lost every rank
170-silently falls back to the same default directory, where markers left by an earlier run can let a
171-barrier pass immediately.
172- 
173## Runtime Constraints161## Runtime Constraints
174 162 
175- Supported chips: Ascend 950PR / Ascend 950DT. CANN version 9.1.0 or later is required.163- Supported chips: Ascend 950PR / Ascend 950DT. CANN version 9.1.0 or later is required.
176-- At least two NPUs on the same host are required to run; single-NPU environments only support164+- At least two NPUs are required to run; single-NPU environments only support compilation verification.
177- compilation verification.165+- The executable captures `SIGINT`/`SIGTERM` and terminates all rank processes, so Ctrl+C does not
166+ leave processes blocked on a host barrier still holding devices.
178- The build relies on CANN ASC CMake utilities and links against the CANN `hcomm` library.167- The build relies on CANN ASC CMake utilities and links against the CANN `hcomm` library.
Rexamples/simt_urma/main.cpp→examples/aicore/hcomm/04_simt_urma/main.cpp+128-278
@@ -13,44 +13,29 @@
13 * \brief Hcomm SIMT URMA功能样例Host侧流程：通信域创建、内存注册、P2P通道创建与结果校验13 * \brief Hcomm SIMT URMA功能样例Host侧流程：通信域创建、内存注册、P2P通道创建与结果校验
14 */14 */
15 15 
16-#include <cerrno>
17#include <chrono>16#include <chrono>
18#include <cstdint>17#include <cstdint>
19#include <cstdio>18#include <cstdio>
20#include <cstdlib>19#include <cstdlib>
21#include <cstring>20#include <cstring>
22-#include <fstream>
23#include <iostream>21#include <iostream>
24#include <string>22#include <string>
25#include <thread>23#include <thread>
26#include <vector>24#include <vector>
27 25 
28-#include <sys/stat.h>
29-#include <sys/types.h>
30- 
31#include "hccl/hccl.h"26#include "hccl/hccl.h"
32#include "hccl/hccl_res.h"27#include "hccl/hccl_res.h"
33#include "hccl/hccl_types.h"28#include "hccl/hccl_types.h"
34#include "hccl/hccl_rank_graph.h"29#include "hccl/hccl_rank_graph.h"
35 30 
31+#include "utils/check.h"
32+#include "utils/hccl_comm_init.h"
33+#include "utils/process_manager.h"
34+#include "utils/rank_sync.h"
35+#include "utils/arg_parser.h"
36+ 
36#include "simt_urma_common.h"37#include "simt_urma_common.h"
37 38 
38-#define ACLCHECK(ret) \
39- do { \
40- if ((ret) != ACL_SUCCESS) { \
41- printf("acl interface return err %s:%d, retcode: %d \n", __FILE__, __LINE__, (ret)); \
42- return (ret); \
43- } \
44- } while (0)
45- 
46-#define HCCLCHECK(ret) \
47- do { \
48- if ((ret) != HCCL_SUCCESS) { \
49- printf("hccl interface return err %s:%d, retcode: %d \n", __FILE__, __LINE__, (ret)); \
50- return (ret); \
51- } \
52- } while (0)
53- 
54extern void LaunchSimtUrma(void* stream, uint64_t channel, uint32_t sendBufIdx, uint32_t recvBufIdx, uint32_t mode);39extern void LaunchSimtUrma(void* stream, uint64_t channel, uint32_t sendBufIdx, uint32_t recvBufIdx, uint32_t mode);
55 40 
56namespace {41namespace {
Rexamples/simt_urma/op_kernel.cpp→examples/aicore/hcomm/04_simt_urma/op_kernel.cpp+0-0
文件重命名但无更改。
Rexamples/simt_urma/simt_urma_common.h→examples/aicore/hcomm/04_simt_urma/simt_urma_common.h+0-0
文件重命名但无更改。
Rexamples/simt_urma_perftest/CMakeLists.txt→examples/aicore/hcomm/05_simt_urma_perftest/CMakeLists.txt+26-10
@@ -10,34 +10,50 @@
10 10 
11cmake_minimum_required(VERSION 3.16)11cmake_minimum_required(VERSION 3.16)
12 12 
13-set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture")13+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
14+ 
14find_package(ASC REQUIRED)15find_package(ASC REQUIRED)
15project(simt_urma_perftest LANGUAGES ASC CXX)16project(simt_urma_perftest LANGUAGES ASC CXX)
16 17 
17-set(CMAKE_CXX_STANDARD 17)
18-set(CMAKE_CXX_STANDARD_REQUIRED ON)
19set(CMAKE_POSITION_INDEPENDENT_CODE ON)18set(CMAKE_POSITION_INDEPENDENT_CODE ON)
20 19 
21set_source_files_properties(op_kernel.cpp PROPERTIES LANGUAGE ASC)20set_source_files_properties(op_kernel.cpp PROPERTIES LANGUAGE ASC)
22-add_executable(simt_urma_perftest main.cpp op_kernel.cpp)21+ 
22+add_executable(simt_urma_perftest
23+ main.cpp
24+ op_kernel.cpp
25+)
23 26 
24target_compile_definitions(simt_urma_perftest PRIVATE27target_compile_definitions(simt_urma_perftest PRIVATE
25 USE_SIMT28 USE_SIMT
26)29)
27-set_target_properties(simt_urma_perftest PROPERTIES LINKER_LANGUAGE ASC)
28 30 
29-target_include_directories(simt_urma_perftest BEFORE PRIVATE31+set_target_properties(simt_urma_perftest PROPERTIES
30- ${CMAKE_CURRENT_SOURCE_DIR}32+ LINKER_LANGUAGE ASC
31- $ENV{ASCEND_HOME_PATH}/include
32- $<$<COMPILE_LANGUAGE:ASC>:$ENV{ASCEND_HOME_PATH}/asc/include/adv_api>
33)33)
34+ 
35+target_include_directories(simt_urma_perftest PRIVATE
36+ ${CMAKE_CURRENT_SOURCE_DIR}
37+ ${CMAKE_CURRENT_SOURCE_DIR}/../..
38+ $ENV{ASCEND_HOME_PATH}/include
39+)
40+ 
34target_link_directories(simt_urma_perftest PRIVATE41target_link_directories(simt_urma_perftest PRIVATE
35 $ENV{ASCEND_HOME_PATH}/lib6442 $ENV{ASCEND_HOME_PATH}/lib64
36 $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/lib6443 $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/lib64
37)44)
45+ 
38target_link_libraries(simt_urma_perftest PRIVATE46target_link_libraries(simt_urma_perftest PRIVATE
39- hcomm hccl ascendcl acl_rt runtime c_sec dl pthread47+ hcomm
48+ hccl
49+ ascendcl
50+ acl_rt
51+ runtime
52+ c_sec
53+ dl
54+ pthread
40)55)
56+ 
41target_compile_options(simt_urma_perftest PRIVATE57target_compile_options(simt_urma_perftest PRIVATE
42 $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>58 $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
43 $<$<COMPILE_LANGUAGE:ASC>:--cce-auto-infer-kernel-type=false>59 $<$<COMPILE_LANGUAGE:ASC>:--cce-auto-infer-kernel-type=false>
Rexamples/simt_urma_perftest/README_en.md→examples/aicore/hcomm/05_simt_urma_perftest/README_en.md+67-35
Rexamples/simt_urma_perftest/main.cpp→examples/aicore/hcomm/05_simt_urma_perftest/main.cpp+137-207
Rexamples/simt_urma_perftest/op_kernel.cpp→examples/aicore/hcomm/05_simt_urma_perftest/op_kernel.cpp+0-0
Rexamples/simt_urma_perftest/simt_urma_perftest_common.h→examples/aicore/hcomm/05_simt_urma_perftest/simt_urma_perftest_common.h+0-0