多卡推理场景，模型权重分散在多张昇腾NPU卡上，前向推理的时候需要把中间激活值在各个卡之间传来传去。通信慢了，GPU/NPU的计算单元就空转等着。

HCCL（Huawei Collective Communication Library）是昇腾CANN提供的集合通信库，专门解决多NPU之间的数据传输问题。它实现了AllReduce、AllGather、ReduceScatter这些标准集合通信原语，底层走的是昇腾的HCCS（Huawei Compute Connectivity System）高速互联。

多卡推理的通信模式

大模型推理常见的并行策略有三种：TP（Tensor Parallelism，张量并行）、PP（Pipeline Parallelism，流水线并行）、EP（Expert Parallelism，专家并行）。这三种策略的通信模式不一样，用的HCCL接口也不一样。

并行策略	通信模式	使用的HCCL原语	通信频率
TP（张量并行）	每层Transformer都要AllReduce	AllReduce (SUM)	极高（每层2次）
PP（流水线并行）	层与层之间传激活值	P2P Send/Recv	中（层边界）
EP（专家并行）	MoE路由，AllGather+ReduceScatter	AllGather + ReduceScatter	高（每层MoE模块）

实际部署，大部分推理框架会混用。比如TP+PP：Transformer层内做张量并行（卡内通信），层间做流水线并行（卡间通信）。

HCCL的环境初始化

用HCCL之前，必须做初始化。这一步很多人容易漏，报错的时候一头雾水。

#include "hccl/hccl.h"
#include "acl/acl.h"

// 初始化ACL运行时
aclRet = aclInit(nullptr);
if (aclRet != ACL_SUCCESS) {
    printf("ACL init failed, ret=%d\n", aclRet);
    return -1;
}

// 获取NPU设备数量和当前进程对应的设备ID
int32_t device_id = 0;  // 当前进程绑定的NPU ID
aclRet = aclrtSetDevice(device_id);
if (aclRet != ACL_SUCCESS) {
    printf("Set device %d failed\n", device_id);
    return -1;
}

// 创建HCCL通信域配置
HcclComm comm;
HcclRootInfo root_info;

// 获取Root Info（ rank 0 进程执行）
if (rank == 0) {
    HcclGetRootInfo(&root_info);
    // 把root_info通过MPI/共享内存等方式广播给所有rank
    MPI_Bcast(&root_info, sizeof(root_info), MPI_BYTE, 0, MPI_COMM_WORLD);
} else {
    MPI_Bcast(&root_info, sizeof(root_info), MPI_BYTE, 0, MPI_COMM_WORLD);
}

// 初始化HCCL通信域
HcclCommInitRootInfo(1, &root_info, &comm);

printf("HCCL init success, rank=%d\n", HcclGetRankId(comm));

这段代码里，HcclGetRootInfo必须在rank 0上调用，然后把结果广播给所有其他rank。这是HCCL的初始化协议，不能省。如果漏了这一步，其他rank会一直阻塞在HcclCommInitRootInfo上。

AllReduce实战：张量并行的通信核心

张量并行的前向计算，每一层的输出需要AllReduce求和（因为输入被切分了，每个卡算一部分，需要汇总）。

以MatMul为例，权重按列切分（Column Parallel），前半部分输出在卡0，后半部分在卡1。需要AllReduce把两部分的输出拼起来。

#include "hccl/hccl.h"
#include "aclnn/aclnn.h"

// 假设有2张卡，rank 0和rank 1
// 每张卡上的MatMul输出是 [batch, seq_len, hidden/2]
// AllReduce后，每张卡拿到完整的 [batch, seq_len, hidden]

void TensorParallelMatMul(
    HcclComm comm,
    aclTensor* input,      // [batch, seq_len, hidden/2]
    aclTensor* weight,     // [hidden/2, hidden/2] (本卡权重)
    aclTensor* output      // [batch, seq_len, hidden/2] (本卡输出)
) {
    // 1. 本卡做MatMul
    aclnnMatMulGetWorkspaceSize(input, weight, output, &ws_size);
    void* workspace = aclrtMalloc(ws_size);
    aclnnMatMul(input, weight, output, workspace);

    // 2. AllReduce：把所有卡的output求和，再写回每张卡
    HcclAllReduce(
        /*sendbuf=*/output,                  // 本卡的输出
        /*recvbuf=*/output,                  // 接收缓冲区（可以in-place）
        /*count=*/batch * seq_len * (hidden / 2),
        /*datatype=*/HCCL_FLOAT16,          // 半精度
        /*op=*/HCCL_REDUCE_SUM,            // 求和操作
        /*comm=*/comm,
        /*stream=*/aclrtStreamDefault        // 默认流
    );

    // 3. 等待AllReduce完成
    aclrtSynchronizeStream(aclrtStreamDefault);

    // 现在output在两张卡上都是完整的结果了
}

