HCCL使用说明

Ascend C提供一组HCCL通信类高阶API，方便算子Kernel开发用户在AI Core侧灵活管理通算融合算子中计算与通信任务的执行顺序。使用HCCL通信类高阶API进行算子开发前，请参考通算融合章节了解必要的背景知识。

HCCL为集合通信任务客户端，主要对外提供了集合通信原语接口（以下统称为Prepare接口），对标集合通信C++接口，详细可参见HCCL接口参考，当前支持AllReduce、AllGather、ReduceScatter、AlltoAll接口等。本章的所有接口运行在AI Core上，且不执行通信任务，而是由用户调用Prepare接口将对应类型的通信任务信息发送给AI CPU或CCU服务端，并在合适的时机通过Commit接口通知AI CPU或CCU上的服务端执行对应的通信任务。注意，当前Ascend 950PR/Ascend 950DT上仅支持CCU服务端。

所谓合适的时机，取决于用户编排的是先通信后计算的任务，还是先计算后通信的任务。对于这两种场景，简述如下：

当后续无通信任务时，调用Finalize接口，通知服务端后续无通信任务，执行结束后退出，客户端检测并等待最后一个通信任务执行结束。以上介绍的AI Core下发HCCL通信任务的机制如下图所示。

图1 AI Core下发HCCL通信任务机制
图2 Ascend 950PR/Ascend 950DT AI Core下发HCCL通信任务机制

对于 Atlas A3 训练系列产品 / Atlas A3 推理系列产品 ，在AI CPU作为服务端的场景中，HCCL通信API的功能依赖开放AI CPU用户态下发调度任务，存在一定的安全风险，用户需要自行确保AI Core自定义算子的安全可靠，防止恶意攻击行为。

实现AI Core下发一个通信任务的具体步骤如下：

  1. 创建HCCL对象，并调用初始化接口InitV2。 

    1
2
3
4
5
6
7
    		 
    // 传initTiling地址的调用方式，推荐使用该方式
GET_TILING_DATA_WITH_STRUCT(AllGatherCustomTilingData, tilingData, tilingGM); // AllGatherCustomTilingData为对应算子头文件定义的结构体

Hccl<HcclServerType::HCCL_SERVER_TYPE_AICPU> hccl; // 通过模板入参的方式选择硬件类型
GM_ADDR contextGM = GetHcclContext<0>();  // AscendC自定义算子kernel中，通过此方式获取HCCL context

hccl.InitV2(contextGM, &tilingData);

    当调用InitV2接口时，必须使用标准C++语法定义TilingData结构体的开发方式，具体请参考使用标准C++语法定义Tiling结构体。如上示例代码中的tilingGM为host侧传入的、作为核函数入参的算子TilingData的GM地址，通过GET_TILING_DATA_WITH_STRUCT获取TilingData。调用InitV2初始化接口时，需要传入通信上下文信息，可以通过框架提供的获取通信上下文的接口GetHcclContext获取。

  2. 设置对应通信算法的Tiling地址。

    通过SetCcTilingV2接口设置对应通信算法的Tiling地址，调用Commit接口后，该地址被发送到服务端由服务端解析。SetCcTilingV2接口必须与InitV2接口配合使用。示例如下。 
     1
 2
 3
 4
 5
 6
 7
 8
 9
10
    		 
    // 传initTiling地址的调用方式
GET_TILING_DATA_WITH_STRUCT(AllGatherCustomTilingData, tilingData, tilingGM);

Hccl hccl;
GM_ADDR contextGM = GetHcclContext<0>();  // AscendC自定义算子kernel中，通过此方式获取HCCL context

hccl.InitV2(contextGM, &tilingData);
if (SetCcTilingV2(offsetof(AllGatherCustomTilingData, mc2CcTiling)) != HCCL_SUCCESS) {
  return;
}

  3. 用户通过对应的Prepare接口异步下发对应类型的通信任务，并获取到该任务的标识handleId，服务端接收到后开始通信任务的展开和下发，示例如下。 

    1
2
3
4
5
6
7
8
    		 
    auto handleId = hccl.ReduceScatter<false>(aGM, cGM, recvCount, 
                                          AscendC::HCCL_DATA_TYPE_FP16,
                                          HCCL_REDUCE_SUM, strideCount, 1);
