Posted Updated 1-AIV35 minutes read (About 5251 words)

AIV Direct Drive Programming

导言

理解“AIV 直驱缩短了通信控制路径”之后，下一步不是立即抄一个 AllGather，而是先闭合一条最小链路：Host 建立资源并分配对称内存，启动一个 AIV Kernel；Kernel 把数据写到对端并发布 signal；接收端等待 signal，最后 Host 确认完成并按依赖顺序销毁资源。

本文基于 cann/shmem@382afa08efa801d7bca6c2645fd17e155111efcc 和 CANN 9.0.X / 9.0.0-beta.2 官方文档。代码是绑定该 revision 的教学归一化代码，不是承诺可在任意 CANN 版本直接编译的通用样例。

先选编程路径

“AIV 直驱”不是单一 API 名称。写代码前，先选清资源由谁管理。

路径 Host 看见什么 AIV 看见什么 适合起点
HCCL 自定义 AIV 通信算子 Engine Context、Notify 内存、Channel、HCCL Buffer、AIV binary GM/LM、远端 buffer、软同步标记、Ascend C 搬运与同步 需要手工控制拓扑和 HCCL 通信资源
cann/shmem 对称内存 bootstrap 属性、Team、对称堆、ACL stream PE、对称地址、Put/Get、Signal、Quiet、Barrier 学习最小 Device 侧单边通信闭环

CANN 9.0.0-beta.2 的通信算子 API 列表把接口分成控制面与数据面；固定 SHMEM revision 则用 shmem.h 条件聚合 Host 和 Device 公共头（行 15-50）。二者可以建立相近的心智模型，却不能拼成一套对象模型：SHMEM 用户不需要伪造一个 ChannelHandle。

