master

cann-robotfeat(examples): add udma_perftest under shmem_perftest

f8011a9d创建于 6 小时前历史提交

文件	最后提交记录	最后更新时间
engine	feat(examples): add udma_perftest under shmem_perftest Co-authored-by: suqwe<sujianjia@huawei.com> # message auto-generated for no-merge-commit merge: !367 merge feat/udma-perftest into master feat(examples): add udma_perftest under shmem_perftest Created-by: suqwe Commit-by: suqwe Merged-by: cann-robot Description: ## 描述 新增 UDMA 引擎专用的性能测试样例 examples/shmem_perftest/udma_perftest/，平行于已有的 mte_perftest： - **新增样例 udma_perftest**：覆盖 aclshmemx_udma_put_nbi / aclshmemx_udma_get_nbi / aclshmemx_udma_put_signal_nbi 三个 UDMA 低阶接口，模式包含 put / bi_put / get / bi_get / put_signal，配套 main.cpp / udma_perftest_kernel.cpp / run.sh / CMakeLists.txt / README.md。 - **目录重命名**：mte_perftest -> shmem_perftest、内层 inner -> mte_perftest，把 MTE / UDMA / AscendC 三个样例统一收编到 shmem_perftest 下。 - **HBM-only 范围**：UDMA 引擎当前未对 Host 侧 DRAM 提供 RMA 路径，本样例仅测 HBM (DEVICE_SIDE)，不再支持 D2H / HOST_SIDE。 - **强制单核**：UDMA 不允许同 peer 多核并发，block_dim=1 固定，-b/--block-size、--block-range 仅形式上保留。 - **BW / Latency 双口径**： - --metric bw（默认）：prof_start → loop(*_nbi) → quiet → prof_end，窗口含 quiet。 - --metric lat：prof_start → loop(put_nbi) → prof_end → quiet，quiet 移到窗口外，仅测下发。 - 两种口径都是单 SHMEMI_PROF_START/END 包整段循环，再除以 loop_count，避免每次循环打点把 pipe_barrier 开销叠到延时数字上。 - **--batch 提交粒度（仅 BW 路径）**： - --batch 0（默认）等价于 --batch <loop_count>，全异步，仅末尾一次 quiet，反映稳态吞吐； - --batch 1 每次 *_nbi 后立刻 quiet，等价于同步提交，反映"提交+完成"端到端开销； - --batch N（1 < N < loop_count）每 N 次 *_nbi 一次 quiet，可观察 batch size 与吞吐的关系；loop_count % N != 0 时 prof_end 之前补一次 quiet。 - --metric lat 不受 --batch 影响。 - **put_signal 行为**：测试启动分配对称信号缓冲并初始化为 0，每次 put_signal_nbi 信号值线性递增；数据点结束后 host 端读回并校验 signal_base + warmup + loop_count - 1。 - **配套修复**：put_signal_nbi 内部传给 write_notify 的指针改为 typed pointer，避免 SOC 上 ABI 错位（src/device/gm2gm/engine/shmem_device_udma.hpp）。 ## 关联的Issue Fixes #311 ## 测试 - 在 Ascend950 上 bash scripts/build.sh -examples -soc_type Ascend950 通过；以 PR 当前 head 重新构建 udma_perftest target 链接成功。 - ./run.sh -t put -d float --exponent-range 8 17 --loop-count 1000 跑通；-t bi_put / get / bi_get / put_signal 各跑通，--metric lat -t put 跑通。 - --batch 1 / --batch 16 / --batch 1000 三档跑通，CSV 列与默认 bw 一致；CLI 校验：--batch -1 被 binary 拒绝，--batch abc 被 run.sh 正则拒绝，--metric lat + -t get 组合直接报错退出。 - put_signal 的远端 signal 槽校验通过（signal_base + warmup + loop_count - 1）。 - 非 Ascend950 SOC 上 device kernel 内置 aclshmemi_kernel_abort，按预期退出。 ## 文档更新 - 新增 examples/shmem_perftest/udma_perftest/README.md（含 CLI 参数表、metric 口径说明、--batch 章节、put_signal 行为说明、CSV 输出说明、已知约束）。 - 更新 examples/shmem_perftest/README.md 增加 udma 子目录索引。 - mte_perftest/README.md / ascendc_perftest/README.md 同步路径调整。 ## 类型标签 - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [x] 文档更新 - [ ] 其他，请描述： See merge request: cann/shmem!367	6 小时前
shmem_device_amo.hpp	原子加高阶接口移除quiet并添加文档说明 Co-authored-by: zhangyunqi<zhangyunqi5@huawei.com> # message auto-generated for no-merge-commit merge: !392 merge fix-opt-atomic-add into master 原子加高阶接口移除quiet并添加文档说明 Created-by: zhangyunqi Commit-by: zhangyunqi Merged-by: cann-robot Description: ## 描述 <!--在这里详细描述你的改动，包括改动的原因和所采取的方法。--> 原子加移除同步并添加文档说明 ## 关联的Issue <!-- 如果这个PR是为了解决特定的Issue，请在这里提供Issue链接。例如：关联Issue #123--> https://gitcode.com/cann/shmem/issues/276 ## 测试 <!--描述进行了哪些测试来验证你的改动。包括但不限于二级冒烟、算子泛化等。-->  ## 文档更新 <!--如果这个PR包含文档的更新，请在这里指出。例如：更新了README.md文件。--> ## 类型标签 <!-- [x] 表示选中 --> - [x] Bug修复 - [ ] 新特性 - [ ] 性能优化 - [x] 文档更新 - [ ] 其他，请描述： See merge request: cann/shmem!392	1 天前
shmem_device_mo.hpp	文件名aclshmem改为shmem Co-authored-by: caixilong<caixilong2@h-partners.com>	5 个月前
shmem_device_p2p_sync.hpp	Added memory detection information reporting Co-authored-by: huangxiaolan<huangxiaolan7@huawei.com> # message auto-generated for no-merge-commit merge: !272 merge add_message_report into master 补充 kernel 侧卡间 barrier 和 signal 语义上报 以及 host 侧共享内存信息上报 Created-by: huangxiaolan Commit-by: huangxiaolan Merged-by: cann-robot Description: ## 描述 <!--在这里详细描述你的改动，包括改动的原因和所采取的方法。--> 补充 kernel 侧卡间 barrier 和 signal 语义上报 以及 host 侧共享内存信息上报 ## 关联的Issue <!-- 如果这个PR是为了解决特定的Issue，请在这里提供Issue链接。例如：关联Issue #123--> <!-- 如果这个PR是为了解决特定的问题单，请在这里描述问题单单号。--> 关联issue #213 ## 测试 <!--描述进行了哪些测试来验证你的改动。包括但不限于二级冒烟、算子泛化等。--> ## 文档更新 <!--如果这个PR包含文档的更新，请在这里指出。例如：更新了README.md文件。--> 不涉及 ## 类型标签 <!-- [x] 表示选中 --> - [ ] Bug修复 - [ ] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [x] 其他，请描述： 补充工具信息上报 See merge request: cann/shmem!272	1 个月前
shmem_device_rma.hpp	UDMA场景补齐AMO接口 Co-authored-by: YeZZzzz1<yezhenni1@huawei.com> # message auto-generated for no-merge-commit merge: !284 merge master into master UDMA场景补齐AMO接口 Created-by: YeZZzzz1 Commit-by: YeZZzzz1 Merged-by: cann-robot Description: ## 描述 UMDA场景补齐AMO接口，包括：fetch, set, swap, fetch inc, inc, fetch and, and, fetch or, or, fetch xor, xor ## 关联的Issue https://gitcode.com/cann/shmem/issues/199 https://gitcode.com/cann/shmem/issues/205 ## 测试 examples： put  put signal   FAA  CAS  fetch  set  swap  fetch inc  fetch and  fetch or  fetch xor  uttests：         ## 文档更新 <!--如果这个PR包含文档的更新，请在这里指出。例如：更新了README.md文件。--> ## 类型标签 <!-- [x] 表示选中 --> - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/shmem!284	1 个月前
shmem_device_so.hpp	支持Ascend950 mte atomic Co-authored-by: vector5<caobingjie@huawei.com> Co-authored-by: zhangyunqi<zhangyunqi5@huawei.com> Co-authored-by: QK_25415<zhuzhiming1@huawei.com> # message auto-generated for no-merge-commit merge: !283 merge mteatomic into master 支持Ascend950 mte atomic Created-by: vector5 Commit-by: zhangyunqi;vector5;QK_25415 Merged-by: cann-robot Description: ## 描述 <!--在这里详细描述你的改动，包括改动的原因和所采取的方法。--> 支持Ascend950 mte atomic的以下接口和UT T aclshmemx_mte_atomic_fetch(gm T *src, int32_t pe); void aclshmemx_mte_atomic_set(gm T *dst, T value, int32_t pe); T aclshmemx_mte_atomic_compare_swap(gm T *dst, T cond, T value, int32_t pe); T aclshmemx_mte_atomic_swap(gm T *dst, T value, int32_t pe); void aclshmemx_mte_atomic_inc(gm T *dst, int32_t pe); void aclshmemx_mte_atomic_add(gm T *dst, T value, int32_t pe); T aclshmemx_mte_atomic_fetch_inc(gm T *dst, int32_t pe); T aclshmemx_mte_atomic_fetch_add(gm T *dst, T value, int32_t pe); ## 关联的Issue <!-- 如果这个PR是为了解决特定的Issue，请在这里提供Issue链接。例如：关联Issue #123--> <!-- 如果这个PR是为了解决特定的问题单，请在这里描述问题单单号。--> https://gitcode.com/cann/shmem/issues/220 ## 测试 <!--描述进行了哪些测试来验证你的改动。包括但不限于二级冒烟、算子泛化等。--> A2验证ShmemAtomic  MteAtomic  A5验证ShmemAtomic  MteAtomic  ## 类型标签 <!-- [x] 表示选中 --> - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/shmem!283	3 天前
shmemi_device_cc.h	feat(issue-280): [Task|任务]: namespace 整改 Co-authored-by: nino888<yinqiran1@huawei.com> # message auto-generated for no-merge-commit merge: !387 merge autodev/issue-280 into master feat(issue-280): [Task|任务]: namespace 整改 Created-by: nino888 Commit-by: nino888 Merged-by: cann-robot Description: ## Summary - Implement issue #280: [Task|任务]: namespace 整改 - Source issue: https://gitcode.com/cann/shmem/issues/280 - Branch: autodev/issue-280 (nino888/shmem -> cann/shmem) ## Changes - examples/dispatch_gmm_combine/include/dispatch_gmm_combine.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_init_routing_quant_v2.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_common.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_expert_token_out.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_fullload_dynamic_quant.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_fullload_quant.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_fullload_quant_base.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_gather_dynamic_quant.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_gather_quant.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_mrgsort.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_mrgsort_out.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_sort_base.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_sort_multi_core.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_sort_one_core.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_src_to_dst_and_gather.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_src_to_dst_op.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_src_to_dst_with_capacity.h - examples/dispatch_gmm_combine/include/moe_token_unpermute.h - examples/dispatch_gmm_combine/include/select_helper.h - examples/dispatch_gmm_combine/include/sync_util.h - examples/dynamic_tiling/impl/kernel/allgather_matmul.h - examples/dynamic_tiling/impl/kernel/allgather_matmul_padding.h - examples/dynamic_tiling/impl/kernel/allgather_matmul_with_gather_result.h - examples/dynamic_tiling/impl/kernel/matmul_allreduce.h - examples/dynamic_tiling/impl/kernel/matmul_reduce_scatter.h - examples/dynamic_tiling/impl/kernel/matmul_reduce_scatter_padding_a.h - examples/dynamic_tiling/impl/kernel/matmul_reduce_scatter_padding_ab.h - examples/dynamic_tiling/impl/kernel/matmul_reduce_scatter_padding_b.h - examples/matmul_allreduce/epilogue/block/epilogue_allreduce.hpp - src/device/gm2gm/shmemi_device_rma.cpp - src/host/bootstrap/shmemi_bootstrap_config_store.cpp - src/host/data_plane/shmem_host_rma.cpp - src/host/entity/mem_entity_default.cpp - src/host/entity/mem_entity_entry.cpp - src/host/init/shmem_init.cpp - src/host/mem/heap/hybm_vmm_based_segment.cpp - src/host/mem/shmem_rma.cpp - src/host/team/shmem_team.cpp - src/host/transport/transport_manager.cpp ## Local Validation - echo 'TODO: replace with real tests, e.g. pytest -q': passed See merge request: cann/shmem!387	16 小时前
shmemi_device_cc_kernel.cpp	移除必须0核执行限制 && 修复当前加上-enable_ascendc_dump编译报错的问题 Co-authored-by: vector5<caobingjie@huawei.com> # message auto-generated for no-merge-commit merge: !220 merge fixato into master 移除必须0核执行限制 && 修复当前加上-enable_ascendc_dump编译报错的问题 Created-by: vector5 Commit-by: vector5 Merged-by: cann-robot Description: ## 描述 <!--在这里详细描述你的改动，包括改动的原因和所采取的方法。--> 移除必须0核执行限制，8.3.RC1版本支持[[bisheng::core_ratio(0,1)]] ## 测试 <!--描述进行了哪些测试来验证你的改动。包括但不限于二级冒烟、算子泛化等。-->   8.2.RC1.alp003版本不支持__global__ __aicore__  8.2.RC1.alp003版本支持[[bisheng::core_ratio(1,1)]] 但不支持[[bisheng::core_ratio(0,1)]]  ## 类型标签 <!-- [x] 表示选中 --> - [x] Bug修复 See merge request: cann/shmem!220	2 个月前
shmemi_device_cc_kernel.h	aclshmem p2p sync interface Co-authored-by: Liwansi<liwansi@huawei.com> # message auto-generated for no-merge-commit merge: !79 merge p2p_sync into master aclshmem p2p sync interface Created-by: liwansi Commit-by: Liwansi Merged-by: cann-robot Description: ## 描述 对标nvshmem，支持P2P同步接口：wait类接口和test类接口 ## 关联的Issue Issue https://gitcode.com/cann/shmem/issues/57 ## 测试 新增UT：  所有UT：  ## 文档更新 /docs/pythonAPI.md ## 类型标签 <!-- [x] 表示选中 --> - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/shmem!79	4 个月前
shmemi_device_mo.h	文件名aclshmem改为shmem Co-authored-by: caixilong<caixilong2@h-partners.com>	5 个月前
shmemi_device_p2p_sync.h	include 目录下 aclshmemi_ 内容整改 Co-authored-by: zhu-wangyi<zhuwangyi@huawei.com> # message auto-generated for no-merge-commit merge: !297 merge bug_fix/sync_bit_rename into master include 目录下 aclshmemi_ 内容整改 Created-by: zhu-wangyi Commit-by: zhu-wangyi Merged-by: cann-robot Description: ## 描述 include 目录下 aclshmemi_ 内容整改 ## 关联的Issue https://gitcode.com/cann/shmem/issues/206 ## 测试 <!--描述进行了哪些测试来验证你的改动。包括但不限于二级冒烟、算子泛化等。--> ## 文档更新 <!--如果这个PR包含文档的更新，请在这里指出。例如：更新了README.md文件。--> ## 类型标签 <!-- [x] 表示选中 --> - [x] Bug修复 - [ ] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/shmem!297	30 天前
shmemi_device_p2p_sync_kernel.cpp	aclshmem p2p sync interface Co-authored-by: Liwansi<liwansi@huawei.com> # message auto-generated for no-merge-commit merge: !79 merge p2p_sync into master aclshmem p2p sync interface Created-by: liwansi Commit-by: Liwansi Merged-by: cann-robot Description: ## 描述 对标nvshmem，支持P2P同步接口：wait类接口和test类接口 ## 关联的Issue Issue https://gitcode.com/cann/shmem/issues/57 ## 测试 新增UT：  所有UT：  ## 文档更新 /docs/pythonAPI.md ## 类型标签 <!-- [x] 表示选中 --> - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/shmem!79	4 个月前
shmemi_device_p2p_sync_kernel.h	aclshmem p2p sync interface Co-authored-by: Liwansi<liwansi@huawei.com> # message auto-generated for no-merge-commit merge: !79 merge p2p_sync into master aclshmem p2p sync interface Created-by: liwansi Commit-by: Liwansi Merged-by: cann-robot Description: ## 描述 对标nvshmem，支持P2P同步接口：wait类接口和test类接口 ## 关联的Issue Issue https://gitcode.com/cann/shmem/issues/57 ## 测试 新增UT：  所有UT：  ## 文档更新 /docs/pythonAPI.md ## 类型标签 <!-- [x] 表示选中 --> - [ ] Bug修复 - [x] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/shmem!79	4 个月前
shmemi_device_rma.cpp	feat(issue-280): [Task|任务]: namespace 整改 Co-authored-by: nino888<yinqiran1@huawei.com> # message auto-generated for no-merge-commit merge: !387 merge autodev/issue-280 into master feat(issue-280): [Task|任务]: namespace 整改 Created-by: nino888 Commit-by: nino888 Merged-by: cann-robot Description: ## Summary - Implement issue #280: [Task|任务]: namespace 整改 - Source issue: https://gitcode.com/cann/shmem/issues/280 - Branch: autodev/issue-280 (nino888/shmem -> cann/shmem) ## Changes - examples/dispatch_gmm_combine/include/dispatch_gmm_combine.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_init_routing_quant_v2.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_common.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_expert_token_out.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_fullload_dynamic_quant.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_fullload_quant.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_fullload_quant_base.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_gather_dynamic_quant.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_gather_quant.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_mrgsort.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_mrgsort_out.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_sort_base.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_sort_multi_core.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_sort_one_core.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_src_to_dst_and_gather.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_src_to_dst_op.h - examples/dispatch_gmm_combine/include/moe_init_routing_quant_v2/moe_v2_src_to_dst_with_capacity.h - examples/dispatch_gmm_combine/include/moe_token_unpermute.h - examples/dispatch_gmm_combine/include/select_helper.h - examples/dispatch_gmm_combine/include/sync_util.h - examples/dynamic_tiling/impl/kernel/allgather_matmul.h - examples/dynamic_tiling/impl/kernel/allgather_matmul_padding.h - examples/dynamic_tiling/impl/kernel/allgather_matmul_with_gather_result.h - examples/dynamic_tiling/impl/kernel/matmul_allreduce.h - examples/dynamic_tiling/impl/kernel/matmul_reduce_scatter.h - examples/dynamic_tiling/impl/kernel/matmul_reduce_scatter_padding_a.h - examples/dynamic_tiling/impl/kernel/matmul_reduce_scatter_padding_ab.h - examples/dynamic_tiling/impl/kernel/matmul_reduce_scatter_padding_b.h - examples/matmul_allreduce/epilogue/block/epilogue_allreduce.hpp - src/device/gm2gm/shmemi_device_rma.cpp - src/host/bootstrap/shmemi_bootstrap_config_store.cpp - src/host/data_plane/shmem_host_rma.cpp - src/host/entity/mem_entity_default.cpp - src/host/entity/mem_entity_entry.cpp - src/host/init/shmem_init.cpp - src/host/mem/heap/hybm_vmm_based_segment.cpp - src/host/mem/shmem_rma.cpp - src/host/team/shmem_team.cpp - src/host/transport/transport_manager.cpp ## Local Validation - echo 'TODO: replace with real tests, e.g. pytest -q': passed See merge request: cann/shmem!387	16 小时前
shmemi_device_rma.h	include 目录下 aclshmemi_ 内容整改 Co-authored-by: zhu-wangyi<zhuwangyi@huawei.com> # message auto-generated for no-merge-commit merge: !297 merge bug_fix/sync_bit_rename into master include 目录下 aclshmemi_ 内容整改 Created-by: zhu-wangyi Commit-by: zhu-wangyi Merged-by: cann-robot Description: ## 描述 include 目录下 aclshmemi_ 内容整改 ## 关联的Issue https://gitcode.com/cann/shmem/issues/206 ## 测试 <!--描述进行了哪些测试来验证你的改动。包括但不限于二级冒烟、算子泛化等。--> ## 文档更新 <!--如果这个PR包含文档的更新，请在这里指出。例如：更新了README.md文件。--> ## 类型标签 <!-- [x] 表示选中 --> - [x] Bug修复 - [ ] 新特性 - [ ] 性能优化 - [ ] 文档更新 - [ ] 其他，请描述： See merge request: cann/shmem!297	30 天前