人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载本指南围绕 asc-devkit 开源仓库中 HCCL 通信类下的 HCCL Context 接口展开系统讲解SetHcclContext/GetHcclContext两个核函数Kernel侧接口的功能定位、函数原型、参数约束、调用方式与底层实现原理并给出其在通算融合算子MC2通信流程中的完整实战用法。读完本文你将掌握如何为算子核函数配置通信域 context消息区地址、如何在 host 侧与核函数侧协同完成通信域初始化以及如何借助 kernel_common.h 与 hccl_control.h 中的实现细节理解 context 的内部结构。HCCL Context 在通算融合通信框架中的位置在 asc-devkit 的 HCCL 通信类体系参见 HCCL 通信类总览中HCCL 接口被划分为四个层次HCCL 核函数侧接口、HCCL Tiling 侧接口、HCCL MC2 接口以及本主题的 HCCL Context。其中HCCL MC2 接口用于在 host 侧定义通算融合算子注册通信算子属性HCCL Tiling 侧接口用于在 Tiling 侧设置通信算法参数与 TilingDataHCCL 核函数侧接口Hccl客户端类及其InitV2、Prepare、Commit、Wait、Finalize等接口负责在 AI Core 上编排通信任务HCCL Context为上述核函数侧接口提供每个通信域对应的 context消息区地址是连接 host 侧通信域配置与核函数侧通信任务执行的关键桥梁。根据 HCCL Context 简介 的说明HCCL Context 系列接口用于在核函数侧设置/获取通算融合算子每个通信域对应的 context消息区地址。这里消息区承载了通信上下文的核心信息包括 rankId、rankNum通信域内卡数、workspace 地址、通信窗口windows等。需要注意的是host 侧必须先调用 HcclGroup 接口配置通信域名称核函数侧调用GetHcclContext才能取到对应的 context 地址。[!NOTE] 说明 本系列接口为试验接口在后续版本中可能会调整或改进不保证后续兼容性。开发者在使用过程中需关注 asc-devkit 后续版本更新见 HCCL_Context_intro.md。产品支持情况SetHcclContext与GetHcclContext两个接口的产品支持矩阵一致汇总如下产品系列支持情况Atlas A3 系列产品支持Atlas A2 系列产品支持Atlas 推理系列产品 AI Core支持Atlas 推理系列产品 Vector Core支持Ascend 950PR950DT 系列产品不支持Atlas 200I/500 A2 推理产品不支持Atlas 训练系列产品不支持与 HCCL 模板参数 中HcclServerType的支持范围相对照可以发现Ascend 950PR950DT 系列仅支持 CCU 服务端Atlas A3/A2 系列仅支持 AI CPU 服务端而 HCCL Context 接口恰恰在这些服务端场景之外的产品上不可用说明该接口与特定服务端实现强相关开发前务必确认目标产品在支持矩阵内。SetHcclContext设置通信域 context 地址功能说明SetHcclContext用于设置通算融合算子每个通信域对应的 context消息区地址。在核函数执行流上该接口通常由框架或用户在通信任务下发前调用将 host 侧准备好的 context 写入核内存储。函数原型template uint32_t index __aicore__ inline void SetHcclContext(__gm__ uint8_t* context)该接口是一个模板函数index为编译期常量最终展开为对块内保留数组的一次赋值操作详见下文源码实现剖析。参数说明参数名描述index模板参数表示要设置的通信域 ID。当前只支持 2 个通信域index 只能为 0/1。context对应通信域的 context消息区地址类型为__gm__ uint8_t*即 Global Memory 中的字节地址。返回值说明无。约束说明当前最多只支持 2 个通信域index超过 1 时接口不会写入任何数据。调用示例// 给GROUP_0设置消息区地址 AscendC::SetHcclContext0(contextGM);源码实现剖析SetHcclContext的实现在 include/basic_api/kernel_common.h#L194-L201template uint32_t index __aicore__ inline void SetHcclContext(__gm__ uint8_t* context) { if constexpr (index 1) { return; } g_hcclContextReserved[index] context; }实现要点使用if constexpr (index 1)在编译期过滤非法 indexindex为编译期常量因此非法分支不会生成任何运行时指令也解释了约束说明中最多 2 个通信域的语义来源将 context 地址写入块内保留数组g_hcclContextReserved其定义为见 kernel_common.h#L87__BLOCK_LOCAL__ __inline__ __gm__ uint8_t* g_hcclContextReserved[2];数组长度为 2与最多 2 个通信域的约束一一对应__BLOCK_LOCAL__表明该数组是块Block级局部变量由该核上的所有执行单元共享。GetHcclContext获取通信域 context 地址功能说明GetHcclContext用于获取指定 index 通信域的 context消息区地址是核函数侧读取通信上下文的标准入口。在实际算子中获取到的地址会被直接传给Hccl.InitV2完成客户端初始化详见 InitV2。函数原型template uint32_t index __aicore__ inline __gm__ uint8_t* __gm__ GetHcclContext(void)模板参数说明参数名描述index模板参数表示要获取的通信域 ID。当前只支持 2 个通信域index 只能为 0/1。返回值说明返回指定通信域的 context消息区地址若index 1返回nullptr。约束说明当前最多只支持 2 个通信域。调用示例GM_ADDR contextGM AscendC::GetHcclContext0();源码实现剖析GetHcclContext的实现在 include/basic_api/kernel_common.h#L203-L210template uint32_t index __aicore__ inline __gm__ uint8_t* __gm__ GetHcclContext(void) { if constexpr (index 1) { return nullptr; } return g_hcclContextReserved[index]; }与SetHcclContext对称合法 index 直接返回g_hcclContextReserved[index]非法 index 在编译期即被替换为nullptr返回避免越界访问。在 hccl_control.h 等下游模块中拿到该地址后还会先做空指针检查HCCL_CHECK_STATUS(context nullptr, return)因此即使未设置 context 也不会导致崩溃而是静默跳过相关逻辑。底层原理context 内部结构与消息区消费context 指向的数据结构HcclCombineOpParamcontext 地址在底层被解释为HcclCombineOpParam结构体其定义见 impl/adv_api/detail/hccl/common/hccl_inner_def.h#L113-L136struct HcclCombineOpParam { uint64_t workSpace; // Address for communication between client and server, // hccl requests and clears uint64_t workSpaceSize; // Space for communication between client and server uint32_t rankId; // id of this rank uint32_t rankNum; // num of ranks in this comm group uint64_t winSize; // size of each windows memory #if defined(__NPU_ARCH__) (__NPU_ARCH__ 3510 || __NPU_ARCH__ 9201 || __NPU_ARCH__ 9202) uint64_t windowsIn[HCCL_MAX_RANK_NUM_V310]; uint64_t windowsOut[HCCL_MAX_RANK_NUM_V310]; GM_ADDR xnOffset; GM_ADDR ckeOffset; uint8_t res[16]; uint32_t opType[MAX_CC_TILING_NUM]; uint8_t algorithmType[MAX_CC_TILING_NUM]; uint8_t res2[7744]; #else uint64_t windowsIn[HCCL_MAX_RANK_NUM]; uint64_t windowsOut[HCCL_MAX_RANK_NUM]; uint8_t res[8328]; #endif uint8_t multiFlag; __gm__ IbVerbsData* data; };关键字段含义workSpace/workSpaceSize客户端与服务端AI CPU 或 CCU之间通信所用的消息区地址及大小rankId/rankNum当前 rank 的编号与通信域内 rank 总数winSize、windowsIn/windowsOut通信窗口内存的大小与收发窗口地址数组multiFlag多通信域标志data指向IbVerbsData的全局地址用于底层网络数据传输。这也印证了 InitV2 中context 包含 rankDim、rankID 等相关信息的说明。下游如何消费 context以 hccl_control.h 为例impl/adv_api/detail/hccl/common/hccl_control.h#L39-L63 展示了 context 在核函数侧被消费的典型路径。例如GetRestartFromContext__aicore__ inline void GetRestartFromContext(GM_ADDR context, uint8_t restart) { HCCL_CHECK_STATUS(context nullptr, return); __gm__ HcclCombineOpParam* hcclContext (__gm__ HcclCombineOpParam*)context; uint64_t msgAddr hcclContext-workSpace; HCCL_CHECK_STATUS(msgAddr 0, return); __gm__ HcclMsgArea* hcclMsgArea (__gm__ HcclMsgArea*)msgAddr; __gm__ ControlHcclMsg* controlMsgGM hcclMsgArea-controlMsg; FlushDataCache(controlMsgGM); restart controlMsgGM-restart; }完整调用链为GetHcclContextindex()取出 context → 强转为HcclCombineOpParam*→ 读取workSpace得到消息区基址 → 转为HcclMsgArea*→ 访问controlMsg中的restart等控制字段并通过FlushDataCache保证数据一致性。ResetRestartFlag则反向执行将restart清零、置位restarting与resetSeq后写回。双通信域支持hccl_control.h#L65-L89 中的GetRestart/SetRestart明确体现了最多 2 个通信域的设计// Maximum support for dual communication domains if (ctxNum SINGLE_COMM_NUM) { GetRestartFromContext(AscendC::GetHcclContext0(), restart); } if (ctxNum MULTI_COMM_NUM) { GetRestartFromContext(AscendC::GetHcclContext1(), restart); }其中SINGLE_COMM_NUM 1、MULTI_COMM_NUM 2见 hccl_control.h#L22-L23。当算子涉及多个通信域时可分别通过GetHcclContext0()与GetHcclContext1()访问各自的 context从而支持单域与双域两种模式。实战在 HCCL 通信流程中使用 HCCL Context第一步host 侧注册算子并配置通信域名称使用 HCCL Context 的前提是 host 侧已完成通信域配置。根据 HcclGroup 的说明算子需先通过 MC2 接口注册为通算融合算子并在算子属性中配置通信域名称示例class MC2Custom : public OpDef { public: MC2Custom(const char* name) : OpDef(name) { this-Input(x).ParamType(REQUIRED).DataType({ge::DT_FLOAT}).Format({ge::FORMAT_ND}); this-Input(y).ParamType(REQUIRED).DataType({ge::DT_FLOAT}).Format({ge::FORMAT_ND}); this-Output(z).ParamType(REQUIRED).DataType({ge::DT_FLOAT}).Format({ge::FORMAT_ND}); this-Attr(group).AttrType(REQUIRED).String(); this-AICore().AddConfig(ascendxxx); this-MC2().HcclGroup(group); // 配置通信域名称为group } }; OP_ADD(MC2Custom);约束要点源自 HcclGroup使用HcclGroup前算子必须已通过 MC2 接口注册为通算融合算子并获得OpMC2Def结构通信域名称必须预先配置为 REQUIRED String 类型的属性属性名即为通信域名称单个通信域使用const char *多通信域使用std::vectorconst char *。第二步核函数侧获取 context 并完成 HCCL 客户端初始化在核函数中典型流程是先通过GetHcclContext0()获取 context再传给Hccl::InitV2完整流程参考 HCCL 使用说明 与 InitV2extern C __global__ __aicore__ void userKernel(GM_ADDR aGM, GM_ADDR workspaceGM, GM_ADDR tilingGM) { REGISTER_TILING_DEFAULT(UserCustomTilingData); GET_TILING_DATA_WITH_STRUCT(UserCustomTilingData, tilingData, tilingGM); GM_ADDR contextGM AscendC::GetHcclContext0(); Hccl hccl; hccl.InitV2(contextGM, tilingData); hccl.SetCcTilingV2(offsetof(UserCustomTilingData, tiling)); // 调用HCCL的Prepare、Commit、Wait、Finalize接口 }其中tilingGM为 host 侧传入的算子 TilingData 的 GM 地址通过 GET_TILING_DATA_WITH_STRUCT 获取initTiling参数不能使用 Global Memory 地址建议使用栈地址。第三步完整通信编排示例以下示例来自 HCCL 使用说明展示了GetHcclContext0()在 ReduceScatter 任务编排中的完整用法extern C __global__ __aicore__ void reduce_scatter_custom( GM_ADDR xGM, GM_ADDR yGM, GM_ADDR workspaceGM, GM_ADDR tilingGM) { auto sendBuf xGM; // xGM为ReduceScatter的输入GM地址 auto recvBuf yGM; // yGM为ReduceScatter的输出GM地址 constexpr size_t rankSize 4U; // 4张卡 constexpr size_t tileCnt 3U; // 卡上的数据均匀分成rankSize份且每份又被切分成3份 constexpr size_t tileLen 100U; // 被切分后的每份数据个数 uint64_t strideCount tileLen * tileCnt; // 表示sendBuf上相邻数据块间的起始地址的偏移量 REGISTER_TILING_DEFAULT(ReduceScatterCustomTilingData); GET_TILING_DATA_WITH_STRUCT(ReduceScatterCustomTilingData, tilingData, tilingGM); Hccl hccl; GM_ADDR contextGM AscendC::GetHcclContext0(); // 通过此方式获取HCCL context if (AscendC::g_coreType AIV) { // 指定AIV核通信 hccl.InitV2(contextGM, tilingData); auto ret hccl.SetCcTilingV2(offsetof(ReduceScatterCustomTilingData, reduceScatterCcTiling)); if (ret ! HCCL_SUCCESS) { return; } auto handleId hccl.ReduceScatter( sendBuf, recvBuf, tileLen, HcclDataType::HCCL_DATA_TYPE_FP32, HcclReduceOp::HCCL_REDUCE_SUM, strideCount, tileCnt); for (int i 0; i tileCnt; i) { hccl.Commit(handleId); auto ret hccl.Wait(handleId); // 执行其他计算逻辑 .... } AscendC::SyncAlltrue(); // 全AIV核同步防止0核执行过快提前调用Finalize导致其他核Wait卡死 hccl.Finalize(); } }流程要点获取 contextAscendC::GetHcclContext0()取得通信域 0 的 context 地址初始化客户端hccl.InitV2(contextGM, tilingData)设置 TilingSetCcTilingV2与InitV2必须配对使用下发通信任务Prepare 接口如ReduceScatter返回handleIdrepeat参数表示该任务的通信次数必须与针对该handleId的Commit、Wait调用次数一致通知并等待Commit(handleId)通知 AI CPU/CCU 服务端执行任务Wait(handleId)阻塞等待执行结束退出Finalize()通知服务端后续无通信任务客户端检测并等待最后一个通信任务结束。若模板参数未指定下发核Prepare 及步骤 25 的接口仅能在 AIC 或 AIV 之一上运行需通过内置常量g_coreType判断并限定执行核如示例中的if (AscendC::g_coreType AIV)也可参照 HCCL 模板参数 通过HcclServerConfig直接指定核如{CoreType::ON_AIV, 10}。调试含有 HCCL 高阶 API 的算子时可在编译工程中增加编译选项-DASCENDC_DEBUG开启异常场景拦截能力。工程验证与测试asc-devkit 的测试工程中保留了 HCCL Context 接口的头文件可用性验证用例见 tests/api/basic_api/ascendc_header_checker/kernel_common.cpp#L42-L47// __aicore__ inline void SetHcclContext(__gm__ uint8_t* context) extern C __global__ __aicore__ void KernelTestSetHcclContext1() { __gm__ uint8_t* context; AscendC::SetHcclContext1(context); }该用例以 kernel 入口形式调用AscendC::SetHcclContext1(context)验证了模板参数index 1第二个通信域场景下接口可正确编译链接。开发者在自定义算子中引入 HCCL Context 前可参考此类用例确认接口在当前头文件与编译链路下的可用性。小结SetHcclContext/GetHcclContext是 asc-devkit 中通算融合算子核函数侧管理通信域 context 的唯一标准入口分别负责写入与读取消息区地址两个接口均为模板接口index仅支持 0/1底层由块内保留数组g_hcclContextReserved[2]承载非法 index 在编译期即被拦截context 实际指向HcclCombineOpParam结构包含workSpace、rankId、rankNum、winSize、windowsIn/Out等字段被 hccl_control.h 等下游模块强转消费host 侧须先经 HcclGroup 配置通信域名称核函数侧才能通过GetHcclContext取得有效 context使用前务必核对目标产品在支持矩阵中并留意该系列接口的试验属性与后续版本变更风险。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐昇腾 CANN asc-devkit 多核 Matmul Tiling 核数设置SetDim 接口详解昇腾 CANN asc devkit 多核 Matmul Tiling 核数设置SetDim 接口详解 本文面向使用 CANN 昇腾 AI 处理器算子开发语言人工智能深度学习算子库CANNAscendasc-devkit 中 GetTilingFunc 接口详解按算子类型获取 Tiling 函数并独立调测asc devkit 中 GetTilingFunc 接口详解按算子类型获取 Tiling 函数并独立调测 本篇以 asc devkit 的 Tiling 调人工智能深度学习算子库CANNAscendasc-devkit 通信算子 Tiling 配置SetGroupName 接口详解与源码剖析asc devkit 通信算子 Tiling 配置SetGroupName 接口详解与源码剖析 导读 SetGroupName 是 asc devkitCA人工智能深度学习算子库CANNAscend上一篇QQ聊天记录解密终极指南跨平台数据库密钥提取完整方案下一篇如何免费获得专业级字体Montserrat开源字体完整使用指南创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
