CANN/asc-devkit HCCL AlltoAll接口文档

AlltoAll产品支持情况产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品xAtlas 推理系列产品 AI CorexAtlas 推理系列产品 Vector CorexAtlas 训练系列产品x功能说明集合通信AlltoAll的任务下发接口返回该任务的标识handleId给用户。AlltoAll的功能为每张卡向通信域内所有卡发送相同数据量的数据并从所有卡接收相同数据量的数据。结合原型中的参数描述接口功能具体为第j张卡接收到来自第i张卡的sendBuf中第j块数据并将该数据存放到本卡recvBuf中第i块的位置。函数原型template bool commit false __aicore__ inline HcclHandle AlltoAll(GM_ADDR sendBuf, GM_ADDR recvBuf, uint64_t dataCount, HcclDataType dataType, uint64_t strideCount 0, uint8_t repeat 1)参数说明表 1模板参数说明参数名输入/输出描述commit输入bool类型。参数取值如下true在调用Prepare接口时Commit同步通知服务端可以执行该通信任务。false在调用Prepare接口时不通知服务端执行该通信任务。表 2接口参数说明参数名输入/输出描述sendBuf输入源数据buffer地址。recvBuf输出目的数据buffer地址集合通信结果输出到此buffer中。dataCount输入本卡向通信域内其它每张卡收发的数据量单位为sizeof(dataType)。例如通信域内共4张卡每张卡的sendBuf中均有4个fp16的数据那么dataCount1。dataType输入AlltoAll操作的数据类型目前支持HcclDataType包含的全部数据类型HcclDataType详细可参考表1。strideCount输入多轮切分场景下一次AlltoAll任务中每张卡内参与通信的数据块间的间隔。默认值为0表示数据块内存连续。strideCount0每张卡内参与通信的数据块内存连续。卡rank_j收到来自卡rank_i的sendBuf中第j块数据且数据块间的偏移数据量为j*dataCount并将该数据存放于本卡recvBuf中第i块的位置且偏移数据量为i*dataCount。strideCount0每张卡内参与通信的相邻数据块的起始地址偏移数据量为strideCount。卡rank_j收到来自卡rank_i的sendBuf中第j块数据且数据块间的偏移数据量为j*strideCount并将该数据存放于本卡recvBuf中第i块的位置且偏移数据量为i*strideCount。注意上述的偏移数据量为数据个数单位为sizeof(dataType)。repeat输入一次下发的AlltoAll通信任务个数。repeat取值≥1默认值为1。当repeat1时每轮AlltoAll任务的sendBuf和recvBuf地址由服务端更新每一轮任务i的更新公式如下sendBuf[i] sendBuf dataCount * sizeof(datatype) * i, i∈[0, repeat)recvBuf[i] recvBuf dataCount * sizeof(datatype) * i, i∈[0, repeat)注意当设置repeat1时须与strideCount参数配合使用规划通信数据地址。返回值说明返回该任务的标识handleIdhandleId大于等于0。调用失败时返回 -1。约束说明调用本接口前确保已调用过InitV2和SetCcTilingV2接口。若HCCL对象的config模板参数未指定下发通信任务的核该接口只能在AIC核或者AIV核两者之一上调用。若HCCL对象的config模板参数中指定了下发通信任务的核则该接口可以在AIC核和AIV核上同时调用接口内部会根据指定的核的类型只在AIC核、AIV核二者之一下发该通信任务。对于Ascend 950PR/Ascend 950DT通信服务端为CCU时单次最大通信数据量不能超过256M。对于Ascend 950PR/Ascend 950DT一个通信域内所有Prepare接口的总调用次数不能超过63。对于Atlas A2 训练系列产品/Atlas A2 推理系列产品一个通信域内所有Prepare接口的总调用次数不能超过63。对于Atlas A3 训练系列产品/Atlas A3 推理系列产品一个通信域内所有Prepare接口和InterHcclGroupSync接口的总调用次数不能超过63。对于Atlas A3 训练系列产品/Atlas A3 推理系列产品一个通信域内最大支持128卡通信。调用示例非多轮切分场景4张卡执行AlltoAll通信任务。非多轮切分场景下每张卡上的数据块和数据量一致如下图中每张卡的A\B\C\D数据块数据量均为dataCount。图 1非多轮切分场景下4卡AlltoAll通信extern C __global__ __aicore__ void alltoall_custom(GM_ADDR xGM, GM_ADDR yGM, GM_ADDR workspaceGM, GM_ADDR tilingGM) { constexpr uint64_t dataCount 128U; // 数据量 auto sendBuf xGM; // xGM为AlltoAll的输入GM地址 auto recvBuf yGM; // yGM为AlltoAll的输出GM地址 REGISTER_TILING_DEFAULT(AllToAllCustomTilingData); //AllToAllCustomTilingData为对应算子头文件定义的结构体 GET_TILING_DATA_WITH_STRUCT(AllToAllCustomTilingData, 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(AllToAllCustomTilingData, alltoallCcTiling)); if (ret ! HCCL_SUCCESS) { return; } HcclHandle handleId hccl.AlltoAlltrue(sendBuf, recvBuf, dataCount, HcclDataType::HCCL_DATA_TYPE_FP16); hccl.Wait(handleId); AscendC::SyncAlltrue(); // AIV核全同步防止0核执行过快提前调用hccl.Finalize()接口导致其他核Wait卡死 hccl.Finalize(); } }多轮切分场景使能多轮切分等效处理上述非多轮切分示例的通信。在每张卡的数据均分成4块A\B\C\D的基础上将每一块继续切分若干块。本例中继续切分3块如下图所示被继续切分成的3块数据包括2个数据量为tileLen的数据块1个数据量为tailLen的尾块。切分后需要分3轮进行AlltoAll通信任务将等效上述非多轮切分的通信结果。图 23轮切分场景下4卡AlltoAll通信具体实现为第1轮通信每个rank上0-0\1-0\2-0\3-0数据块进行AlltoAll处理同一个卡上参与通信的相邻数据块的间隔为参数strideCount的取值。第2轮通信每个rank上0-1\1-1\2-1\3-1数据块进行AlltoAll处理。第3轮通信每个rank上0-2\1-2\2-2\3-2数据块进行AlltoAll处理。第1轮通信的图示及代码示例如下。图 3第一轮4卡AlltoAll示意图extern C __global__ __aicore__ void alltoall_custom(GM_ADDR xGM, GM_ADDR yGM, GM_ADDR workspaceGM, GM_ADDR tilingGM) { constexpr uint32_t tileNum 2U; // 首块数量 constexpr uint64_t tileLen 128U; // 首块数据个数 constexpr uint32_t tailNum 1U; // 尾块数量 constexpr uint64_t tailLen 100U; // 尾块数据个数 auto sendBuf xGM; // xGM为AlltoAll的输入GM地址 auto recvBuf yGM; // yGM为AlltoAll的输出GM地址 REGISTER_TILING_DEFAULT(AllToAllCustomTilingData); //AllToAllCustomTilingData为对应算子头文件定义的结构体 GET_TILING_DATA_WITH_STRUCT(AllToAllCustomTilingData, 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(AllToAllCustomTilingData, alltoallCcTiling)); if (ret ! HCCL_SUCCESS) { return; } uint64_t strideCount tileLen * tileNum tailLen * tailNum; // 2个首块处理 HcclHandle handleId1 hccl.AlltoAlltrue(sendBuf, recvBuf, tileLen, HcclDataType::HCCL_DATA_TYPE_FP16, strideCount, tileNum); // 1个尾块处理 constexpr uint32_t kSizeOfFloat16 2U; sendBuf tileLen * tileNum * kSizeOfFloat16; recvBuf tileLen * tileNum * kSizeOfFloat16; HcclHandle handleId2 hccl.AlltoAlltrue(sendBuf, recvBuf, tailLen, HcclDataType::HCCL_DATA_TYPE_FP16, strideCount, tailNum); for (uint8_t i0; itileNum; i) { hccl.Wait(handleId1); } hccl.Wait(handleId2); AscendC::SyncAlltrue(); // 全AIV核同步防止0核执行过快提前调用hccl.Finalize()接口导致其他核Wait卡死 hccl.Finalize(); } }创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考