Ascend C HCCL跨域同步接口
InterHcclGroupSync【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况Ascend 950PR/Ascend 950DT不支持Atlas A3 训练系列产品/Atlas A3 推理系列产品支持Atlas A2 训练系列产品/Atlas A2 推理系列产品不支持Atlas 200I/500 A2 推理产品不支持Atlas 推理系列产品AI Core不支持Atlas 推理系列产品Vector Core不支持Atlas 训练系列产品不支持功能说明用于等待一个跨通信域的通信任务完成。调用该接口后本通信域后续下发的通信任务均等待指定的srcGroupID通信域中的srcHandleID通信任务执行完成才开始执行。函数原型__aicore__ inline void InterHcclGroupSync(int8_t srcGroupID, HcclHandle srcHandleID)参数说明参数名输入/输出描述srcGroupID输入通信域编号即后续通信任务所等待的通信任务所在的通信域编号。srcHandleID输入通信任务即后续通信任务所等待的通信任务的标识HcclHandle。返回值说明无约束说明调用本接口前确保已调用过InitV2和SetCcTilingV2接口。本接口在AIC核或者AIV核上调用必须与对应的Prepare接口的调用核保持一致。一个通信域内所有Prepare接口和InterHcclGroupSync接口的总调用次数不能超过63。调用示例本示例构造一个通信融合算子该算子有1个输入xGM2个输出alltoallGM和allgatherGM另有tilingGM传入tiling数据。算子内有2个通信域首先通信域0对输入进行AlltoAll通信将结果输出至alltoallGM。当结果数据输出到alltoallGM完成后通信域1将该结果alltoallGM作为AllGather通信的输入并将通信结果输出至allgatherGM。extern C __global__ __aicore__ void alltoall_allgather_custom(GM_ADDR xGM, GM_ADDR alltoallGM, GM_ADDR allgatherGM, GM_ADDR tilingGM) { REGISTER_TILING_DEFAULT( AlltoAllAllGatherCustomTilingData); // AlltoAllAllGatherCustomTilingData为对应算子头文件定义的结构体 GET_TILING_DATA_WITH_STRUCT(AlltoAllAllGatherCustomTilingData, tilingData, tilingGM); GM_ADDR contextGM0 AscendC::GetHcclContext0(); GM_ADDR contextGM1 AscendC::GetHcclContext1(); Hccl hccl0; Hccl hccl1; HcclDataType dtype HcclDataType::HCCL_DATA_TYPE_FP16; const uint64_t dataCount 10U; const uint64_t strideCount 0U; if (AscendC::g_coreType AIV) { // 仅使用AIV核进行通信 hccl0.InitV2(contextGM0, tilingData); hccl1.InitV2(contextGM1, tilingData); hccl0.SetCcTilingV2(offsetof(AlltoAllAllGatherCustomTilingData, alltoallTiling)); hccl1.SetCcTilingV2(offsetof(AlltoAllAllGatherCustomTilingData, allgatherTiling)); // 通信域0下发1个AlltoAll任务 auto group0_handle hccl0.AlltoAll(xGM, alltoallGM, dataCount, dtype, strideCount); // 通信域1下发跨域依赖任务保证通信域1后续的AllGather任务在通信域0的AlltoAll执行结束后才开始执行 hccl1.InterHcclGroupSync(0, group0_handle); // 通信域1下发1个AllGather任务 auto group1_handle hccl1.AllGather(alltoallGM, allgatherGM, dataCount, dtype, strideCount); hccl0.Commit(group0_handle); hccl1.Commit(group1_handle); hccl0.Wait(group0_handle); hccl1.Wait(group1_handle); AscendC::SyncAlltrue(); // 全AIV核同步防止0核执行过快提前调用hccl.Finalize()接口导致其他核Wait卡死 hccl0.Finalize(); hccl1.Finalize(); } }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
