AlltoAllV
产品支持情况
功能说明
集合通信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模板参数说明
|
表 2接口参数说明
返回值说明
返回该任务的标识handleId，handleId大于等于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图示
![](https://raw.gitcode.com/cann/asc-devkit/raw/bb4de8edfa8363a842738c29971dd5671b794a6f/docs/api/figures/非切分场景下4卡AlltoAllV图示.png "非切分场景下4卡AlltoAllV图示"?utm_source=gitcode_repo_files)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::GetHcclContext<0>(); // 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.AlltoAllV<true>(sendBuf, sendCounts, sDisplacements, dtype, recvBuf, recvCounts, rDisplacements, dtype); hccl.Wait(handleId1); AscendC::SyncAll<true>(); // 全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卡不均匀收发
![](https://raw.gitcode.com/cann/asc-devkit/raw/bb4de8edfa8363a842738c29971dd5671b794a6f/docs/api/figures/非切分场景下4卡不均匀收发.png "非切分场景下4卡不均匀收发"?utm_source=gitcode_repo_files)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::GetHcclContext<0>(); // 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.AlltoAllV<true>(sendBuf, sendCounts, sDisplacements, dtype, recvBuf, recvCounts, rDisplacements, dtype); hccl.Wait(handleId); AscendC::SyncAll<true>(); // 全AIV核同步，防止0核执行过快，提前调用hccl.Finalize()接口，导致其他核Wait卡死 hccl.Finalize(); } }
创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
