CANN/asc-devkit HCCL使用说明

HCCL使用说明【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkitAscend C提供一组HCCL通信类高阶API方便算子Kernel开发用户在AI Core侧灵活管理通算融合算子中计算与通信任务的执行顺序。HCCL为集合通信任务客户端主要对外提供了集合通信原语接口以下统称为Prepare接口对标集合通信C接口当前支持AllReduce、AllGather、ReduceScatter、AlltoAll、AlltoAllV、AlltoAllvWrite、BatchWrite接口。本章的所有接口运行在AI Core上且不执行通信任务而是由用户调用Prepare接口将对应类型的通信任务信息发送给AI CPU或CCU服务端并在合适的时机通过Commit接口通知AI CPU或CCU上的服务端执行对应的通信任务。 注意当前Ascend 950PR/Ascend 950DT上仅支持CCU服务端。所谓合适的时机取决于用户编排的是先通信后计算的任务还是先计算后通信的任务。对于这两种场景简述如下先通信后计算的任务典型的如AllGather通信Matmul计算任务编排。此场景下用户在调用AllGather接口下发通信任务之后通过AllGather接口返回的该通信任务标识handleId可立即调用Commit接口通知服务端执行该handleId对应的任务同时用户调用Wait阻塞接口等待服务端通知handleId对应的通信任务执行结束待该通信任务结束后再执行计算任务。先计算后通信的任务典型的如Matmul计算AllReduce通信任务编排。此场景下用户可以先调用AllReduce接口通知服务端先下发通信任务再调用Matmul计算接口进行计算这样AllReduce任务的组装和任务下发及执行过程可以被Matmul的计算流水所掩盖待计算任务完成后调用Commit接口通知服务端执行AllReduce任务无须调用Wait接口等待通信任务执行结束。当后续无通信任务时调用Finalize接口通知服务端后续无通信任务执行结束后退出客户端检测并等待最后一个通信任务执行结束。以上介绍的AI Core下发HCCL通信任务的机制如下图所示。图1AI Core下发HCCL通信任务机制图2Ascend 950PR/Ascend 950DT AI Core下发HCCL通信任务机制[!CAUTION]注意 对于Atlas A3 训练系列产品/Atlas A3 推理系列产品在AI CPU作为服务端的场景中HCCL通信API的功能依赖开放AI CPU用户态下发调度任务存在一定的安全风险用户需要自行确保AI Core自定义算子的安全可靠防止恶意攻击行为。实现AI Core下发一个通信任务的具体步骤如下创建HCCL对象并调用初始化接口InitV2。// 传initTiling地址的调用方式推荐使用该方式 GET_TILING_DATA_WITH_STRUCT( AllGatherCustomTilingData, tilingData, tilingGM); // AllGatherCustomTilingData为对应算子头文件定义的结构体 HcclHcclServerType::HCCL_SERVER_TYPE_AICPU hccl; // 通过模板入参的方式选择硬件类型 GM_ADDR contextGM GetHcclContext0(); // AscendC自定义算子kernel中通过此方式获取HCCL context hccl.InitV2(contextGM, tilingData);当调用InitV2接口时必须使用标准C语法定义TilingData结构体的开发方式。如上示例代码中的tilingGM为host侧传入的、作为核函数入参的算子TilingData的GM地址通过GET_TILING_DATA_WITH_STRUCT获取TilingData。调用InitV2初始化接口时需要传入通信上下文信息可以通过框架提供的获取通信上下文的接口GetHcclContext获取。设置对应通信算法的Tiling地址。通过SetCcTilingV2接口设置对应通信算法的Tiling地址调用Commit接口后该地址被发送到服务端由服务端解析。SetCcTilingV2接口必须与InitV2接口配合使用。示例如下。// 传initTiling地址的调用方式 GET_TILING_DATA_WITH_STRUCT(AllGatherCustomTilingData, tilingData, tilingGM); Hccl hccl; GM_ADDR contextGM GetHcclContext0(); // AscendC自定义算子kernel中通过此方式获取HCCL context hccl.InitV2(contextGM, tilingData); if (SetCcTilingV2(offsetof(AllGatherCustomTilingData, mc2CcTiling)) ! HCCL_SUCCESS) { return; }用户通过对应的Prepare接口异步下发对应类型的通信任务并获取到该任务的标识handleId服务端接收到后开始通信任务的展开和下发示例如下。auto handleId hccl.ReduceScatterfalse(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。表1HcclDataType参数说明| 数据类型 | 说明 | | --- | --- | | HcclDataType | HCCL任务的数据类型。HcclDataType为枚举类型定义如下代码所示。HCCL_DATA_TYPE_INT8int8HCCL_DATA_TYPE_INT16int16HCCL_DATA_TYPE_INT32int32HCCL_DATA_TYPE_FP16 half或float16HCCL_DATA_TYPE_FP32 floatHCCL_DATA_TYPE_INT64int64HCCL_DATA_TYPE_UINT64uint64HCCL_DATA_TYPE_UINT8uint8HCCL_DATA_TYPE_UINT16uint16HCCL_DATA_TYPE_UINT32uint32HCCL_DATA_TYPE_FP64float64HCCL_DATA_TYPE_BFP16bfloat16HCCL_DATA_TYPE_INT128int128预留类型暂不支持HCCL_DATA_TYPE_HIF8hif8HCCL_DATA_TYPE_FP8E4M3fp8e4m3HCCL_DATA_TYPE_FP8E5M2fp8e5m2HCCL_DATA_TYPE_FP8E8M0fp8e8m0HCCL_DATA_TYPE_RESERVED暂不支持使用 |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 */ };表2HcclReduceOp参数说明| 数据类型 | 说明 | | --- | --- | | HcclReduceOp | Reduce操作类型。HcclReduceOp为枚举类型定义如下代码所示。HCCL_REDUCE_SUMsumHCCL_REDUCE_PRODprodHCCL_REDUCE_MAXmaxHCCL_REDUCE_MINminHCCL_REDUCE_RESERVED暂不支持使用 |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 */ }用户调用Commit接口通知服务端执行handleId对应的通信任务。// 等待通信任务执行时机成熟调用Commit接口通知服务端执行 hccl.Commit(handleId);用户调用Wait阻塞接口等待服务端执行完对应的通信任务。auto ret hccl.Wait(handleId); // 对于Wait和Query接口在调试时可增加异常值校验和PRINTF打印 // if (ret HCCL_FAILED) { // PRINTF([ERROR] call Wait for handleId[%d] failed., handleId); // return; // } // 调用核间同步接口防止部分核执行较快退出触发Hccl析构影响执行较慢的核 // 开发者可根据实际的业务场景选择调用SyncAll、CrossCoreSetFlag、CrossCoreWaitFlag接口保证全部核的任务完成后再退出执行用户调用Finalize接口通知服务端后续无通信任务执行结束后退出客户端检测并等待最后一个通信任务执行结束。hccl.Finalize();注意若HCCL对象的模板参数未指定下发通信任务的核则Prepare接口仅能在AIC或AIV之一上运行调用步骤2到步骤5的接口前必须指定接口代码运行在AIC或AIV核上实现时如下代码所示。// 通过内置常量g_coreType来判断AIC核或者AIV核 if (g_coreType AIV) { // if (g_coreType AIC) { 调用HCCL接口 }基于以上对单个通信任务下发的了解介绍Prepare接口中repeat参数的灵活使用方式。一次Prepare接口的调用对应一个handleIdPrepare接口中的参数repeat代表这次Prepare的通信任务次数该值必须和针对该handleId调用Commit接口的次数、Wait接口的次数一致。以图3 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接口对应代码片段如下。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::GetHcclContext0(); // 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只调用了repeat1次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::SyncAlltrue(); // 全AIV核同步防止0核执行过快提前调用hccl.Finalize()接口导致其他核Wait卡死 hccl.Finalize(); } }由于每张卡上3份数据的源地址SendBuf是连续的且每张卡中目的地址recvBuf用来存储3份通信结果数据的内存也是连续的因此以上代码可以优化将ReduceScatter接口中的repeat参数设置为3从而调用1次ReduceScatter接口达到下发3个通信任务的效果。此时只有1个handleId的任务但是需要调用3次Commit和Wait接口对应代码片段如下。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::GetHcclContext0(); // 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::SyncAlltrue(); // 全AIV核同步防止0核执行过快提前调用hccl.Finalize()接口导致其他核Wait卡死 hccl.Finalize(); } }图3ReduceScatter通信示例表3MC2_BUFFER_LOCATION参数说明数据类型说明MC2_BUFFER_LOCATION预留参数。计算和通信中间结果的Buffer存放位置。用户在Tiling侧可设置该字段。提示调试含有HCCL高阶API的算子时在算子编译工程中增加编译选项-DASCENDC_DEBUG可以开启异常场景拦截的能力具体内容请参考并使用assert接口。【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考