// 调用示例
HcclComm comm;  // 假设已经初始化好
aclTensor* input = /* ... */;
aclTensor* weight = /* 本卡的权重切片 */;
aclTensor* output = aclCreateTensor(...);

TensorParallelMatMul(comm, input, weight, output);

这段代码是in-place的AllReduce（sendbuf和recvbuf是同一个指针）。HCCL支持in-place和out-of-place两种模式。推理场景一般用in-place，省显存。

AllGather + ReduceScatter：专家并行的通信对

MoE（Mixture of Experts）模型，专家网络分散在不同卡上。前向的时候，每个token需要把自己的激活值发给所有专家所在的卡（AllGather），专家算完后再把结果汇总回来（ReduceScatter）。

#include "hccl/hccl.h"

// 假设有4个专家，分布在4张卡上
// rank 0: 专家0,1   rank 1: 专家2,3   rank 2: 专家4,5   rank 3: 专家6,7

void ExpertParallelMoE(
    HcclComm comm,
    float* token_acts,      // 本卡token的激活值 [num_tokens, hidden]
    int num_tokens,
    int hidden,
    float* expert_output    // 输出 [num_tokens, hidden]
) {
    // 1. AllGather：把所有卡的token激活值收集到每张卡
    //    收集后，每张卡都有全部token的激活值
    int world_size = HcclGetWorldSize(comm);
    float* gathered_acts = (float*)aclrtMalloc(
        world_size * num_tokens * hidden * sizeof(float)
    );

    HcclAllGather(
        /*sendbuf=*/token_acts,
        /*recvbuf=*/gathered_acts,
        /*count=*/num_tokens * hidden,
        /*datatype=*/HCCL_FLOAT16,
        /*comm=*/comm,
        /*stream=*/aclrtStreamDefault
    );

    aclrtSynchronizeStream(aclrtStreamDefault);

    // 2. 本卡计算自己负责的专家
    //    （假设本卡负责专家0和1）
    float* local_expert_out = ComputeExpert(
        /*expert_id_start=*/rank * 2,
        /*expert_id_end=*/rank * 2 + 1,
        /*input=*/gathered_acts,   // 全部token激活值
        /*num_tokens=*/world_size * num_tokens,
        /*hidden=*/hidden
    );

    // 3. ReduceScatter：把各专家的输出按token归约到对应卡
    //    比如token 0~31的结果归约到rank 0，token 32~63归约到rank 1...
    HcclReduceScatter(
        /*sendbuf=*/local_expert_out,  // 本卡专家的输出
        /*recvbuf=*/expert_output,     // 本卡负责的token结果
        /*count=*/num_tokens * hidden, // 本卡负责的token数
        /*datatype=*/HCCL_FLOAT16,
        /*op=*/HCCL_REDUCE_SUM,
        /*comm=*/comm,
        /*stream=*/aclrtStreamDefault
    );

    aclrtSynchronizeStream(aclrtStreamDefault);

    // 现在expert_output里是本卡负责的token的最终结果
    aclrtFree(gathered_acts);
}

AllGather和ReduceScatter是一对。AllGather是"把所有卡的数据收集到每张卡"，ReduceScatter是"把所有卡的数据按归约操作分发到对应卡"。这两个操作合起来，实现了"全局计算+局部归约"的专家并行模式。

P2P通信：流水线并线的激活值传递

流水线并行，模型按层切分到不同卡上。前向的时候，卡0算完第0层的输出，需要传给卡1作为第1层的输入。这个用P2P（Point-to-Point）通信。

#include "hccl/hccl.h"

// 流水线并行：4层模型，4张卡，每张卡算1层
// 卡0 → 卡1 → 卡2 → 卡3

