- CANN
- Ascend
- 人工智能
- 任务调度
【免费下载链接】runtime
本项目提供CANN运行时组件和维测功能组件。
导读
本文以 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_BLOCKING
3.1 功能与取值
环境变量ASCEND_RT_LAUNCH_BLOCKING用于控制 Kernel Launch 和模型执行任务采用同步模式或异步模式,主要面向算子调试场景。仓库中的环境变量说明文档 docs/zh/env_vars/ASCEND_RT_LAUNCH_BLOCKING.md 给出了精确的取值语义:
| 取值 | 行为 |
|---|---|
0(默认) | 接口采用异步模式,完成任务下发后返回,不等待任务执行完成 |
1 | 以下接口采用同步模式,任务执行完成后返回 |
| 其他值 | 与配置为0时相同 |
开启同步模式的接口清单包括:
aclrtLaunchKernelaclrtLaunchKernelV2aclrtLaunchKernelWithConfigaclrtLaunchKernelWithHostArgsaclrtLaunchKernelWithArgsArrayaclrtLaunchSIMTKernelWithArgsArrayaclrtLaunchSIMTKernelWithHostArgsaclmdlRIExecuteAsync
配置示例:
export ASCEND_RT_LAUNCH_BLOCKING=13.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(枚举值 6),aclrtStreamAttrValue联合体中的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并记录耗时; - 立即查询
aclrtStreamQuery:
const 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-control | 0 | Kernel Launch 异步返回;显式同步后结果正确 |
env-control | 1 | Kernel Launch 等待流完成后返回;结果正确 |
stream-mode | 0和1 | CTRL_BY_ENV跟随环境变量,NON_BLOCKING强制异步,BLOCKING强制同步 |
non-blocking-section | 1 | 嵌套区间内保持异步,内层End不同步,最外层End恢复同步策略并等待流完成 |
场景一:env-control(环境变量控制)
对应RunEnvControl(main.cpp):
const 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(流属性覆盖)
对应RunStreamMode(main.cpp),在同一个进程内依次切换四种状态,验证覆盖语义:
CTRL_BY_ENV → 期望 = 环境变量值(跟随) NON_BLOCKING → 期望 = 异步(无论环境变量是 0 还是 1,都强制异步) BLOCKING → 期望 = 同步(无论环境变量是 0 还是 1,都强制同步) restore CTRL_BY_ENV → 期望 = 恢复跟随环境变量注意:流属性是可以在同一进程内动态修改的(区别于环境变量),这是流级策略相对环境变量的最大优势。
场景三:non-blocking-section(嵌套非阻塞区间)
对应RunNonBlockingSection(main.cpp),前提是ASCEND_RT_LAUNCH_BLOCKING=1(同步模式)。核心流程:
// 外层 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(..., expectBlocking=true))。
七、构建与运行
7.1 前置条件
样例支持以下产品(与 环境变量支持型号 一致):
| 产品 | 是否支持 |
|---|---|
| Ascend 950PR / Ascend 950DT | √ |
| Atlas A3 训练系列产品 / Atlas A3 推理系列产品 | √ |
| Atlas A2 训练系列产品 / Atlas A2 推理系列产品 | √ |
环境要求:已安装 CANN 的开发环境,且能解析出SOC_VERSION与ASCENDC_CMAKE_DIR(run.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,编译选项含-std=c++17 -D_GLIBCXX_USE_CXX11_ABI=0 -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 VAR=value cmd的方式只对单个子进程注入环境变量,天然满足"每个场景独立进程"的要求。
7.5 参考输出
========== ASCEND_RT_LAUNCH_BLOCKING=0, scenario=env-control ========== [ENV] ASCEND_RT_LAUNCH_BLOCKING=0 [STATUS] environment control: NOT_READY [PASS] environment control [SUCCESS] env-control ... ========== ASCEND_RT_LAUNCH_BLOCKING=1, scenario=non-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/aclrtResetDevice - Stream 与 Launch Blocking 控制:
aclrtCreateStream/aclrtDestroyStream、aclrtSetStreamAttribute/aclrtGetStreamAttribute、aclrtNonBlockingLaunchBegin/aclrtNonBlockingLaunchEnd、aclrtStreamQuery/aclrtSynchronizeStreamWithTimeout - Kernel 加载与执行:
aclrtBinaryLoadFromFile/aclrtBinaryGetFunction/aclrtBinaryUnLoad、aclrtLaunchKernel - Notify 控制:
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必须为0:aclrtNonBlockingLaunchBegin/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.h(
ACL_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
- CANN
- Ascend
- 人工智能
- 任务调度
【免费下载链接】runtime
本项目提供CANN运行时组件和维测功能组件。
相关推荐
DDrawCompat深度解析:现代Windows系统上经典DirectX游戏兼容性终极解决方案
DDrawCompat深度解析:现代Windows系统上经典DirectX游戏兼容性终极解决方案 DDrawCompat是一个针对DirectX 1 7图形AP
CANNAscend人工智能任务调度CANN Runtime 日志不丢失策略:ASCEND_LOG_SYNC_SAVE 环境变量解析与实现原理
CANN Runtime 日志不丢失策略:ASCEND_LOG_SYNC_SAVE 环境变量解析与实现原理 ASCEND_LOG_SYNC_SAVE 是 CAN
CANNAscend人工智能任务调度CANN Runtime 错误码 W40010 排查指南:环境变量值非法(Config Error Invalid Environment Variable)
CANN Runtime 错误码 W40010 排查指南:环境变量值非法(Config Error Invalid Environment Variable)
CANNAscend人工智能任务调度
创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考