CANN asc-devkit AlltoAllV API

AlltoAllV【免费下载链接】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 推理系列产品x功能说明集合通信AlltoAllV的任务下发接口返回该任务的标识handleId给用户。AlltoAll是AlltoAllV的一个子集AlltoAll要求所有卡的收发数据量相同AlltoAllV则不需要数据量相同使用上更加灵活。AlltoAllV的功能为通信域内的卡互相发送和接收数据并且定制每张卡给其它卡发送的数据量和从其它卡接收的数据量以及定制发送和接收的数据在内存中的偏移。结合原型中的参数描述接口功能具体为第i张卡发送sendBuf内存中第j块数据到第j张卡该块数据在sendBuf中偏移为sdispls[j]的位置且数据量为sendCounts[j]第j张卡将该数据存放到本卡recvBuf中偏移为rdispls[i]的位置接收数据量为recvCounts[i]这里的sendCounts[j]与recvCounts[i]需要相等。注意这里的偏移和数据量均为数据的个数单位为sizeof(sendType)。函数原型template bool commit false __aicore__ inline HcclHandle AlltoAllV(GM_ADDR sendBuf, void *sendCounts, void *sdispls, HcclDataType sendType, GM_ADDR recvBuf, void *recvCounts, void *rdispls, HcclDataType recvType, uint8_t repeat 1)参数说明表 1模板参数说明参数名输入/输出描述commit输入bool类型。参数取值如下true在调用Prepare接口时Commit同步通知服务端可以执行该通信任务。false在调用Prepare接口时不通知服务端执行该通信任务。表 2接口参数说明参数名输入/输出描述sendBuf输入源数据buffer地址。sendCounts输入本卡向通信域内其它每张卡发送的数据量sendCounts[i]表示本卡发送到第i张卡rank_i的数据量单位是sizeof(sendType)。sendCounts是一个uint64_t类型的数组数组长度必须为通信域内总rank数。例如sendBuf内数据的数据类型为fp16sendCounts[0]1sendCounts[1]2表示本rank给rank0发送1个fp16数据给rank1发送2个fp16数据。sdispls输入本卡向其它卡发送的数据在sendBuf中的偏移sdispls[i]n表示本卡发送给rank_i的数据块在sendBuf中的偏移数据量为n。sdispls是一个uint64_t类型的数组数组长度必须为通信域内总rank数。sendType输入sendBuf内数据的数据类型目前支持HcclDataType包含的全部数据类型HcclDataType详细可参考表1。recvBuf输出目的数据buffer地址集合通信结果输出到此buffer中。recvCounts输入本卡从其它卡接收的数据量recvCounts[i]表示本rank接收到的来自rank_i的数据个数单位是sizeof(recvType)。recvCounts是一个uint64_t类型的数组数组长度必须为通信域内总rank数。例如recvBuf内数据的数据类型为fp16recvCounts[0]1recvCounts[1]2表示本rank接收到rank0的1个fp16数据接收到rank1的2个fp16数据。rdispls输入本卡接收的数据存放在recvBuf中的偏移rdispls[i]n表示本卡接收到的rank_i的数据块在recvBuf中的偏移数据量为n。rdispls是一个uint64_t类型的数组数组长度必须为通信域内总rank数。recvType输入recvBuf内数据的数据类型目前支持HcclDataType包含的全部数据类型HcclDataType详细可参考表1。注意recvType和sendType必须一致。repeat输入一次下发的AlltoAllV通信任务个数。repeat取值≥1默认值为1。当repeat1时每个AlltoAllV任务的sendBuf\sendCounts\recvBuf\recvCounts参数不变sdispls和rdispls由服务端更新每一轮任务i的更新公式如下sdispls[i] sdispls[i] sendCounts[i], i∈[0, sdispls.size())rdispls[i] rdispls[i] recvCounts[i], i∈[0, rdispls.size())注意当设置repeat1时须根据此计算公式规划通信内存。返回值说明返回该任务的标识handleIdhandleId大于等于0。调用失败时返回 -1。约束说明调用本接口前确保已调用过InitV2和SetCcTilingV2接口。若HCCL对象的config模板参数未指定下发通信任务的核该接口只能在AIC核或者AIV核两者之一上调用。若HCCL对象的config模板参数中指定了下发通信任务的核则该接口可以在AIC核和AIV核上同时调用接口内部会根据指定的核的类型只在AIC核、AIV核二者之一下发该通信任务。一个通信域内所有Prepare接口和InterHcclGroupSync接口的总调用次数不能超过63。每张卡发送给卡rank_j的数据量sendCounts[j]与rank_j接收对应卡rank_i的数据量recvCounts[i]必须相等。对于Atlas A3 训练系列产品/Atlas A3 推理系列产品一个通信域内最大支持128卡通信。对于Ascend 950PR/Ascend 950DT一个通信域内所有Prepare接口的总调用次数不能超过63。对于Ascend 950PR/Ascend 950DT通信服务端为CCU时单次最大通信数据量不能超过256M。调用示例使用AlltoAllV接口等效实现4卡间的AlltoAll通信4张卡调用AlltoAllV接口。非多轮切分场景下每张卡上的数据块和数据量一致如下图中每张卡的A\B\C\D数据块数据量均为dataCount。图 1非切分场景下4卡AlltoAllV图示extern C __global__ __aicore__ void alltoallv_custom(GM_ADDR xGM, GM_ADDR yGM, GM_ADDR workspaceGM, GM_ADDR tilingGM) { constexpr uint32_t rankNum 4U; constexpr uint32_t dataCount 10U; // 假设图中A\B\C\D数据块的个数均为10个 uint64_t sendCounts[rankNum] {0}; uint64_t sDisplacements[rankNum] {0}; uint64_t recvCounts[rankNum] {0}; uint64_t rDisplacements[rankNum] {0}; for (uint32_t i 0U; i rankNum; i) { sendCounts[i] dataCount; sDisplacements[i] i * dataCount; recvCounts[i] dataCount; rDisplacements[i] i * dataCount; } auto sendBuf xGM; // xGM为AlltoAllV的输入GM地址 auto recvBuf yGM; // yGM为AlltoAllV的输出GM地址 auto dtype HcclDataType::HCCL_DATA_TYPE_FP16; REGISTER_TILING_DEFAULT(AllToAllVCustomTilingData); //AllToAllVCustomTilingData为对应算子头文件定义的结构体 GET_TILING_DATA_WITH_STRUCT(AllToAllVCustomTilingData, 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(AllToAllVCustomTilingData, alltoallvCcTiling)); if (ret ! HCCL_SUCCESS) { return; } auto handleId1 hccl.AlltoAllVtrue(sendBuf, sendCounts, sDisplacements, dtype, recvBuf, recvCounts, rDisplacements, dtype); hccl.Wait(handleId1); AscendC::SyncAlltrue(); // 全AIV核同步防止0核执行过快提前调用hccl.Finalize()接口导致其他核Wait卡死 hccl.Finalize(); } }使用AlltoAllV接口实现4卡间不同数据量的数据收发如下图所示每个rank下的方格中数字表示发送或接收的数据个数以rank1为例进行说明rank1分别向rank0、rank1、rank2、rank3发送2、2、3、2个数据并分别从rank0、rank1、rank2、rank3接收3、2、4、3个数据对应的代码示例如下。图 2非切分场景下4卡不均匀收发extern C __global__ __aicore__ void alltoallv_custom(GM_ADDR xGM, GM_ADDR yGM, GM_ADDR workspaceGM, GM_ADDR tilingGM) { constexpr uint32_t rankNum 4U; uint64_t sendCounts[rankNum] {0}; uint64_t sDisplacements[rankNum] {0}; uint64_t recvCounts[rankNum] {0}; uint64_t rDisplacements[rankNum] {0}; auto sendBuf xGM; // xGM为AlltoAllV的输入GM地址 auto recvBuf yGM; // yGM为AlltoAllV的输出GM地址 auto dtype HcclDataType::HCCL_DATA_TYPE_FP16; REGISTER_TILING_DEFAULT(AllToAllVCustomTilingData); //AllToAllVCustomTilingData为对应算子头文件定义的结构体 GET_TILING_DATA_WITH_STRUCT(AllToAllVCustomTilingData, 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(AllToAllVCustomTilingData, alltoallvCcTiling)); if (ret ! HCCL_SUCCESS) { return; } if(hccl.GetRankId() 0) { sendCounts[0] 3; sendCounts[1] 3; sendCounts[2] 3; sendCounts[3] 3; sDisplacements[1] 3; sDisplacements[2] 6; sDisplacements[2] 9; recvCounts[0] 3; recvCounts[1] 2; recvCounts[2] 1; recvCounts[3] 3; rDisplacements[1] 3; rDisplacements[2] 5; rDisplacements[3] 6; } else if(hccl.GetRankId() 1) { sendCounts[0] 2; sendCounts[1] 2; sendCounts[2] 3; sendCounts[3] 2; sDisplacements[1] 2; sDisplacements[2] 4; sDisplacements[2] 7; recvCounts[0] 3; recvCounts[1] 2; recvCounts[2] 4; recvCounts[3] 3; rDisplacements[1] 3; rDisplacements[2] 5; rDisplacements[3] 9; } else if(hccl.GetRankId() 2) { sendCounts[0] 1; sendCounts[1] 4; sendCounts[2] 4; sendCounts[3] 4; sDisplacements[1] 1; sDisplacements[2] 5; sDisplacements[2] 9; recvCounts[0] 3; recvCounts[1] 3; recvCounts[2] 4; recvCounts[3] 3; rDisplacements[1] 3; rDisplacements[2] 6; rDisplacements[3] 10; } else if(hccl.GetRankId() 3) { sendCounts[0] 3; sendCounts[1] 3; sendCounts[2] 3; sendCounts[3] 3; sDisplacements[1] 3; sDisplacements[2] 6; sDisplacements[2] 9; recvCounts[0] 3; recvCounts[1] 2; recvCounts[2] 4; recvCounts[3] 3; rDisplacements[1] 3; rDisplacements[2] 5; rDisplacements[3] 9; } auto handleId hccl.AlltoAllVtrue(sendBuf, sendCounts, sDisplacements, dtype, recvBuf, recvCounts, rDisplacements, dtype); hccl.Wait(handleId); AscendC::SyncAlltrue(); // 全AIV核同步防止0核执行过快提前调用hccl.Finalize()接口导致其他核Wait卡死 hccl.Finalize(); } }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考