// 对于Prepare接口，在调试时可增加异常值校验和PRINTF打印
// if (handleId == INVALID_HANDLE_ID) {
// 	PRINTF("[ERROR] call ReduceScatter failed, handleId is -1.");
//	return;
// }

    示例的Prepare接口为ReduceScatter，其他接口可参考后续章节的内容。其中的参数AscendC::HCCL_DATA_TYPE_FP16是HCCL任务的数据类型，其数据结构为HcclDataType，对应的参数说明参考表1；参数HCCL_REDUCE_SUM是一种Reduce操作，AllReduce和ReduceScatter归约操作支持的Reduce操作类型参见表2。

    表1 HcclDataType参数说明

    数据类型

    说明

    HcclDataType

    HCCL任务的数据类型。

     1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
    		 
    enum HcclDataType {
HCCL_DATA_TYPE_INT8 = 0,   /* int8 */
HCCL_DATA_TYPE_INT16 = 1,  /* int16 */
HCCL_DATA_TYPE_INT32 = 2,  /* int32 */
HCCL_DATA_TYPE_FP16 = 3,   /* half或float16 */
HCCL_DATA_TYPE_FP32 = 4,   /* float */
HCCL_DATA_TYPE_INT64 = 5,  /* int64 */
HCCL_DATA_TYPE_UINT64 = 6, /* uint64 */
HCCL_DATA_TYPE_UINT8 = 7,  /* uint8 */
HCCL_DATA_TYPE_UINT16 = 8, /* uint16 */
HCCL_DATA_TYPE_UINT32 = 9, /* uint32 */
HCCL_DATA_TYPE_FP64 = 10,  /* float64 */
HCCL_DATA_TYPE_BFP16 = 11, /* bfloat16 */
HCCL_DATA_TYPE_INT128 = 12, /* int128 预留类型，暂不支持 */
HCCL_DATA_TYPE_HIF8 = 14,  /* hif8 */
HCCL_DATA_TYPE_FP8E4M3 = 15,  /* fp8e4m3 */
HCCL_DATA_TYPE_FP8E5M2 = 16,  /* fp8e5m2 */
HCCL_DATA_TYPE_FP8E8M0 = 17,  /* fp8e8m0 */
HCCL_DATA_TYPE_RESERVED    /* reserved */
}
    表2 HcclReduceOp参数说明

    数据类型

    说明

    HcclReduceOp

    Reduce操作类型。

    1
2
3
4
5
6
7
    		 
    enum HcclReduceOp {
HCCL_REDUCE_SUM = 0,  /* sum */
HCCL_REDUCE_PROD = 1, /* prod */
HCCL_REDUCE_MAX = 2,  /* max */
HCCL_REDUCE_MIN = 3,  /* min */
HCCL_REDUCE_RESERVED  /* reserved */
}

  4. 用户调用Commit接口通知服务端执行handleId对应的通信任务。 

    1
2
    		 
    // 等待通信任务执行时机成熟，调用Commit接口通知服务端执行
hccl.Commit(handleId);

  5. 用户调用Wait阻塞接口，等待服务端执行完对应的通信任务。 

    1
2
3
4
5
6
7
8
9
    		 
    auto ret = hccl.Wait(handleId);
// 对于Wait和Query接口，在调试时可增加异常值校验和PRINTF打印
// if (ret == HCCL_FAILED) {
//	PRINTF("[ERROR] call Wait for handleId[%d] failed.", handleId);
//	return;
// }

// 调用核间同步接口，防止部分核执行较快退出，触发Hccl析构，影响执行较慢的核
// 开发者可根据实际的业务场景，选择调用、、接口，保证全部核的任务完成后再退出执行

  6. 用户调用Finalize接口，通知服务端后续无通信任务，执行结束后退出；客户端检测并等待最后一个通信任务执行结束。 

    1
    		 
    hccl.Finalize();

注意：若HCCL对象的模板参数未指定下发通信任务的核，则Prepare接口仅能在AIC或AIV之一上运行，调用步骤2到步骤5的接口前，必须指定接口代码运行在AIC或AIV核上，实现时如下代码所示。

1
2
3
4
5
 
// 通过内置常量g_coreType来判断AIC核或者AIV核
if (g_coreType == AIV) {
// if (g_coreType == AIC) {
        调用HCCL接口
}

基于以上对单个通信任务下发的了解，介绍Prepare接口中repeat参数的灵活使用方式。一次Prepare接口的调用对应一个handleId，Prepare接口中的参数repeat代表这次Prepare的通信任务次数，该值必须和针对该handleId调用Commit接口的次数、Wait接口的次数一致。以图2 ReduceScatter通信示例进行说明，假设共4张卡，每张卡上源数据首先按照rankId均匀分成4份，每份数据被切分成3份，最终被切分后的每份数据的个数为TileLen，每次ReduceScatter通信仅通信一组切分数据（如图中数据0-0、1-0、2-0、3-0为一组切分数据），因此需要做3次ReduceScatter操作，全部数据才能通信完。

这种场景下，可以调用3次repeat参数为1的ReduceScatter接口，下发3个通信任务，同时更新每个ReduceScatter任务的收发地址，得到3个任务的handleId，每个handleId任务调用1次Commit和Wait接口，对应代码片段如下。 
 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
 
