hw-native-sys/pto-isa/pto-comm-operator-develop

Host 侧                          Device 侧
┌─────────────────┐              ┌─────────────────────────┐
│ main.cpp        │              │ comm_kernel.cpp         │
│ - MPI 初始化    │   启动       │ - __global__ AICORE     │
│ - ACL 初始化    │──kernel──→   │ - TPUT/TGET/TNOTIFY/... │
│ - HCCL 通信域   │              │ - 信号同步逻辑          │
│ - 内存分配      │              ├─────────────────────────┤
│ - Kernel 启动   │              │ compute_kernel.cpp      │
│ - 结果验证      │   启动       │ - __global__ AICORE     │
│                 │──kernel──→   │ - TMATMUL/TADD/...      │
└─────────────────┘              │ - 计算逻辑              │
                                 └─────────────────────────┘

需要 NPU 间通信？
├── 仅需基本 P2P 传输
│   └── 使用 TPUT/TGET → 参考 "开发模式" 之 P2P 模式
│
├── 需要集合通信（AllReduce/AllGather/ReduceScatter 等）
│   ├── 可用内置集合指令完成？
│   │   └── 使用 TGATHER/TSCATTER/TBROADCAST/TREDUCE → 参考 "开发模式" 之集合通信模式
│   └── 需要自定义算法（如 RS+AG 组合 AllReduce）？
│       └── 使用 TPUT<AtomicAdd> + TNOTIFY/TWAIT 组合 → 参考 "开发模式" 之自定义集合通信
│
├── 需要通算融合（计算+通信重叠）
│   └── 使用双 kernel（cube + vec）+ 队列/信号同步 → 参考 "开发模式" 之通算融合模式
│
└── 需要异步大块传输
    └── 使用 TPUT_ASYNC/TGET_ASYNC → 参考 pto-comm-isa-reference

kernels/manual/<platform>/<operator_name>/
├── comm_kernel.cpp           # 通信 kernel（Vec 架构）
├── compute_kernel.cpp        # 计算 kernel（Cube 架构，如需融合）
├── config.h                  # Tiling 配置、Block 数量、常量定义
├── kernel_launchers.h        # Host 侧 kernel 启动函数声明
├── common.hpp                # 远端地址计算等共享工具
├── main.cpp                  # Host 侧：初始化、启动、验证
├── CMakeLists.txt            # 构建配置
├── run.sh                    # 运行脚本
└── README_zh.md              # 算子文档