CANN Runtime 资源限制下下发 Kernellaunch_kernel_with_reslimit 样例全解析【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime导读本篇文章以 CANN runtime 开源仓库中的 1_launch_kernel_with_reslimit 样例 为核心讲解如何在进程级为 Device 设置计算资源Cube Core上限并在该限制下下发打印 Kernel 的完整流程。读完本文你将掌握aclrtSetDeviceResLimit/aclrtGetDeviceResLimit这对运行时配置接口的用法、与之配套的初始化与资源清理调用链以及如何通过脚本自动化校验资源限制是否真正生效。样例概述在资源限制下执行 Kernel当多个进程或任务共享同一个 Device 时如果不加约束某个任务可能抢占全部 AI Core 资源影响其他任务运行。CANN Runtime 提供进程级 Device 资源限制机制允许业务进程在启动 Kernel 前为当前进程指定可使用的核心数量上限。本样例演示的正是这一能力的基础闭环初始化 ACL选择 Device 0 并创建 Stream调用aclrtSetDeviceResLimit为当前进程设置Cube Core 资源限制值为 8调用aclrtGetDeviceResLimit查询当前限制值以查询到的限制值作为 blockDim 下发打印 Kernel每个 block 打印一条Hello World因此成功运行后共输出 8 条脚本自动校验输出条数是否为 8从而验证资源限制配置在真实下发路径中生效。从源码结构看该机制属于 Runtime 的“运行时配置”能力与 Stream 级资源限制aclrtSetStreamResLimit、线程级资源查询aclrtGetResInCurrentThread同属一套资源限额控制体系。产品支持情况本样例支持以下产品来自 README_en.md产品是否支持Ascend 950PR/Ascend 950DT支持Atlas A3 训练系列产品/Atlas A3 推理系列产品支持Atlas A2 训练系列产品/Atlas A2 推理系列产品支持编译与运行步骤1. 准备环境并切换目录将样例代码下载到已安装 CANN 软件的环境中切换到样例目录cd ${git_clone_path}/example/2_advanced_features/kernel/1_launch_kernel_with_reslimit其中${git_clone_path}为仓库 clone 到本地的路径。2. 设置环境变量# ${install_root} 替换为 CANN 安装根目录默认安装在 /usr/local/Ascend 目录 source ${install_root}/cann/set_env.sh # 自动识别 SOC_VERSION 和 ASCENDC_CMAKE_DIR source ${git_clone_path}/example/set_sample_env.sh第一行导入 CANN 基础环境第二行来自仓库根目录下的 set_sample_env.sh用于自动探测当前芯片的SOC_VERSION以及昇腾算子编译器AscendC的ASCENDC_CMAKE_DIR路径。3. 执行样例bash run.shrun.sh 的完整执行逻辑如下set -euo pipefail开启严格模式任何一步失败立即退出source ${EXAMPLE_DIR}/common/resolve_cann_env.sh解析 CANN 安装路径若环境变量未设置则自动加载set_sample_env.sh探测SOC_VERSION与ASCENDC_CMAKE_DIR两者缺失或ASCENDC_CMAKE_DIR/ascendc.cmake不存在时给出明确报错并退出通过 CMake 配置并编译编译命令可参考 CMakeLists.txt其通过include(${ASCENDC_CMAKE_DIR}/ascendc.cmake)引入昇腾编译框架将kernel_print.cpp编译为kernels静态库再与libacl_rt.so链接生成main可执行文件运行main并将输出通过tee同时写入output_msg.txt统计Hello World出现次数必须恰好为 8才判定成功否则打印失败原因并以非零状态退出。核心 APIDevice 资源限制的读取与设置资源限制类型枚举在 acl_rt.h 中定义了资源限制类型typedef enum { ACL_RT_DEV_RES_CUBE_CORE 0, // Cube CoreAI Core 中的矩阵计算核心资源 ACL_RT_DEV_RES_VECTOR_CORE, // Vector Core向量计算核心资源 } aclrtDevResLimitType;本样例使用ACL_RT_DEV_RES_CUBE_CORE值为 0即限制当前进程可用的 Cube Core 数量。设置资源限制aclrtSetDeviceResLimit接口声明见 acl_rt.h/** * brief Set the value of the current devices limited resources * param [in] deviceId the device id * param [in] type resource type * param [in] value resource limit value * retval ACL_SUCCESS The function is successfully executed. * retval OtherValues Failure */ ACL_FUNC_VISIBILITY aclError aclrtSetDeviceResLimit(int32_t deviceId, aclrtDevResLimitType type, uint32_t value);参数说明deviceId目标 Device 的 IDtype资源类型取aclrtDevResLimitType枚举值value要设置的资源上限值本样例设置为 8返回值ACL_SUCCESS表示成功其他值为失败。查询资源限制aclrtGetDeviceResLimit接口声明见 acl_rt.h/** * brief Get the value of the current devices limited resources * param [in] deviceId the device id * param [in] type resources type * param [out] value resources limit value * retval ACL_SUCCESS The function is successfully executed. * retval OtherValues Failure */ ACL_FUNC_VISIBILITY aclError aclrtGetDeviceResLimit(int32_t deviceId, aclrtDevResLimitType type, uint32_t* value);其中value为出参用于接收当前进程在该 Device 上对应资源的限制值。同族相关接口在 acl_rt.h 中还声明了同一资源限额体系下的其他接口便于读者按需扩展aclrtResetDeviceResLimit(deviceId)重置当前 Device 的资源限制aclrtGetStreamResLimit(stream, type, value)/aclrtSetStreamResLimit(stream, type, value)以 Stream 为粒度设置/获取资源限制aclrtUseStreamResInCurrentThread(stream)/aclrtUnuseStreamResInCurrentThread(stream)在当前线程启用/停用某 Stream 的资源配额aclrtGetResInCurrentThread(type, value)查询当前线程可用的资源限制值。主程序逐段精读资源限制 Kernel 下发完整源码见 main.cpp。整体结构清晰主流程放在一个立即执行的 lambda 中若中途出错则提前返回随后按创建顺序逆序做资源清理。资源限制查询与 Kernel 下发函数int32_t LaunchPrintKernelWithCurrentResLimit(int32_t deviceId, aclrtDevResLimitType type, aclrtStream stream) { // Query the resource limit and use it as the print kernel block dimension. uint32_t coreDim 0; CHECK_ERROR(aclrtGetDeviceResLimit(deviceId, type, coreDim)); INFO_LOG(Current device resource limit type %d is %u., static_castint32_t(type), coreDim); PrintDo(coreDim, stream); CHECK_ERROR(aclrtSynchronizeStream(stream)); return 0; }这里体现了一个关键设计用查询到的限制值直接作为 Kernel 的 blockDim。也就是说下发的 Kernel 展开的 block 数量严格等于进程被允许使用的核心数从运行结果可以反推资源限制是否真实生效。打印 Kernel 的实现打印 Kernel 定义在 kernel_print.cppextern C __global__ __aicore__ void KernelPrint() { AscendC::printf(Hello World\n); } void PrintDo(uint32_t blockDim, void* stream) { KernelPrintblockDim, nullptr, stream(); }其声明位于 kernel_ops.h。PrintDo将blockDim个 block 下发给指定 Stream每个 block 执行一次AscendC::printf从而输出 blockDim 条Hello World。main 主流程const int32_t deviceId 0; const aclrtDevResLimitType resLimitType ACL_RT_DEV_RES_CUBE_CORE; const uint32_t blockDim 8;在 Device 0 上执行资源类型为 Cube Core资源上限设置为 8。主流程依次执行初始化aclInit(nullptr)初始化 ACL 配置Device 管理aclrtSetDevice(deviceId)指定用于运算的 DeviceStream 管理aclrtCreateStream(stream)创建 Stream设置资源限制aclrtSetDeviceResLimit(deviceId, resLimitType, blockDim)将当前进程的 Cube Core 限制设为 8查询并下发调用LaunchPrintKernelWithCurrentResLimit查询限制值并据此下发 Kernel同步aclrtSynchronizeStream(stream)阻塞等待 Stream 上任务执行完成确保打印输出全部落盘后再继续。逆序清理与错误处理主流程结束后按“后创建、先释放”的顺序做清理任何一步失败都会将最终结果置为失败aclrtDestroyStreamForce(stream)强制销毁StreamaclrtResetDeviceForce(deviceId)强制复位当前运算 Device回收 Device 上的资源aclFinalize()去初始化。UpdateFinalResultOnError封装了对每个清理接口返回码的检查任一接口返回非ACL_SUCCESS即打印错误日志Operation failed: %s returned error code %d并将finalResult置为 -1。示例输出与校验逻辑正常运行时的输出如下来自 README_en.md[INFO]: Current compile soc version is ... Configuring CMake... Building... [INFO] ACL initialized. [INFO] Device 0 selected. [INFO] Stream created. [INFO] Device resource limit type 0 set to 8. [INFO] Current device resource limit type 0 is 8. Hello World Hello World Hello World Hello World Hello World Hello World Hello World Hello World [INFO] Run the launch_kernel_with_reslimit sample successfully. [SUCCESS]: Launch kernels under resource limits successfully.日志解读[INFO] Device resource limit type 0 set to 8.类型 0 即ACL_RT_DEV_RES_CUBE_CORE设置成功[INFO] Current device resource limit type 0 is 8.查询结果与设置值一致读写闭环成立8 条Hello World说明 Kernel 实际以 8 个 block 展开执行资源限制被 Kernel 下发路径真正采纳末尾的[SUCCESS]来自 run.sh 的自动化校验grep -c Hello World统计结果为 8 时判定成功否则输出[FAILURE]: Expected 8 occurrences of ...并以失败退出。总结launch_kernel_with_reslimit样例以最小的代码量串联起了 CANN Runtime 资源限额能力的完整链路通过aclrtSetDeviceResLimit设置进程级 Cube Core 上限通过aclrtGetDeviceResLimit回读验证再以限制值驱动 Kernel 下发最后由脚本对输出做确定性断言。这套“设置 → 查询 → 使用 → 校验”的模式可直接迁移到多任务共享 Device 的资源隔离场景也可结合 acl_rt.h 中同族的 Stream 级与线程级资源限制接口构造更细粒度的资源管控方案。【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
CANN Runtime 资源限制下下发 Kernellaunch_kernel_with_reslimit 样例全解析【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime导读本篇文章以 CANN runtime 开源仓库中的 1_launch_kernel_with_reslimit 样例 为核心讲解如何在进程级为 Device 设置计算资源Cube Core上限并在该限制下下发打印 Kernel 的完整流程。读完本文你将掌握aclrtSetDeviceResLimit/aclrtGetDeviceResLimit这对运行时配置接口的用法、与之配套的初始化与资源清理调用链以及如何通过脚本自动化校验资源限制是否真正生效。样例概述在资源限制下执行 Kernel当多个进程或任务共享同一个 Device 时如果不加约束某个任务可能抢占全部 AI Core 资源影响其他任务运行。CANN Runtime 提供进程级 Device 资源限制机制允许业务进程在启动 Kernel 前为当前进程指定可使用的核心数量上限。本样例演示的正是这一能力的基础闭环初始化 ACL选择 Device 0 并创建 Stream调用aclrtSetDeviceResLimit为当前进程设置Cube Core 资源限制值为 8调用aclrtGetDeviceResLimit查询当前限制值以查询到的限制值作为 blockDim 下发打印 Kernel每个 block 打印一条Hello World因此成功运行后共输出 8 条脚本自动校验输出条数是否为 8从而验证资源限制配置在真实下发路径中生效。从源码结构看该机制属于 Runtime 的“运行时配置”能力与 Stream 级资源限制aclrtSetStreamResLimit、线程级资源查询aclrtGetResInCurrentThread同属一套资源限额控制体系。产品支持情况本样例支持以下产品来自 README_en.md产品是否支持Ascend 950PR/Ascend 950DT支持Atlas A3 训练系列产品/Atlas A3 推理系列产品支持Atlas A2 训练系列产品/Atlas A2 推理系列产品支持编译与运行步骤1. 准备环境并切换目录将样例代码下载到已安装 CANN 软件的环境中切换到样例目录cd ${git_clone_path}/example/2_advanced_features/kernel/1_launch_kernel_with_reslimit其中${git_clone_path}为仓库 clone 到本地的路径。2. 设置环境变量# ${install_root} 替换为 CANN 安装根目录默认安装在 /usr/local/Ascend 目录 source ${install_root}/cann/set_env.sh # 自动识别 SOC_VERSION 和 ASCENDC_CMAKE_DIR source ${git_clone_path}/example/set_sample_env.sh第一行导入 CANN 基础环境第二行来自仓库根目录下的 set_sample_env.sh用于自动探测当前芯片的SOC_VERSION以及昇腾算子编译器AscendC的ASCENDC_CMAKE_DIR路径。3. 执行样例bash run.shrun.sh 的完整执行逻辑如下set -euo pipefail开启严格模式任何一步失败立即退出source ${EXAMPLE_DIR}/common/resolve_cann_env.sh解析 CANN 安装路径若环境变量未设置则自动加载set_sample_env.sh探测SOC_VERSION与ASCENDC_CMAKE_DIR两者缺失或ASCENDC_CMAKE_DIR/ascendc.cmake不存在时给出明确报错并退出通过 CMake 配置并编译编译命令可参考 CMakeLists.txt其通过include(${ASCENDC_CMAKE_DIR}/ascendc.cmake)引入昇腾编译框架将kernel_print.cpp编译为kernels静态库再与libacl_rt.so链接生成main可执行文件运行main并将输出通过tee同时写入output_msg.txt统计Hello World出现次数必须恰好为 8才判定成功否则打印失败原因并以非零状态退出。核心 APIDevice 资源限制的读取与设置资源限制类型枚举在 acl_rt.h 中定义了资源限制类型typedef enum { ACL_RT_DEV_RES_CUBE_CORE 0, // Cube CoreAI Core 中的矩阵计算核心资源 ACL_RT_DEV_RES_VECTOR_CORE, // Vector Core向量计算核心资源 } aclrtDevResLimitType;本样例使用ACL_RT_DEV_RES_CUBE_CORE值为 0即限制当前进程可用的 Cube Core 数量。设置资源限制aclrtSetDeviceResLimit接口声明见 acl_rt.h/** * brief Set the value of the current devices limited resources * param [in] deviceId the device id * param [in] type resource type * param [in] value resource limit value * retval ACL_SUCCESS The function is successfully executed. * retval OtherValues Failure */ ACL_FUNC_VISIBILITY aclError aclrtSetDeviceResLimit(int32_t deviceId, aclrtDevResLimitType type, uint32_t value);参数说明deviceId目标 Device 的 IDtype资源类型取aclrtDevResLimitType枚举值value要设置的资源上限值本样例设置为 8返回值ACL_SUCCESS表示成功其他值为失败。查询资源限制aclrtGetDeviceResLimit接口声明见 acl_rt.h/** * brief Get the value of the current devices limited resources * param [in] deviceId the device id * param [in] type resources type * param [out] value resources limit value * retval ACL_SUCCESS The function is successfully executed. * retval OtherValues Failure */ ACL_FUNC_VISIBILITY aclError aclrtGetDeviceResLimit(int32_t deviceId, aclrtDevResLimitType type, uint32_t* value);其中value为出参用于接收当前进程在该 Device 上对应资源的限制值。同族相关接口在 acl_rt.h 中还声明了同一资源限额体系下的其他接口便于读者按需扩展aclrtResetDeviceResLimit(deviceId)重置当前 Device 的资源限制aclrtGetStreamResLimit(stream, type, value)/aclrtSetStreamResLimit(stream, type, value)以 Stream 为粒度设置/获取资源限制aclrtUseStreamResInCurrentThread(stream)/aclrtUnuseStreamResInCurrentThread(stream)在当前线程启用/停用某 Stream 的资源配额aclrtGetResInCurrentThread(type, value)查询当前线程可用的资源限制值。主程序逐段精读资源限制 Kernel 下发完整源码见 main.cpp。整体结构清晰主流程放在一个立即执行的 lambda 中若中途出错则提前返回随后按创建顺序逆序做资源清理。资源限制查询与 Kernel 下发函数int32_t LaunchPrintKernelWithCurrentResLimit(int32_t deviceId, aclrtDevResLimitType type, aclrtStream stream) { // Query the resource limit and use it as the print kernel block dimension. uint32_t coreDim 0; CHECK_ERROR(aclrtGetDeviceResLimit(deviceId, type, coreDim)); INFO_LOG(Current device resource limit type %d is %u., static_castint32_t(type), coreDim); PrintDo(coreDim, stream); CHECK_ERROR(aclrtSynchronizeStream(stream)); return 0; }这里体现了一个关键设计用查询到的限制值直接作为 Kernel 的 blockDim。也就是说下发的 Kernel 展开的 block 数量严格等于进程被允许使用的核心数从运行结果可以反推资源限制是否真实生效。打印 Kernel 的实现打印 Kernel 定义在 kernel_print.cppextern C __global__ __aicore__ void KernelPrint() { AscendC::printf(Hello World\n); } void PrintDo(uint32_t blockDim, void* stream) { KernelPrintblockDim, nullptr, stream(); }其声明位于 kernel_ops.h。PrintDo将blockDim个 block 下发给指定 Stream每个 block 执行一次AscendC::printf从而输出 blockDim 条Hello World。main 主流程const int32_t deviceId 0; const aclrtDevResLimitType resLimitType ACL_RT_DEV_RES_CUBE_CORE; const uint32_t blockDim 8;在 Device 0 上执行资源类型为 Cube Core资源上限设置为 8。主流程依次执行初始化aclInit(nullptr)初始化 ACL 配置Device 管理aclrtSetDevice(deviceId)指定用于运算的 DeviceStream 管理aclrtCreateStream(stream)创建 Stream设置资源限制aclrtSetDeviceResLimit(deviceId, resLimitType, blockDim)将当前进程的 Cube Core 限制设为 8查询并下发调用LaunchPrintKernelWithCurrentResLimit查询限制值并据此下发 Kernel同步aclrtSynchronizeStream(stream)阻塞等待 Stream 上任务执行完成确保打印输出全部落盘后再继续。逆序清理与错误处理主流程结束后按“后创建、先释放”的顺序做清理任何一步失败都会将最终结果置为失败aclrtDestroyStreamForce(stream)强制销毁StreamaclrtResetDeviceForce(deviceId)强制复位当前运算 Device回收 Device 上的资源aclFinalize()去初始化。UpdateFinalResultOnError封装了对每个清理接口返回码的检查任一接口返回非ACL_SUCCESS即打印错误日志Operation failed: %s returned error code %d并将finalResult置为 -1。示例输出与校验逻辑正常运行时的输出如下来自 README_en.md[INFO]: Current compile soc version is ... Configuring CMake... Building... [INFO] ACL initialized. [INFO] Device 0 selected. [INFO] Stream created. [INFO] Device resource limit type 0 set to 8. [INFO] Current device resource limit type 0 is 8. Hello World Hello World Hello World Hello World Hello World Hello World Hello World Hello World [INFO] Run the launch_kernel_with_reslimit sample successfully. [SUCCESS]: Launch kernels under resource limits successfully.日志解读[INFO] Device resource limit type 0 set to 8.类型 0 即ACL_RT_DEV_RES_CUBE_CORE设置成功[INFO] Current device resource limit type 0 is 8.查询结果与设置值一致读写闭环成立8 条Hello World说明 Kernel 实际以 8 个 block 展开执行资源限制被 Kernel 下发路径真正采纳末尾的[SUCCESS]来自 run.sh 的自动化校验grep -c Hello World统计结果为 8 时判定成功否则输出[FAILURE]: Expected 8 occurrences of ...并以失败退出。总结launch_kernel_with_reslimit样例以最小的代码量串联起了 CANN Runtime 资源限额能力的完整链路通过aclrtSetDeviceResLimit设置进程级 Cube Core 上限通过aclrtGetDeviceResLimit回读验证再以限制值驱动 Kernel 下发最后由脚本对输出做确定性断言。这套“设置 → 查询 → 使用 → 校验”的模式可直接迁移到多任务共享 Device 的资源隔离场景也可结合 acl_rt.h 中同族的 Stream 级与线程级资源限制接口构造更细粒度的资源管控方案。【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