void PipelineStageForward(
    HcclComm comm,
    int rank,
    int world_size,
    float* layer_input,    // 本层输入 [batch, seq_len, hidden]
    int batch,
    int seq_len,
    int hidden
) {
    float* layer_output = (float*)aclrtMalloc(
        batch * seq_len * hidden * sizeof(float)
    );

    // 1. 本层计算
    ComputeTransformerLayer(layer_input, layer_output, rank);

    // 2. 如果不是最后一层，把输出发给下一层
    if (rank < world_size - 1) {
        HcclSend(
            /*buf=*/layer_output,
            /*count=*/batch * seq_len * hidden,
            /*datatype=*/HCCL_FLOAT16,
            /*peer_rank=*/rank + 1,   // 发给下一卡
            /*comm=*/comm,
            /*stream=*/aclrtStreamDefault
        );
    }

    // 3. 如果不是第一层，接收上一层的输入
    if (rank > 0) {
        HcclRecv(
            /*buf=*/layer_input,     // 覆盖成本层的输入
            /*count=*/batch * seq_len * hidden,
            /*datatype=*/HCCL_FLOAT16,
            /*peer_rank=*/rank - 1,   // 从上一卡收
            /*comm=*/comm,
            /*stream=*/aclrtStreamDefault
        );
    }

    aclrtSynchronizeStream(aclrtStreamDefault);

    // 下一轮：用layer_output作为输入，传给下层
    // （实际代码里这里会是循环，为了清晰这里只写一轮）
}

P2P通信需要显式指定peer_rank。发送方调HcclSend，接收方调HcclRecv，两张卡要配对。如果rank搞错了（比如发送方写成了rank+2），通信会卡死，因为接收方永远等不到数据。

性能调优的几个关键点

HCCL的性能调优，主要看三个指标：带宽利用率、延迟、同步开销。

1. 用异步通信隐藏延迟

HCCL的通信操作默认是异步的（放到ACL流里执行）。你可以先计算、后同步，把通信和计算重叠起来。

// 不好的写法：等通信完成再计算
HcclAllReduce(...);
aclrtSynchronizeStream(...);
ComputeNextLayer(...);

// 好的写法：通信和计算重叠
HcclAllReduce(...);  // 放到流里，立即返回
ComputeNextLayer(...);  // 计算和通信并行
aclrtSynchronizeStream(...);  // 用数据前再同步

2. 选择合适的通信域

HCCL支持创建多个通信域（Communicator）。如果模型有多个并行维度（比如TP+PP），可以给每个维度创建独立的通信域，避免干扰。

// 创建TP通信域（卡内高速通信）
HcclComm tp_comm;
HcclCommInitRootInfo(/*tp_size=*/4, /*...*/, &tp_comm);

// 创建PP通信域（卡间通信）
HcclComm pp_comm;
HcclCommInitRootInfo(/*pp_size=*/8, /*...*/, &pp_comm);

3. 用小包聚合减少通信次数

如果每次通信的数据量很小（比如<1MB），HCCS带宽利用率很低。可以把多个小包聚合成一个大数据包再通信。

// 不好的写法：频繁小包通信
for (int i = 0; i < 10; i++) {
    HcclAllReduce(small_tensor[i], ...);  // 每次10KB
}

// 好的写法：聚合后一次通信
ConcatTensors(small_tensors, large_tensor);  // 10×10KB = 100KB
HcclAllReduce(large_tensor, ...);
SplitTensors(large_tensor, small_tensors);  // 再拆回来

4. 检查HCCS链路状态

通信慢，有时候不是软件问题，是硬件链路没起来。用npu-smi看HCCS状态：

npu-smi info -t hccs

如果看到HCCS Link Status: Down，说明物理链路断了，需要检查NPU卡的互联线缆。这个问题在机房部署的时候很常见，软件层面看不出来，只会表现为通信带宽远低于理论值。

常见错误排查

错误1：HCCL_INIT_FAILED

原因：ACL运行时没初始化，或者NPU设备没设对。

解决：检查aclInit和aclrtSetDevice的返回值。

错误2：AllReduce卡死

原因：不同rank调用的AllReduce参数不一致（比如count不一样）。

解决：确保所有rank的AllReduce调用参数完全一致。

错误3：P2P通信超时

原因：发送方和接收方的peer_rank不匹配。

解决：画一张rank关系图，确保Send和Recv配对。

错误4：带宽远低于理论值

原因：小包通信，或者HCCS链路没起来。

解决：用小包聚合，或者用npu-smi检查HCCS状态。

---

仓库地址：https://atomgit.com/cann/hccl