CANNAscend人工智能任务调度【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址https://gitcode.com/cann/runtime点击查看免费下载导读本文以 CANN/runtime 仓库中 4_launch_blocking 样例 为线索系统讲解 CANN Runtime 中控制 Kernel Launch 同步/异步行为的完整机制进程级默认策略ASCEND_RT_LAUNCH_BLOCKING环境变量、流级三态策略ACL_STREAM_LAUNCH_BLOCKING_MODE_*流属性以及临时非阻塞区间aclrtNonBlockingLaunchBegin/End嵌套。读完本文你将掌握这套控制体系的优先级规则、如何用可验证的样例代码观测策略是否生效以及在实际算子调试场景中避免死锁、按需切换同步/异步下发的实战方法。一、为什么需要 Kernel Launch Blocking 控制在 CANN Runtime 的默认模型下aclrtLaunchKernel将 Kernel 任务下发到 Stream 后立即返回任务的真正执行由 Device 侧异步完成。这种异步模型吞吐高但在算子调试、错误定位等场景中开发者往往希望任务真正执行完再返回从而基于可靠的计算结果进行排查。CANN Runtime 为此提供了一套分层、可嵌套、可覆盖的控制体系让同步/异步行为既可以由进程级环境变量统一管控也可以细粒度到单条 Stream甚至可以在同步模式下临时开一个异步窗口进程级默认策略环境变量ASCEND_RT_LAUNCH_BLOCKING流级策略aclrtSetStreamAttribute设置ACL_STREAM_LAUNCH_BLOCKING_MODE属性临时非阻塞区间aclrtNonBlockingLaunchBegin/aclrtNonBlockingLaunchEnd配对使用。三者之间的关系优先级从高到低是本文的核心也是实际工程中最容易踩坑的地方。二、控制体系总览与优先级规则根据 4_launch_blocking 样例文档 的Control Rules and Notes章节Kernel Launch 是否阻塞按下述优先级决定临时非阻塞区间优先流处于aclrtNonBlockingLaunchBegin和aclrtNonBlockingLaunchEnd标记的非阻塞区间时无论环境变量或流属性如何Kernel Launch 都保持异步。显式流属性次之流属性被设置为ACL_STREAM_LAUNCH_BLOCKING_MODE_NON_BLOCKING强制异步或ACL_STREAM_LAUNCH_BLOCKING_MODE_BLOCKING强制同步时覆盖环境变量。环境变量兜底流属性为ACL_STREAM_LAUNCH_BLOCKING_MODE_CTRL_BY_ENV跟随环境变量时行为完全由ASCEND_RT_LAUNCH_BLOCKING决定。用一张判定流程图可以直观表达流是否处于非阻塞区间 ├── 是 → 异步最高优先 └── 否 → 检查流属性 ├── NON_BLOCKING → 异步 ├── BLOCKING → 同步 └── CTRL_BY_ENV → 看 ASCEND_RT_LAUNCH_BLOCKING ├── 1 → 同步 └── 0 → 异步这一优先级在样例的main.cpp三个场景中被逐一验证下文会结合源码拆解。三、第一层进程级默认策略ASCEND_RT_LAUNCH_BLOCKING3.1 功能与取值环境变量ASCEND_RT_LAUNCH_BLOCKING用于控制 Kernel Launch 和模型执行任务采用同步模式或异步模式主要面向算子调试场景。仓库中的环境变量说明文档 docs/zh/env_vars/ASCEND_RT_LAUNCH_BLOCKING.md 给出了精确的取值语义取值行为0默认接口采用异步模式完成任务下发后返回不等待任务执行完成1以下接口采用同步模式任务执行完成后返回其他值与配置为0时相同开启同步模式的接口清单包括aclrtLaunchKernelaclrtLaunchKernelV2aclrtLaunchKernelWithConfigaclrtLaunchKernelWithHostArgsaclrtLaunchKernelWithArgsArrayaclrtLaunchSIMTKernelWithArgsArrayaclrtLaunchSIMTKernelWithHostArgsaclmdlRIExecuteAsync配置示例export ASCEND_RT_LAUNCH_BLOCKING13.2 关键使用约束易踩坑初始化时读取不支持动态修改Runtime 在初始化阶段读取该变量调用任意 Runtime 接口前需完成配置初始化后修改不生效必须重启进程。这正是样例run.sh要为每个场景单独启动进程的根本原因。性能影响开启后所有 Kernel Launch 都要等任务执行完才返回会明显影响业务性能官方建议仅在算子调试场景下使用。死锁风险开启后较早版本的算子可能未适配本功能。若算子中已下发的任务依赖尚未下发的任务同步等待会导致后续任务无法继续下发可能发生死锁。对应的解法正是本文后面要讲的临时非阻塞区间或流级NON_BLOCKING属性。特殊 Stream 不生效通过aclrtCreateStreamWithConfig且 flag 为ACL_STREAM_PERSISTENT、ACL_STREAM_CPU_SCHEDULE或ACL_STREAM_DEVICE_USE_ONLY创建的 Stream本功能不生效。四、第二层流级三态策略4.1 三个模式的定义流属性ACL_STREAM_LAUNCH_BLOCKING_MODE定义了三种取值定义位于 include/external/acl/acl_rt.h#define ACL_STREAM_LAUNCH_BLOCKING_MODE_CTRL_BY_ENV 0x00000000U // 跟随环境变量 #define ACL_STREAM_LAUNCH_BLOCKING_MODE_NON_BLOCKING 0x00000001U // 强制异步 #define ACL_STREAM_LAUNCH_BLOCKING_MODE_BLOCKING 0x00000002U // 强制同步设置/读取该属性的接口为aclError aclrtSetStreamAttribute(aclrtStream stream, aclrtStreamAttr stmAttrType, aclrtStreamAttrValue* value); aclError aclrtGetStreamAttribute(aclrtStream stream, aclrtStreamAttr stmAttrType, aclrtStreamAttrValue* value);其中stmAttrType取ACL_STREAM_LAUNCH_BLOCKING_MODE枚举值 6aclrtStreamAttrValue联合体中的launchBlockingMode字段携带模式值见 include/external/acl/acl_rt.h。4.2 样例中的读写验证样例 main.cpp 中SetAndCheckLaunchBlockingMode函数演示了写入-回读-比对的标准做法aclrtStreamAttrValue value{}; value.launchBlockingMode mode; // 传入 CTRL_BY_ENV / NON_BLOCKING / BLOCKING aclrtSetStreamAttribute(context.Stream(), ACL_STREAM_LAUNCH_BLOCKING_MODE, value); aclrtStreamAttrValue actual{}; aclrtGetStreamAttribute(context.Stream(), ACL_STREAM_LAUNCH_BLOCKING_MODE, actual); // 校验 actual.launchBlockingMode mode注意ACL_STREAM_LAUNCH_BLOCKING_MODE 6这一枚举序号位于 include/external/acl/acl_rt.h与ACL_STREAM_ATTR_PRIORITY等属性并列属于 Stream 属性体系的一部分见同文件 L3902、L3917 的说明。五、第三层临时非阻塞区间5.1 接口定义接口声明位于 include/external/acl/acl_rt.h// begin a non-blocking kernel launch section on the specified stream aclError aclrtNonBlockingLaunchBegin(aclrtStream stream, uint64_t flag); // end a non-blocking kernel launch section on the specified stream aclError aclrtNonBlockingLaunchEnd(aclrtStream stream, uint64_t flag);参数语义stream后续 Kernel Launch 使用的 Stream传nullptr表示当前 Context 的默认流。flag保留参数必须传0。5.2 嵌套语义核心非阻塞区间支持嵌套内层aclrtNonBlockingLaunchEnd只减少嵌套深度不做任何同步区间内保持异步最外层aclrtNonBlockingLaunchEnd真正退出非阻塞区间并在恢复后的策略要求同步执行时等待流完成。这意味着Begin 与 End 必须在同一条 Stream 上配对调用且退出时机的同步行为完全由退出后恢复的策略决定。5.3 底层实现线索从源码结构看该功能在 Runtime 内部有完整的调用链支撑。以rtNonBlockingLaunchBegin为例其分层如下对外 C 接口src/runtime/api/api_c_standard_soc.cc负责把aclrtStream转成内部Stream*后调用apiInstance-NonBlockingLaunchBegin统一抽象接口src/runtime/api/api.hpp 定义了virtual rtError_t NonBlockingLaunchBegin(Stream* const stream, const uint64_t flag) 0装饰器错误码/参数校验src/runtime/api/impl/api_decorator.cc、src/runtime/api/impl/api_error_standard_soc.cc具体实现不同形态的产品走不同实现如 src/runtime/api/impl/api_impl_standard_soc.cc 中StreamLaunchBlocking::NonBlockingLaunchBegin(targetStm)以及 api_impl_arch5162.cc、api_impl_tiny.cc 等。有兴趣的读者可以沿StreamLaunchBlocking继续追踪嵌套深度的计数与最外层 End 后等待流的实现逻辑。六、样例工程解剖如何验证策略是否生效样例代码位于 example/2_advanced_features/kernel/4_launch_blocking包含四个文件文件作用main.cpp三个场景的完整实现含断言与日志输出run.sh构建一次样例为每个场景分别启动新进程并设置环境变量CMakeLists.txt构建配置使用ascendc_fatbin_library编译核函数README.md/README_en.md中文/英文说明文档6.1 可验证的门闩设计核心观测手段样例要判断策略是否生效靠的是在 Kernel Launch 返回后立刻查询流状态。为此它设计了一个 Notify 门闩先在目标 Stream 上调用aclrtWaitAndResetNotify(notify_, stream_, timeout)让目标 Stream等待一个 Notify另起一个gateStream在延迟 1500ms的线程中才aclrtRecordNotify(notify_, gateStream_)释放门闩随后调用aclrtLaunchKernel并记录耗时立即查询aclrtStreamQueryconst aclrtStreamStatus expectedStatus expectBlocking ? ACL_STREAM_STATUS_COMPLETE : ACL_STREAM_STATUS_NOT_READY; context.CheckStreamStatus(expectedStatus, tag);若策略生效为同步aclrtLaunchKernel会等待门闩释放并等任务跑完才返回此时流状态应为ACL_STREAM_STATUS_COMPLETE若为异步Kernel Launch 立即返回门闩尚未释放流状态应为ACL_STREAM_STATUS_NOT_READY。之后再释放门闩、aclrtSynchronizeStreamWithTimeout同步并做输出正确性校验VerifyOutput逐元素比对 half 计算结果确保策略正确与结果正确两个维度都被验证。耗时仅打印用于观察不作为判定依据。6.2 核函数带可调计算量的 Ascend C Kernel核函数定义在 example/kernel_func/launch_blocking_kernel.cpp入口为extern C __global__ __aicore__ void launch_blocking_kernel(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR loop)它使用TPipe 双缓冲TQue流水CopyIn → Compute → CopyOut分 8 个 tile 处理 2048 个 half 元素执行z x y。关键设计是loop参数本样例固定为 1可放大计算量让同步等待的耗时差异更容易观测。输入为 half(1.0) 与 half(2.0)期望输出 half(3.0)对应main.cpp中的常量kInputX 0x3c00U、kInputY 0x4000U、kExpected 0x4200U。6.3 三种场景与预期行为场景环境变量预期行为env-control0Kernel Launch 异步返回显式同步后结果正确env-control1Kernel Launch 等待流完成后返回结果正确stream-mode0和1CTRL_BY_ENV跟随环境变量NON_BLOCKING强制异步BLOCKING强制同步non-blocking-section1嵌套区间内保持异步内层End不同步最外层End恢复同步策略并等待流完成场景一env-control环境变量控制对应RunEnvControlmain.cppconst bool envEnabled IsLaunchBlockingEnabledByEnv(); // 读取 ASCEND_RT_LAUNCH_BLOCKING SetAndCheckLaunchBlockingMode(context, ACL_STREAM_LAUNCH_BLOCKING_MODE_CTRL_BY_ENV, CTRL_BY_ENV); LaunchAndCheck(context, environment control, envEnabled);把流属性显式设为CTRL_BY_ENV然后以环境变量的值作为expectBlocking期望。run.sh会分别以0、1启动两个独立进程验证两种模式。场景二stream-mode流属性覆盖对应RunStreamModemain.cpp在同一个进程内依次切换四种状态验证覆盖语义CTRL_BY_ENV → 期望 环境变量值跟随 NON_BLOCKING → 期望 异步无论环境变量是 0 还是 1都强制异步 BLOCKING → 期望 同步无论环境变量是 0 还是 1都强制同步 restore CTRL_BY_ENV → 期望 恢复跟随环境变量注意流属性是可以在同一进程内动态修改的区别于环境变量这是流级策略相对环境变量的最大优势。场景三non-blocking-section嵌套非阻塞区间对应RunNonBlockingSectionmain.cpp前提是ASCEND_RT_LAUNCH_BLOCKING1同步模式。核心流程// 外层 Begin 内层 Begin嵌套 aclrtNonBlockingLaunchBegin(stream_, 0); aclrtNonBlockingLaunchBegin(stream_, 0); // 两次 Kernel Launch —— 都在非阻塞区间内应保持异步 first Launch(); second Launch(); CheckStreamStatus(ACL_STREAM_STATUS_NOT_READY, inside nested section); // 仍为 NOT_READY // 内层 End只减少嵌套深度不同步 aclrtNonBlockingLaunchEnd(stream_, 0); CheckStreamStatus(ACL_STREAM_STATUS_NOT_READY, after inner end); // 仍为 NOT_READY // 外层 End退出区间恢复同步策略等待流完成 aclrtNonBlockingLaunchEnd(stream_, 0); CheckStreamStatus(ACL_STREAM_STATUS_COMPLETE, after outer end); // COMPLETE退出区间后再执行一次普通 Launch 验证同步策略已恢复LaunchAndCheck(..., expectBlockingtrue)。七、构建与运行7.1 前置条件样例支持以下产品与 环境变量支持型号 一致产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√环境要求已安装 CANN 的开发环境且能解析出SOC_VERSION与ASCENDC_CMAKE_DIRrun.sh会依次调用 example/common/resolve_cann_env.sh 与 example/set_sample_env.sh 自动探测也可手动 source。7.2 编译运行步骤# 1. 切换到样例目录 cd ${git_clone_path}/example/2_advanced_features/kernel/4_launch_blocking # 2. 设置 CANN 环境变量${install_root} 替换为 CANN 安装根目录默认 /usr/local/Ascend source ${install_root}/cann/set_env.sh # 3. 构建并运行全部场景 bash run.sh7.3 构建要点CMakeLists.txt 展示了两个关键点ascendc_fatbin_library(launch_blocking_kernel ../../../kernel_func/launch_blocking_kernel.cpp)用 CANN 的ascendc.cmake把 Ascend C 核函数编译为 fatbin产物位于./out/fatbin/launch_blocking_kernel/launch_blocking_kernel.o与main.cpp中的kKernelPath对应可执行程序链接libacl_rt.so编译选项含-stdc17 -D_GLIBCXX_USE_CXX11_ABI0 -Wall -Werror。7.4run.sh的关键逻辑run.sh 的核心是构建一次、按场景分别起进程run_case() { local env_value$1 local scenario$2 env ASCEND_RT_LAUNCH_BLOCKING${env_value} ${OUTPUT_DIR}/bin/launch_blocking ${scenario} } run_case 0 env-control run_case 1 env-control run_case 0 stream-mode run_case 1 stream-mode run_case 1 non-blocking-section两个值得注意的细节不要在执行run.sh前固定导出ASCEND_RT_LAUNCH_BLOCKING脚本内部会为每个场景单独设置该变量外层导出会干扰脚本逻辑更重要的是该变量只在进程启动时生效无法在同一进程内切换。脚本使用env VARvalue cmd的方式只对单个子进程注入环境变量天然满足每个场景独立进程的要求。7.5 参考输出 ASCEND_RT_LAUNCH_BLOCKING0, scenarioenv-control [ENV] ASCEND_RT_LAUNCH_BLOCKING0 [STATUS] environment control: NOT_READY [PASS] environment control [SUCCESS] env-control ... ASCEND_RT_LAUNCH_BLOCKING1, scenarionon-blocking-section [STATUS] inside nested section: NOT_READY [STATUS] after inner end: NOT_READY [STATUS] after outer end: COMPLETE [PASS] nested non-blocking section [STATUS] blocking restored after outer end: COMPLETE [SUCCESS] non-blocking-section All launch blocking scenarios passed.[STATUS] ... NOT_READY/COMPLETE与门闩设计一一对应同步模式下 Kernel Launch 返回时流已完成异步模式下流仍在等待门闩。全部场景通过后打印All launch blocking scenarios passed.。八、完整 API 清单本样例涉及的关键 API 如下初始化与 Device 管理aclInit/aclFinalize、aclrtSetDevice/aclrtResetDeviceStream 与 Launch Blocking 控制aclrtCreateStream/aclrtDestroyStream、aclrtSetStreamAttribute/aclrtGetStreamAttribute、aclrtNonBlockingLaunchBegin/aclrtNonBlockingLaunchEnd、aclrtStreamQuery/aclrtSynchronizeStreamWithTimeoutKernel 加载与执行aclrtBinaryLoadFromFile/aclrtBinaryGetFunction/aclrtBinaryUnLoad、aclrtLaunchKernelNotify 控制aclrtCreateNotify/aclrtWaitAndResetNotify、aclrtRecordNotify/aclrtDestroyNotify内存管理与数据传输aclrtMalloc/aclrtFree、aclrtMemcpy/aclrtMemset这些接口的流管理章节说明可进一步参考 docs/zh/api_ref/06_stream_management.md含aclrtNonBlockingLaunchBegin、aclrtNonBlockingLaunchEnd、aclrtSetStreamAttribute的详细文档。九、工程注意事项与最佳实践综合样例文档、环境变量文档与源码实现总结如下工程要点环境变量是进程级一次性配置ASCEND_RT_LAUNCH_BLOCKING在 Runtime 初始化时读取修改后必须重启进程需要进程内动态切换时改用流级aclrtSetStreamAttribute。flag必须为0aclrtNonBlockingLaunchBegin/End的flag是保留参数传非 0 值会导致行为未定义或报错。Begin/End 必须同流配对、可嵌套内层End只减深度只有最外层End才可能触发等待流的同步动作。nullptr即默认流stream参数传nullptr表示当前 Context 的默认流注意与显式创建流混用时的语义。不是所有流都支持流级控制模型流、绑定流、Capture 阶段的流以及用ACL_STREAM_PERSISTENT/ACL_STREAM_CPU_SCHEDULE/ACL_STREAM_DEVICE_USE_ONLY等 flag 创建的特殊流流级 Launch Blocking 控制不生效。作用范围仅限 Kernel Launch该功能只影响 Kernel Launch以及aclmdlRIExecuteAsync等模型执行接口的同步/异步行为异步内存复制和 Event 等接口不会因此自动变为同步调用。需要同步内存拷贝时仍要显式使用同步接口或aclrtSynchronizeStream。死锁防护的正规姿势开启全局同步模式做算子调试时若存在已下发任务依赖未下发任务的场景旧算子常见应使用aclrtNonBlockingLaunchBegin/End圈出临界区或将相关流属性设为NON_BLOCKING而不是修改全局环境变量。结果校验优先于耗时观测判定策略是否生效应像样例一样以返回后的流状态 最终计算结果为准耗时仅作参考受机器负载影响波动大。十、进一步阅读样例源码example/2_advanced_features/kernel/4_launch_blocking/main.cpp、run.sh、CMakeLists.txt核函数实现example/kernel_func/launch_blocking_kernel.cpp环境变量说明docs/zh/env_vars/ASCEND_RT_LAUNCH_BLOCKING.md接口声明include/external/acl/acl_rt.hACL_STREAM_LAUNCH_BLOCKING_MODE宏、aclrtNonBlockingLaunchBegin/End、aclrtSetStreamAttribute/GetStreamAttribute、aclrtStreamAttrValue运行时底层实现入口src/runtime/api/api_c_standard_soc.cc、src/runtime/api/impl/api_impl_standard_soc.cc赞分享CANNAscend人工智能任务调度【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址https://gitcode.com/cann/runtime点击查看免费下载相关推荐DDrawCompat深度解析现代Windows系统上经典DirectX游戏兼容性终极解决方案DDrawCompat深度解析现代Windows系统上经典DirectX游戏兼容性终极解决方案 DDrawCompat是一个针对DirectX 1 7图形APCANNAscend人工智能任务调度CANN Runtime 日志不丢失策略ASCEND_LOG_SYNC_SAVE 环境变量解析与实现原理CANN Runtime 日志不丢失策略ASCEND_LOG_SYNC_SAVE 环境变量解析与实现原理 ASCEND_LOG_SYNC_SAVE 是 CANCANNAscend人工智能任务调度CANN Runtime 错误码 W40010 排查指南环境变量值非法Config Error Invalid Environment VariableCANN Runtime 错误码 W40010 排查指南环境变量值非法Config Error Invalid Environment VariableCANNAscend人工智能任务调度上一篇Asspp 项目使用与启动教程下一篇Astro Theme Pure为博客赋予简约与性能的双重魅力创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