extern "C" __global__ __aicore__ void reduce_scatter_custom(GM_ADDR xGM, GM_ADDR yGM, GM_ADDR workspaceGM, GM_ADDR tilingGM) {
    auto sendBuf = xGM;  // xGM为ReduceScatter的输入GM地址
    auto recvBuf = yGM;  // yGM为ReduceScatter的输出GM地址
    constexpr size_t rankSize = 4U; // 4张卡
    constexpr size_t tileCnt = 3U;  // 卡上的数据均匀分成rankSize份，且每份又被切分成3份
    constexpr size_t tileLen = 100U;  // 被切分后的每份数据个数
    uint64_t strideCount = tileLen*tileCnt;  // 表示sendBuf上相邻数据块间的起始地址的偏移量

    REGISTER_TILING_DEFAULT(ReduceScatterCustomTilingData); //ReduceScatterCustomTilingData为对应算子头文件定义的结构体
    GET_TILING_DATA_WITH_STRUCT(ReduceScatterCustomTilingData, tilingData, tilingGM);

    Hccl hccl;
    GM_ADDR contextGM = AscendC::GetHcclContext<0>();  // AscendC自定义算子kernel中，通过此方式获取HCCL context
    if (AscendC::g_coreType == AIV) {  // 指定AIV核通信   
        hccl.InitV2(contextGM, &tilingData);
        auto ret = hccl.SetCcTilingV2(offsetof(ReduceScatterCustomTilingData, reduceScatterCcTiling));
        if (ret != HCCL_SUCCESS) {
          return;
        }
	// for循环中生成了3个handleId，每个handleId只调用了repeat=1次Commit和Wait接口
        for (int i = 0; i < tileCnt; ++i) {
	    auto handleId = hccl.ReduceScatter(sendBuf, recvBuf, tileLen, HcclDataType::HCCL_DATA_TYPE_FP32, HcclReduceOp::HCCL_REDUCE_SUM, strideCount, 1); //具体参数参见ReduceScatter接口说明
	    hccl.Commit(handleId);
	    auto ret = hccl.Wait(handleId);
	    // 执行其他计算逻辑 ....
			
	    // 更新ReduceScatter的收发地址
	    sendBuf += tileLen * sizeof(float);
	    recvBuf += tileLen * sizeof(float);
	}
        AscendC::SyncAll<true>();  // 全AIV核同步，防止0核执行过快，提前调用hccl.Finalize()接口，导致其他核Wait卡死   
        hccl.Finalize();
    }
}

由于每张卡上3份数据的源地址SendBuf是连续的，且每张卡中目的地址recvBuf用来存储3份通信结果数据的内存也是连续的，因此以上代码可以优化，将ReduceScatter接口中的repeat参数设置为3，从而调用1次ReduceScatter接口，达到下发3个通信任务的效果。此时，只有1个handleId的任务，但是需要调用3次Commit和Wait接口，对应代码片段如下。

 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
 
extern "C" __global__ __aicore__ void reduce_scatter_custom(GM_ADDR xGM, GM_ADDR yGM, GM_ADDR workspaceGM, GM_ADDR tilingGM) {
    auto sendBuf = xGM;  // xGM为ReduceScatter的输入GM地址
    auto recvBuf = yGM;  // yGM为ReduceScatter的输出GM地址
    constexpr size_t rankSize = 4U; // 4张卡
    constexpr size_t tileCnt = 3U;  // 卡上的数据均匀分成rankSize份，且每份又被切分成3份
    constexpr size_t tileLen = 100U;  // 被切分后的每份数据个数
    uint64_t strideCount = tileLen*tileCnt;  // 表示sendBuf上相邻数据块间的起始地址的偏移量

    REGISTER_TILING_DEFAULT(ReduceScatterCustomTilingData); //ReduceScatterCustomTilingData为对应算子头文件定义的结构体
    GET_TILING_DATA_WITH_STRUCT(ReduceScatterCustomTilingData, tilingData, tilingGM);

    Hccl hccl;
    GM_ADDR contextGM = AscendC::GetHcclContext<0>();  // AscendC自定义算子kernel中，通过此方式获取HCCL context
    if (AscendC::g_coreType == AIV) {  // 指定AIV核通信   
        hccl.InitV2(contextGM, &tilingData);
        auto ret = hccl.SetCcTilingV2(offsetof(ReduceScatterCustomTilingData, reduceScatterCcTiling));
        if (ret != HCCL_SUCCESS) {
          return;
        }
		
	auto handleId = hccl.ReduceScatter(sendBuf, recvBuf, tileLen, HcclDataType::HCCL_DATA_TYPE_FP32, HcclReduceOp::HCCL_REDUCE_SUM, strideCount, tileCnt); //具体参数参见ReduceScatter接口说明
	for (int i = 0; i < tileCnt; ++i) {
	    hccl.Commit(handleId);
	    auto ret = hccl.Wait(handleId);
	    // 执行其他计算逻辑 ....
	}
        AscendC::SyncAll<true>();  // 全AIV核同步，防止0核执行过快，提前调用hccl.Finalize()接口，导致其他核Wait卡死   
        hccl.Finalize();
    }
}
图3 ReduceScatter通信示例
表3 MC2_BUFFER_LOCATION参数说明

数据类型

说明

MC2_BUFFER_LOCATION

预留参数。计算和通信中间结果的Buffer存放位置。用户在Tiling侧可设置该字段。

提示：调试含有HCCL高阶API的算子时，在算子编译工程中增加编译选项-DASCENDC_DEBUG，可以开启异常场景拦截的能力，具体内容请参考并使用assert接口。

父主题： HCCL Kernel侧接口