Iterate【免费下载链接】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 训练系列产品不支持功能说明在某些算法下一次完整的集合通信任务可以细分为多个步骤对每个步骤的数据完成点对点的通信任务称为细粒度通信。以通信算法AlltoAlllevel0:fullmesh;level1:pairwise、通信步长为1的AlltoAllV通信任务为例这里参数level0代表配置Server昇腾AI Server通常是8卡或16卡的昇腾NPU设备组成的服务器形态的统称内通信算法参数level1代表配置Server间通信算法fullmesh为全连接通信算法pairwise为逐对通信算法如下图所示该示例展示了AlltoAllV通信的所有待发送数据、每一步通信完成后各卡收到的数据。图1使用pairwise算法的AlltoAllV通信步骤示意图在通算融合算子中通过调用本接口结合对应的Prepare原语获取通信算法每一步的输入或输出让计算、通信实现更精细粒度的流水排布从而获得更好的性能收益。函数原型template bool sync true __aicore__ inline int32_t Iterate(HcclHandle handleId, uint16_t* seqSlices, uint16_t seqSliceLen)参数说明表1模板参数说明参数名输入/输出描述sync输入bool类型。是否需要等待当前通信步骤完成再进行后续计算或通信任务参数取值如下true默认值表示阻塞并等待当前通信步骤完成。该参数取值为true时无需再调用Wait接口等待通信任务完成。false表示不等待当前通信步骤完成。表2接口参数说明参数名输入/输出描述handleId输入对应通信任务的标识ID只能使用Prepare原语接口的返回值。using HcclHandle int8_t;seqSlices输出由用户申请的栈空间用于保存当前通信步骤的输入或输出数据块的索引下标。在先计算后通信场景该参数返回当前通信步骤需要的输入数据块索引在先通信后计算场景该参数返回当前通信步骤的输出数据块索引。seqSliceLen输入seqSlices数组的长度。根据算法的通信步长及算法逻辑取每一步通信需要保存的数据块索引个数为该数组长度。返回值说明当通信任务未结束时在先计算后通信场景返回值为当前通信步骤需要的输入数据块数量与参数seqSliceLen数值相同。在先通信后计算场景返回值为当前通信步骤产生的输出数据块数量与参数seqSliceLen数值相同。当通信任务结束后返回值为0。约束说明调用本接口前确保已调用过InitV2和SetCcTilingV2接口。入参handleId只能使用Prepare原语对应接口的返回值。本接口当前支持的通信算法为AlltoAlllevel0:fullmesh;level1:pairwise。调用示例extern C __global__ __aicore__ void alltoallv_custom( GM_ADDR sendBuf, GM_ADDR recvBuf, GM_ADDR workspaceGM, GM_ADDR tilingGM) { // 指定AIV核通信 if (AscendC::g_coreType ! AIV) { return; } constexpr uint32_t RANK_NUM 4U; constexpr uint32_t STEP_SIZE 1U; // 细粒度通信步长通常使用SetStepSize接口设置示例代码简化成1 constexpr uint64_t sendCounts[RANK_NUM][RANK_NUM] {{3, 3, 3, 3}, {2, 2, 3, 2}, {1, 4, 4, 4}, {3, 3, 3, 3}}; constexpr uint64_t sDisplacements[RANK_NUM][RANK_NUM] {{0, 3, 6, 9}, {0, 2, 4, 7}, {0, 1, 5, 9}, {0, 3, 6, 9}}; constexpr uint64_t recvCounts[RANK_NUM][RANK_NUM] {{3, 2, 1, 3}, {3, 2, 4, 3}, {3, 3, 4, 3}, {3, 2, 4, 3}}; constexpr uint64_t rDisplacements[RANK_NUM][RANK_NUM] {{0, 3, 5, 6}, {0, 3, 5, 9}, {0, 3, 6, 10}, {0, 3, 5, 9}}; HcclDataType dtype HcclDataType::HCCL_DATA_TYPE_FP16; REGISTER_TILING_DEFAULT(AllToAllVCustomTilingData); // AllToAllVCustomTilingData为对应算子头文件定义的结构体 GET_TILING_DATA_WITH_STRUCT(AllToAllVCustomTilingData, tilingData, tilingGM); GM_ADDR contextGM AscendC::GetHcclContext0(); // AscendC自定义算子kernel中通过此方式获取HCCL context Hccl hccl; hccl.InitV2(contextGM, tilingData); auto ret hccl.SetCcTilingV2(offsetof(AllToAllVCustomTilingData, alltoallvCcTiling)); if (ret ! HCCL_SUCCESS) { return; } const uint32_t selfRankId hccl.GetRankId(); // 当通信任务为AlltoAlllevel0:fullmesh;level1:pairwise时 // 1. 每步通信产生的数据块数量等于STEP_SIZE // 2. 总的通信步数为RANK_NUM/STEP_SIZE*repeat uint16_t sliceInfo[STEP_SIZE]; if (TILING_KEY_IS(1000UL)) { // 通算融合中的“先通信后计算”场景即每一步都是先通信再将通信的输出作为计算的输入并执行计算 const auto handleId hccl.AlltoAllVtrue( sendBuf, sendCounts[selfRankId], sDisplacements[selfRankId], dtype, recvBuf, recvCounts[selfRankId], rDisplacements[selfRankId], dtype); // 模板参数sync true表示该接口会阻塞等待每一步通信结果并将输出数据块的下标索引填入sliceInfo中 while (hccl.Iteratetrue(handleId, sliceInfo, sizeof(sliceInfo) / sizeof(sliceInfo[0]))) { // 每一步通信的输出数据块的下标索引保存在sliceInfo中可以插入相应的计算流程实现细粒度的通算融合 } // Iterate已经会阻塞等待因此不再需要Wait // hccl.Wait(handleId); } else if (TILING_KEY_IS(1001UL)) { // 通算融合中的“先计算后通信”场景即每一步都是先计算再将计算的结果作为通信的输入并提交通信事务 const uint8_t tileNum 2U; const auto handleId hccl.AlltoAllVfalse( sendBuf, sendCounts[selfRankId], sDisplacements[selfRankId], dtype, recvBuf, recvCounts[selfRankId], rDisplacements[selfRankId], dtype, tileNum); for (uint8_t i 0; i tileNum; i) { for (uint8_t j 0; j RANK_NUM; j) { // 模板参数sync false表示该接口不会阻塞等待只会将当前这一步通信的输入数据块填入sliceInfo中 if (hccl.Iteratefalse(handleId, sliceInfo, sizeof(sliceInfo) / sizeof(sliceInfo[0])) 0) { break; } // sliceInfo表示相对地址偏移需要结合sDisplacements进行GM地址的运算保证通信的输入正确 // 计算完之后需要核间同步再通过Commit接口通知服务端进行集合通信 hccl.Commit(handleId); } } for (uint8_t i 0; i tileNum * RANK_NUM; i) { hccl.Wait(handleId); } } AscendC::SyncAlltrue(); hccl.Finalize(); }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
)
)
