news 2026/7/17 15:19:51

cann/asc-devkit HCCL AlltoAll接口

作者头像

张小明

前端开发工程师

1.2k 24
文章封面图
cann/asc-devkit HCCL AlltoAll接口

AlltoAll

【免费下载链接】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 训练系列产品：不支持

功能说明

集合通信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的数据，那么dataCount=1。
dataType输入AlltoAll操作的数据类型，目前支持HcclDataType包含的全部数据类型，HcclDataType详细可参考表1。
strideCount输入多轮切分场景下，一次AlltoAll任务中，每张卡内参与通信的数据块间的间隔。默认值为0，表示数据块内存连续。
strideCount=0，每张卡内参与通信的数据块内存连续。卡rank_j收到来自卡rank_i的sendBuf中第j块数据，且数据块间的偏移数据量为j*dataCount，并将该数据存放于本卡recvBuf中第i块的位置，且偏移数据量为i*dataCount。
strideCount>0，每张卡内参与通信的相邻数据块的起始地址偏移数据量为strideCount。卡rank_j收到来自卡rank_i的sendBuf中第j块数据，且数据块间的偏移数据量为j*strideCount，并将该数据存放于本卡recvBuf中第i块的位置，且偏移数据量为i*strideCount。

注意：上述的偏移数据量为数据个数，单位为sizeof(dataType)。
repeat输入一次下发的AlltoAll通信任务个数。repeat取值≥1，默认值为1。当repeat>1时，每轮AlltoAll任务的sendBuf和recvBuf地址由服务端更新，每一轮任务i的更新公式如下：

sendBuf[i] = sendBuf + dataCount * sizeof(datatype) * i, i∈[0, repeat)

recvBuf[i] = recvBuf + dataCount * sizeof(datatype) * i, i∈[0, repeat)

注意：当设置repeat>1时，须与strideCount参数配合使用，规划通信数据地址。

返回值说明

返回该任务的标识handleId，handleId大于等于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通信
    ![](https://raw.gitcode.com/cann/asc-devkit/raw/1cd1d1192e01ae4a8684cd07552c35e13e5c4bdf/docs/zh/api/figures/非多轮切分场景下4卡AlltoAll通信.png "非多轮切分场景下4卡AlltoAll通信"?utm_source=gitcode_repo_files)

    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::GetHcclContext<0>(); // 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.AlltoAll<true>(sendBuf, recvBuf, dataCount, HcclDataType::HCCL_DATA_TYPE_FP16); hccl.Wait(handleId); AscendC::SyncAll<true>(); // AIV核全同步，防止0核执行过快，提前调用hccl.Finalize()接口，导致其他核Wait卡死 hccl.Finalize(); } }

  • 多轮切分场景

    开启多轮切分，等效处理上述非多轮切分示例的通信。在每张卡的数据均分成4块（A\B\C\D）的基础上，将每一块继续切分若干块。本例中继续切分3块，如下图所示，被继续切分成的3块数据包括，2个数据量为tileLen的数据块，1个数据量为tailLen的尾块。切分后，需要分3轮进行AlltoAll通信任务，将等效上述非多轮切分的通信结果。

    图23轮切分场景下4卡AlltoAll通信
    ![](https://raw.gitcode.com/cann/asc-devkit/raw/1cd1d1192e01ae4a8684cd07552c35e13e5c4bdf/docs/zh/api/figures/3轮切分场景下4卡AlltoAll通信.png "3轮切分场景下4卡AlltoAll通信"?utm_source=gitcode_repo_files)

    具体实现为，第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示意图
    ![](https://raw.gitcode.com/cann/asc-devkit/raw/1cd1d1192e01ae4a8684cd07552c35e13e5c4bdf/docs/zh/api/figures/第一轮4卡AlltoAll示意图.png "第一轮4卡AlltoAll示意图"?utm_source=gitcode_repo_files)

    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::GetHcclContext<0>(); // 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.AlltoAll<true>(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.AlltoAll<true>(sendBuf, recvBuf, tailLen, HcclDataType::HCCL_DATA_TYPE_FP16, strideCount, tailNum); for (uint8_t i = 0; i < tileNum; i++) { hccl.Wait(handleId1); } hccl.Wait(handleId2); AscendC::SyncAll<true>(); // 全AIV核同步，防止0核执行过快，提前调用hccl.Finalize()接口，导致其他核Wait卡死 hccl.Finalize(); } }

【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit

创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考

版权声明: 本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若内容造成侵权/违法违规/事实不符，请联系邮箱：809451989@qq.com进行投诉反馈，一经查实，立即删除！
网站建设 2026/7/17 15:19:26

自己写的小说也被判AI？把误判的AI率降到检测合格

自己写的小说也被判AI&#xff1f;把误判的AI率降到检测合格 你大概是被气到了才点进来的。一部小说从人物到情节全是你自己脑子里长出来的&#xff0c;一个字都没让AI代笔&#xff0c;投到平台却弹出一句"疑似AI生成&#xff0c;AI率偏高"。你甚至想截图去跟编辑理…

作者头像 李华
网站建设 2026/7/17 15:18:13

flow.ci进阶技巧：条件执行、定时任务与Web终端调试实用教程

flow.ci进阶技巧&#xff1a;条件执行、定时任务与Web终端调试实用教程 【免费下载链接】docs flow.ci user guide 项目地址: https://gitcode.com/gh_mirrors/docs17/docs flow.ci是一款功能强大的持续集成工具&#xff0c;本文将为你介绍三个提升工作效率的进阶技巧&a…

作者头像 李华
网站建设 2026/7/17 15:17:10

5分钟快速上手Dialogic 2：打造专业游戏对话系统的终极指南

5分钟快速上手Dialogic 2&#xff1a;打造专业游戏对话系统的终极指南 【免费下载链接】dialogic &#x1f4ac; Create Dialogs, Visual Novels, RPGs, and manage Characters with Godot to create your Game! 项目地址: https://gitcode.com/gh_mirrors/dia/dialogic …

作者头像 李华
网站建设 2026/7/17 15:16:46

如何让AMD显卡原生运行CUDA程序：ZLUDA兼容层完整指南

如何让AMD显卡原生运行CUDA程序&#xff1a;ZLUDA兼容层完整指南 【免费下载链接】ZLUDA CUDA on non-NVIDIA GPUs 项目地址: https://gitcode.com/GitHub_Trending/zl/ZLUDA 想让你的AMD显卡也能运行CUDA程序吗&#xff1f;ZLUDA是一个革命性的开源项目&#xff0c;它为…

作者头像 李华