AlltoAllvWrite【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品xAtlas A2 训练系列产品 / Atlas A2 推理系列产品x功能说明集合通信AlltoAllvWrite的任务下发接口返回该任务的标识handleId给用户。AlltoAllvWrite的功能为通信域内的卡互相发送和接收数据并且定制每张卡给其它卡发送的数据量和从其它卡接收的数据量以及定制发送和接收的数据在内存中的偏移。结合原型中的参数描述接口功能具体为本卡发送地址偏移为sendOffsets[i]字节且大小为sendSizes[i]字节的数据给第i张卡remoteWinOffset表示对端卡发送数据的地址偏移localDataSize表示发送给本卡的数据大小。注意这里的偏移和数据量均为字节数。函数原型template bool commit false __aicore__ inline HcclHandle AlltoAllvWrite(GM_ADDR usrIn, GM_ADDR sendOffsets, GM_ADDR sendSizes, uint64_t remoteWinOffset, uint64_t localDataSize)参数说明表 1模板参数说明参数名输入/输出描述commit输入bool类型。参数取值如下true在调用Prepare接口时Commit同步通知服务端可以执行该通信任务。false在调用Prepare接口时不通知服务端执行该通信任务。表 2接口参数说明参数名输入/输出描述usrIn输入源数据buffer地址。sendOffsets输入待发送的每个分片的数据大小以字节为单位。sendSizes输入待发送的每个分片的偏移以字节为单位。remoteWinOffset输入对端卡发送的数据偏移以字节为单位。localDataSize输入发送给本卡的数据大小以字节为单位。返回值说明返回该任务的标识handleIdhandleId大于等于0。调用失败时返回 -1。约束说明调用本接口前确保已调用过InitV2和SetCcTilingV2接口。若HCCL对象的模板参数config未指定下发通信任务的核则该接口只能在AIC核或者AIV核两者之一上调用。若HCCL对象的模板参数config指定了下发通信任务的核则该接口可以在AIC核和AIV核上同时调用接口内部根据指定的核的类型在对应的AIC核、AIV核二者之一下发该通信任务。一个通信域内所有Prepare接口和InterHcclGroupSync接口的总调用次数不能超过63。对于Ascend 950PR/Ascend 950DT通信服务端为CCU时单次最大通信数据量不能超过256M。调用示例extern C __global__ __aicore__ void alltoallvwrite_custom(GM_ADDR xGM, GM_ADDR yGM, GM_ADDR workspaceGM, GM_ADDR tilingGM) { REGISTER_TILING_DEFAULT(AllToAllVWriteCustomTilingData); //AllToAllVWriteCustomTilingData为对应算子头文件定义的结构体 GET_TILING_DATA_WITH_STRUCT(AllToAllVWriteCustomTilingData, tilingData, tilingGM); auto cfg tilingData.param; uint32_t M cfg.M; uint32_t K cfg.K; uint32_t dataType cfg.dataType; uint32_t dataTypeSize cfg.dataTypeSize; KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2); HcclHcclServerType::HCCL_SERVER_TYPE_CCU hccl; GM_ADDR context GetHcclContextHCCL_GROUP_ID_0(); hccl.InitV2(context, tilingData); hccl.SetCcTilingV2(offsetof(AllToAllVCustomV3TilingData, mc2CcTiling)); uint32_t rankDim hccl.GetRankDim(); uint32_t rankId hccl.GetRankId(); uint64_t perRankDataSize_ M * K * dataTypeSize / rankDim; GM_ADDR sendSizeGM_ workspaceGM; GM_ADDR sendOffsetGM_ sendSizeGM_ rankDim * sizeof(uint64_t) * 2; __gm__ uint64_t *sendSizes reinterpret_cast__gm__ uint64_t *(sendSizeGM_); __gm__ uint64_t *sendOffsets reinterpret_cast__gm__ uint64_t *(sendOffsetGM_); for (uint32_t i 0U; i rankDim; i) { // 当前ccu通信都是双die所以sendSize和sendOffset需要等分切成die0和die1的数据 sendSizes[i] perRankDataSize_ / 2; sendSizes[i rankDim] perRankDataSize_ - perRankDataSize_ / 2; sendOffsets[i] i * perRankDataSize_; sendOffsets[i rankDim] i * perRankDataSize_ sendSizes[i]; } uint64_t remoteWinOffset rankId * perRankDataSize_; uint64_t localDataSize perRankDataSize_; if (TILING_KEY_IS(1000UL)) { if ASCEND_IS_AIV { AscendC::HcclHandle handleId -1; handleId hccl.AlltoAllvWritetrue(xGM, sendOffsetGM_, sendSizeGM_, remoteWinOffset, localDataSize); 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），仅供参考