![AIV 直驱最小编程闭环](https://pic.shaojiemike.top/shaojiemike/2026/08/d9e7e4741adcc61269331ffa32fefc38.png){ width=100% }
自绘技术图：Host 控制面负责初始化、对称分配、Kernel 下发与销毁；AIV 通过 PE、对称地址和 RMA 发起远端操作；quiet、signal/wait 与 stream synchronize 共同封闭完成语义。图中接口取自固定 cann/shmem revision，表达语义依赖而非跨版本 ABI 承诺。

这张图只表达依赖关系。SHMEM 初始化的实现还会建立 bootstrap、映射对称堆并同步 device state；这些细节由库封装，业务代码不应调用 shmemi_* 内部符号。固定源码：docs/principles/init_finalize.md:1-23, 333-361

先读懂六类对象

Channel 与 Notify

在 HCCL 手工资源路径中：

  • Channel 是跨 rank 的连接和远端资源可达关系，不是 payload。
  • Notify 是同步状态的承载资源。AIV 软同步会把标记放在通信可见内存中，通过 GM 与 LM 之间的数据搬运完成 Record/Wait。
  • HCCL Buffer 承载通信数据；Channel 让本端获得对端 buffer 或已注册 Notify 内存的描述。

官方创建资源页面按 Engine Context -> Notify 内存 -> HcclCommMemReg -> HcclChannelAcquire -> 查询本端/远端内存 展示流程；任务编排页面进一步说明 Record、Wait 和 GM-to-GM 搬运都会借助 Ascend C DataCopy 在 Local Memory 中转。

官方片段不等于完整工程

页面中的 CpGM2GM、Record1vN、WaitNv1 是示例封装；Context 大小、tag 管理、错误回滚和完整销毁也被省略。它们不能当作安装 CANN 后即可链接的基础 API。

对称内存与 PE

SHMEM 路径不让每次 Put 都携带远端虚拟地址描述，而是要求各 PE 建立布局一致的对称堆：

  1. 所有 PE 用相同参数参加 init，只有 my_pe 不同。
  2. 所有 PE 以相同顺序、相同大小调用对称分配和释放。
  3. Kernel 把“本地对称地址 + 目标 pe”交给 SHMEM，库据此换算远端对应地址。

这个不变量见固定源码 docs/principles/init_finalize.md:13-23, 70-74，公开分配接口见 include/host/mem/shmem_host_heap.h:21-60。

最小对象账本

下面的闭环只处理两个 PE 和 64 B payload。先把物理对象与指针视图分开，代码会更不容易写错。

对象 类型 所在域 容量 生命周期与约束
attr Host 配置 每个进程 一个结构体 init 前创建；除 my_pe 外保持一致
stream ACL 句柄 Host 一条 stream launch 前创建；free/finalize 前同步
payload 物理对称分配 Device GM 64 B 每个 PE 同序同大小分配；PE 0 是源，PE 1 的对应偏移是目标
signals 物理对称分配 Device GM 2 * 8 B 每个 signal slot 使用 ACLSHMEM_SIGNAL_SIZE；固定源码定义为 8 B
data_ready signals 的代数视图 AIV slot 0 PE 0 更新 PE 1；PE 1 在本地等待
ack signals 的代数视图 AIV slot 1 PE 1 更新 PE 0；PE 0 在本地等待
pe 代数索引 AIV 标量 由 aclshmem_my_pe() 查询；本例只接受 0 或 1

Signal slot 的固定 revision 定义见 include/host_device/shmem_common_types.h:142-144，PE 查询接口见 include/device/team/shmem_device_team.h:22-36。

最小 Host + Kernel 闭环

代码身份

以下代码不使用省略号，但仍是教学归一化代码。公共 SHMEM API 来自固定 revision；LaunchAivOneWay、CheckOrAbortAll 和 CMake target 是教学工程自己的符号。要编译，必须把 Kernel 与 wrapper 接入该 revision 的 bisheng/CMake 构建，并匹配实际 SoC、CANN、ops 包和驱动固件。

AIV Kernel

选择同步 typed put-with-signal，是因为公共头明确规定它先复制数据，再更新远端 signal，并且只支持单核调用。这个限制见 include/device/gm2gm/shmem_device_so.h:55-128。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
// aiv_one_way_kernel.cpp：教学归一化代码，绑定 shmem@382afa08...
#include "kernel_operator.h"
#include "shmem.h"

extern "C" [[bisheng::core_ratio(0, 1)]] __global__ __aicore__
void AivOneWay(GM_ADDR payload_addr, GM_ADDR signals_addr,
               uint32_t payload_bytes)
{
    if (GetBlockIdx() != 0 || aclshmem_n_pes() != 2) {
        return;
    }

    auto *payload = reinterpret_cast<__gm__ uint8_t *>(payload_addr);
    auto *signal_bytes = reinterpret_cast<__gm__ uint8_t *>(signals_addr);
    auto *data_ready = reinterpret_cast<__gm__ int32_t *>(signal_bytes);
    auto *ack = reinterpret_cast<__gm__ int32_t *>(
        signal_bytes + ACLSHMEM_SIGNAL_SIZE);

    const int pe = aclshmem_my_pe();
    if (pe == 0) {
        // dst 和 data_ready 都是本地对称地址；最后一个参数选择远端 PE 1。
        aclshmem_uint8_put_signal(
            payload, payload, payload_bytes,
            data_ready, 1, ACLSHMEM_SIGNAL_SET, 1);

        // 发送端只有收到 ACK 才允许 Kernel 结束。
        aclshmem_signal_wait_until(ack, ACLSHMEM_CMP_EQ, 1);
    } else {
        // put-with-signal 保证 signal 发布在 payload 复制之后。
        aclshmem_signal_wait_until(data_ready, ACLSHMEM_CMP_EQ, 1);

        // 通知 PE 0：PE 1 已经越过 data-ready 边界。
        aclshmemx_signal_op(ack, 1, ACLSHMEM_SIGNAL_SET, 0);
        aclshmem_quiet();
    }
}

void LaunchAivOneWay(uint32_t block_dim, void *stream,
                     uint8_t *payload, uint8_t *signals,
                     uint32_t payload_bytes)
{
    AivOneWay<<<block_dim, nullptr, stream>>>(
        payload, signals, payload_bytes);
}

固定仓库采用相同的 AIV Kernel 声明与 <<<block_dim, nullptr, stream>>> wrapper 形态。examples/allgather/allgather_kernel.cpp:288-330 中的 allgather_demo 和 examples/udma_demo/udma_demo_kernel.cpp:66-105 中的 launch 函数都是示例 wrapper，不是 SHMEM 公共 API。

Host 生命周期

下面固定 n_pes = 2，每个进程由 launcher 传入 pe_id、device_id 和 rank 0 的 tcp://host:port。正常路径展示完整销毁顺序；异常路径采用集群 fail-fast，要求 launcher 同时终止两个 PE，以免只有一个 PE 退出、另一个卡在 collective。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
// main.cpp：教学归一化代码，绑定 shmem@382afa08...
#include <acl/acl.h>
#include <algorithm>
#include <array>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include "shmem.h"

extern void LaunchAivOneWay(uint32_t block_dim, void *stream,
                            uint8_t *payload, uint8_t *signals,
                            uint32_t payload_bytes);

constexpr int kPeCount = 2;
constexpr uint32_t kPayloadBytes = 64;
constexpr uint64_t kSymmetricHeapBytes = 1UL << 20;

[[noreturn]] void CheckOrAbortAll(const char *stage, int code)
{
    std::fprintf(stderr, "%s failed, code=%d; launcher must stop both PEs\n",
                 stage, code);
    std::abort();
}

void Check(const char *stage, int code, int success)
{
    if (code != success) {
        CheckOrAbortAll(stage, code);
    }
}

int main(int argc, char **argv)
{
    if (argc != 4) {
        std::fprintf(stderr,
            "usage: %s <pe_id:0|1> <device_id> <tcp://rank0-host:port>\n",
            argv[0]);
        return 2;
    }

    const int pe_id = std::atoi(argv[1]);
    const int device_id = std::atoi(argv[2]);
    if (pe_id < 0 || pe_id >= kPeCount) {
        return 2;
    }

    Check("aclInit", aclInit(nullptr), ACL_SUCCESS);
    Check("aclrtSetDevice", aclrtSetDevice(device_id), ACL_SUCCESS);

    aclrtStream stream = nullptr;
    Check("aclrtCreateStream", aclrtCreateStream(&stream), ACL_SUCCESS);

    aclshmemx_init_attr_t attr{};
    attr.my_pe = pe_id;
    attr.n_pes = kPeCount;
    attr.local_mem_size = kSymmetricHeapBytes;
    const int written = std::snprintf(
        attr.ip_port, sizeof(attr.ip_port), "%s", argv[3]);
    if (written <= 0 || static_cast<size_t>(written) >= sizeof(attr.ip_port)) {
        CheckOrAbortAll("ip_port", ACLSHMEM_INVALID_VALUE);
    }

    Check("aclshmemx_init_attr",
          aclshmemx_init_attr(ACLSHMEMX_INIT_WITH_DEFAULT, &attr),
          ACLSHMEM_SUCCESS);

    auto *payload = static_cast<uint8_t *>(
        aclshmem_calloc(1, kPayloadBytes));
    auto *signals = static_cast<uint8_t *>(
        aclshmem_calloc(2, ACLSHMEM_SIGNAL_SIZE));
    if (payload == nullptr || signals == nullptr) {
        CheckOrAbortAll("aclshmem_calloc", ACLSHMEM_MALLOC_FAILED);
    }

    std::array<uint8_t, kPayloadBytes> host_data{};
    if (pe_id == 0) {
        host_data.fill(0x5A);
        Check("payload H2D",
              aclrtMemcpy(payload, kPayloadBytes,
                          host_data.data(), host_data.size(),
                          ACL_MEMCPY_HOST_TO_DEVICE),
              ACL_SUCCESS);
    }

    LaunchAivOneWay(1, stream, payload, signals, kPayloadBytes);
    Check("aclrtSynchronizeStream",
          aclrtSynchronizeStream(stream), ACL_SUCCESS);

    if (pe_id == 1) {
        Check("payload D2H",
              aclrtMemcpy(host_data.data(), host_data.size(),
                          payload, kPayloadBytes,
                          ACL_MEMCPY_DEVICE_TO_HOST),
              ACL_SUCCESS);
        if (!std::all_of(host_data.begin(), host_data.end(),
                         [](uint8_t x) { return x == 0x5A; })) {
            CheckOrAbortAll("payload validation", ACLSHMEM_INNER_ERROR);
        }
    }

    // 所有 PE 必须以相同顺序释放；这里按分配的逆序释放。
    aclshmem_free(signals);
    aclshmem_free(payload);
    Check("aclshmem_finalize", aclshmem_finalize(), ACLSHMEM_SUCCESS);
    Check("aclrtDestroyStream", aclrtDestroyStream(stream), ACL_SUCCESS);
    Check("aclrtResetDevice", aclrtResetDevice(device_id), ACL_SUCCESS);
    Check("aclFinalize", aclFinalize(), ACL_SUCCESS);
    return 0;
}

这里直接构造 aclshmemx_init_attr_t，依赖固定 revision 的默认 option_attr 为 MTE；结构定义和公开 init/finalize 入口见 include/host/shmem_host_def.h:145-195 与 include/host/init/shmem_host_init.h:94-208。仓库的 test_set_attr 位于 examples/utils/utils.h:141-160，只是样例 helper，不能当作公共 API。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
sequenceDiagram
    participant H0 as Host PE 0
    participant V0 as AIV PE 0
    participant V1 as AIV PE 1
    participant H1 as Host PE 1
    H0->>V0: launch one AIV core
    H1->>V1: launch one AIV core
    V0->>V1: put payload, then set data_ready
    V1->>V1: wait data_ready == 1
    V1->>V0: set ack = 1, quiet
    V0->>V0: wait ack == 1
    V0-->>H0: kernel return + stream sync
    V1-->>H1: kernel return + stream sync
    H0->>H0: free, finalize, destroy ACL resources
    H1->>H1: free, finalize, destroy ACL resources

基础 API 地图

Host 侧

API 作用 完成或集体性 证据
aclshmemx_init_attr 按 bootstrap 属性建立 SHMEM 资源 所有 PE 对称参加；参数布局一致 include/host/init/shmem_host_init.h:137-147
aclshmemx_get_uniqueid、aclshmemx_set_attr_uniqueid_args 使用 UID 模式准备属性 UID 的跨 PE 广播由应用负责 同文件行 94-118
aclshmem_malloc/calloc/align/free 管理对称堆 各 PE 同序同大小 include/host/mem/shmem_host_heap.h:21-60
aclshmem_barrier(_all) Host 参与者会合并完成此前 CPU 侧 remote updates 不替 NPU 侧完成 include/host/data_plane/shmem_host_cc.h:28-49
aclshmemx_barrier_all_on_stream 把 barrier 排入指定 stream 用于 stream 顺序，但仍需按版本确认场景 同文件行 51-68
aclshmem_finalize 释放当前 instance 所有 PE 对称调用；之后不得再调用 SHMEM API docs/principles/init_finalize.md:70-74

AIV 侧

API 作用 最容易误解的边界 证据
aclshmem_my_pe/n_pes 查询本 PE 与参与者数 PE 是进程/通信参与者，不是 AIV 核 ID include/device/team/shmem_device_team.h:22-36
aclshmem_*_put/get 同步 Put/Get 远端操作数必须是对称地址；RDMA 对两端范围要求更严 include/device/gm2gm/shmem_device_rma.h:154-173, 315-335
aclshmem_*_put_nbi/get_nbi 非阻塞发起 RMA 返回不代表可消费或可复用 同文件行 477-525, 670-715
aclshmem_putmem_signal、typed variants Put 后更新远端 signal 只支持单核或单 writer include/device/gm2gm/shmem_device_so.h:55-128
aclshmemx_signal_op 远端 signal SET/ADD 分离的数据移动必须先建立 data-before-signal include/device/gm2gm/shmem_device_p2p_sync.h:23-36
aclshmem_signal_wait_until 阻塞等待本地 signal 条件 看见 signal 不自动证明任意独立 RMA 已完成 同文件行 38-51
aclshmem_quiet 完成本 PE 在 NPU 侧发起的 SHMEM 操作 Host 仍需 stream/device synchronize include/device/gm2gm/shmem_device_mo.h:23-35
aclshmem_barrier PE 会合，并完成此前 remote updates 全核 API 对 MIX Kernel 有限制；CPU/NPU 完成域分离 include/device/gm2gm/shmem_device_cc.h:15-77
aclshmem_sync、aclshmemx_sync_vec 同步普通 memory stores 不完成 SHMEM remote updates 同文件行 79-124

固定 revision 中 aclshmem_fence 因当前硬件实现与 quiet 相同，同时提供排序和完成；这只是 include/device/gm2gm/shmem_device_mo.h:37-48 的 revision 行为，不能外推为所有版本或其他 SHMEM 实现的规范。

完成语义

把“函数返回”拆成四级，能避开大多数直驱错误：

  1. Issued：*_nbi 已经发起，源和目标仍可能在使用。
  2. Device complete：aclshmem_quiet 让调用 PE 的 NPU 侧 SHMEM 操作完成。
  3. Remote consumable：生产者通过 put-with-signal，或先 quiet 再 signal；消费者 wait 到匹配 signal 后才能越过数据依赖。
  4. Host observed：Host 通过 aclrtSynchronizeStream 或 device synchronize 确认 Kernel 完成。
1
2
3
4
5
6
7
stateDiagram-v2
    [*] --> Issued: put_nbi / get_nbi
    Issued --> DeviceComplete: quiet
    DeviceComplete --> Published: signal
    Published --> Consumable: remote wait satisfied
    Consumable --> HostObserved: stream synchronize
    HostObserved --> [*]: free / finalize

三个错误等式

  • _nbi return == remote data ready：错误，仍需 quiet 或等价完成机制。
  • signal visible == 任意 payload 完成：错误，必须先建立 data-before-signal。
  • sync == barrier == quiet：错误；固定 revision 的 vec sync 明确不完成 SHMEM remote updates。

固定仓库 allgather 小数据路径在 examples/allgather/allgather_kernel.cpp:263-285 中展示了 put_nbi -> quiet -> SyncAll -> signal/wait -> get_nbi。但公共同步头又提示不要在同一 Kernel 混用 ACLSHMEM inter-PE synchronization 与 SyncAll。本文优先遵循公共头限制，最小例只用一个 AIV 核和 SHMEM P2P 同步，不机械复制该多核示例。

Tiling 与下发

SHMEM 直启

固定仓库示例把元素数、buffer 指针和 FFTS 地址作为普通 Kernel 参数，wrapper 直接使用 <<<block_dim, nullptr, stream>>>。这种写法的“tiling”只是应用自己计算参数和 block 数，不等于框架注册的 TilingFunction。

框架算子 Tiling

框架自定义算子通常由 Host TilingFunction 计算 TilingData，Kernel 再用 GET_TILING_DATA 取出。CANN 9.0.X 的 GET_TILING_DATA 官方页面明确写明：该宏当前不支持 Kernel Launch 工程。

迁移原则：

  1. 先让直启最小闭环用普通参数跑通。
  2. 接入框架时新增算子注册、Host TilingFunction 与 TilingData 序列化。
  3. 不要在未经目标 CANN 版本验证的直启工程中直接加入 GET_TILING_DATA。

HCCL AIV binary 下发

CANN 9.0.0-beta.2 的算子下发页面给出另一条版本化路径：查询 Vector Core 资源，加载独立 Kernel .o，取得函数句柄，设置 ACL_RT_ENGINE_TYPE_AIV，再调用 aclrtLaunchKernelWithHostArgs。该片段还要求 Device binary 带 AIV meta section。

这条路径适合 HCCL 自定义通信算子，但页面没有给出完整参数 ABI、binary 构建 target 与失败回滚，不能把片段改名后当作完整 launch helper。

手工 Channel 路线

如果业务必须接管 HCCL 通信资源，Host 侧最小资源关系是：

1
2
3
4
5
6
7
8
9
10
11
12
flowchart TD
    COMM[已有 HCCL 通信域] --> CTX[HcclEngineCtxCreate: AIV Notify 区]
    CTX --> ZERO[aclrtMemset 清零]
    ZERO --> REG[HcclCommMemReg]
    REG --> DESC[HcclChannelDescInit + remoteRank]
    DESC --> CH[HcclChannelAcquire]
    CH --> LOCAL[HcclGetHcclBuffer]
    CH --> REMOTE[HcclChannelGetHcclBuffer]
    CH --> TAG[HcclChannelGetRemoteMems]
    LOCAL --> BIN[加载并下发 AIV binary]
    REMOTE --> BIN
    TAG --> BIN

这里要严格区分：

  • HcclEngineCtxCreate、HcclCommMemReg、HcclChannelAcquire 和查询函数列在 beta.2 公开 API 页面中，属于版本化公开接口。
  • CpGM2GM、Record1vN、WaitNv1 是文档示例封装。
  • HcommEndpointCreate/HcommChannelCreate 是另一组较新的基础资源 API；不能在没有产品与 CANN 版本验证时替换 legacy HCCL Channel 路线。
  • 官方页面提供 HcclEngineCtxDestroy，但没有在同一流程中给出与 HcclChannelAcquire 配对的完整 Channel 回收代码。销毁必须服从目标版本通信域的所有权规则，本文不猜测缺失接口。

错误与销毁边界

必须拒绝的输入

  • n_pes 与实际启动进程数不一致，或 my_pe 越界。
  • 各 PE 的 local_mem_size、bootstrap 模式、对称分配顺序或大小不一致。
  • Put/Get 的远端地址不在对称堆中；启用 RDMA 时，完整源/目标范围不满足对称内存约束。
  • 多个 AIV 核并发调用只支持单核的 put-with-signal。
  • free 或 finalize 时仍有 Kernel 或 stream 在访问对称对象。
  • 一个 PE 局部返回，其他 PE 继续进入 collective。

正常销毁偏序

1
2
3
4
5
6
7
停止新 launch
    -> aclrtSynchronizeStream
    -> 所有 PE 同序 aclshmem_free
    -> 所有 PE aclshmem_finalize
    -> aclrtDestroyStream
    -> aclrtResetDevice
    -> aclFinalize

固定仓库的直接通信示例采用 free -> finalize -> destroy stream -> reset device -> aclFinalize，见 examples/udma_demo/main.cpp:123-135；更强的不变量是 finalize 前 Kernel 已完成，finalize 在 device reset 和 ACL finalize 前，所有 PE 对称参与。

失败路径不能只写 goto cleanup

多 PE 程序中，本地 cleanup 可能本身就是 collective。生产实现要用 launcher、MPI 或独立 control plane 汇总错误，再让全部 PE 一致退出；不能让一个 PE 执行 aclshmem_free/finalize，另一个 PE 已经返回。

编译环境

固定 revision 的 docs/quickstart.md:28-75, 195-229, 323-347给出以下边界：

  • 硬件：Atlas 800I/800T A2/A3、Ascend 950；Host 为 aarch64 或 x86。
  • 工具链：gcc/g++ 不低于 7.3 且版本一致，CMake 不低于 3.19，GLIBC 不低于 2.28，Python 不低于 3.9；Device Kernel 使用 CANN toolkit 随附的 bisheng。
  • 环境：先加载 CANN set_env.sh；源码构建后加载 SHMEM install/set_env.sh。
  • 版本：HDK 25.0.RC1.1 配 CANN 8.5 以上可用 MTE/RDMA；CANN 9.0.0-beta.2 以上增加表内 SDMA 能力；HDK 26.0.RC1 配 CANN 9.1 以上时，A3 可用 SDMA，Ascend 950 可用 UDMA。实际通路仍受 SoC 与 ops 包限制。
  • 构建：A2/A3 示例使用 bash scripts/build.sh -examples；Ascend 950 使用 bash scripts/build.sh -soc_type Ascend950 -examples。

本文没有宣称已经编译

当前草稿在 macOS 本地完成，只做了固定源码与 Markdown 静态核验；本机没有 Ascend NPU、CANN toolkit、匹配 ops 包和 bisheng。要把教学例变成可执行样例，仍需接入 pinned SHMEM 的 example/CMake target，在目标 SoC 上完成双 PE 编译、运行、结果校验和失败注入。

仓库 README 也明确指出 examples 只供学习参考，生产使用前需完成功能与性能测试，并建议锁定 CANN 版本。README.md:260-264

总结

最基础的 AIV 直驱程序可以压缩为五个不变量：

  1. Host 先建资源：ACL device/stream、bootstrap、Team 与对称堆都在 Kernel 前完成。
  2. 地址必须可解释：SHMEM 用“对称地址 + PE”，HCCL 手工路线用 Channel 交换的远端资源。
  3. 数据与状态分开：Put/Get 搬 payload，Signal/Notify 发布依赖；二者必须有明确顺序。
  4. 完成分层：NBI 发起、Device quiet、远端可消费、Host stream observed 是不同边界。
  5. 销毁服从偏序：先完成 Kernel，再同序 free 和 collective finalize，最后释放 ACL 资源。

最小闭环跑通后，才适合增加多核切分、框架 Tiling、RDMA/SDMA/UDMA 引擎、Team 集合通信与性能流水。否则，复杂示例中的每一个 SyncAll、signal offset 和 block 数都可能掩盖真正的生命周期错误。

参考资料

  1. cann/shmem 固定源码，commit 382afa08efa801d7bca6c2645fd17e155111efcc
  2. SHMEM 对外头文件与库，docs/public-headers-and-libraries.md:21-73
  3. SHMEM 初始化与终止，docs/principles/init_finalize.md:1-108, 333-361
  4. CANN 9.0.0-beta.2：通信算子开发 API 列表
  5. CANN 9.0.0-beta.2：AIV 创建资源
  6. CANN 9.0.0-beta.2：AIV 任务编排
  7. CANN 9.0.0-beta.2：AIV 算子下发
  8. CANN 9.0.X：GET_TILING_DATA

AIV Direct Drive Programming

http://icarus.shaojiemike.top/2026/08/25/Work/Artificial Intelligence/AIInfra/AIV-Direct-Drive-Programming/

Author

Shaojie Tan

Posted on

2026-08-25

Updated on

2026-08-25

Licensed under

#

Like this article? Support the author with

Follow

Categories

Subscribe for updates

follow.it

Links

Recents

Archives

Tags

0/1 Adam1
1-bit Adam1
1-bit LAMB1
3FS1
4+1 View Model1
5G1
64bits vs 32bits1
ADR1
AGRS1
AI10
AI AGENT1
AI Agent1
AI Cluster1
AI Engineering1
AI INFRA1
AI Infrastructure1
AI STARTUP1
AI Systems1
AIInfra1
AISBench2
×