尧图网站建设 
CANN Runtime 资源限制下 Kernel 下发实战基于 aclrtSetDeviceResLimit/aclrtGetDeviceResLimit 的进程级 Device 资源配额示例解析【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime导读本文以 CANN runtime 开源仓库中的1_launch_kernel_with_reslimit样例为线索完整讲解如何在进程内设置 Device 级资源限制Resource Limit并在此限制下下发 Ascend 算子Kernel的完整流程。读者将掌握aclrtSetDeviceResLimit与aclrtGetDeviceResLimit两个核心 API 的用法与枚举含义理解设置 Cube Core 资源限制 → 查询当前限制值 → 按限制值设置核函数 blockDim → 下发打印 Kernel的实战链路并能够独立编译、运行和校验该样例从而在多任务、多进程共享 Device 的场景下合理规划 AI Core 资源配额。1. 样例概述与适用产品1.1 样例要解决的问题现代异构计算中一个物理 Device 可能同时被多个进程、多个任务使用。若不加以限制单个进程可能抢占过多 AI Core 计算资源导致其他进程无法按时获得执行机会。CANN Runtime 提供了进程级Process-levelDevice 资源限制能力进程可以在初始化后、下发 Kernel 前通过aclrtSetDeviceResLimit接口为当前进程声明对指定类型计算核心如 Cube Core的资源配额后续 Kernel 下发即可依据该配额运行。本样例演示的正是这一基础流程在当前进程设置Cube Core资源限制通过aclrtGetDeviceResLimit查询当前限制值将查询到的限制值作为打印 Kernel 的 blockDim块维度下发出打印 Kernel运行后观察 Kernel 输出的Hello World日志并由脚本校验输出数量是否符合预期。1.2 样例在产品支持情况根据样例文档 README.md本样例支持以下产品产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品/Atlas A3 推理系列产品√Atlas A2 训练系列产品/Atlas A2 推理系列产品√注意不同产品对资源限制类型例如 Cube Core / Vector Core的支持能力可能不同实际使用前请确认当前环境的 SOC 版本支持情况。2. 编译运行全流程2.1 环境准备编译运行该样例需要满足以下前提已安装 CANN 软件包默认安装根目录为/usr/local/Ascend下文以${install_root}代替已下载 CANN runtime 开源仓库代码下文以${git_clone_path}代替仓库根目录目标环境为上述支持的昇腾产品。2.2 编译运行步骤步骤 1切换到样例目录。cd ${git_clone_path}/example/2_advanced_features/kernel/1_launch_kernel_with_reslimit步骤 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其中set_env.sh由 CANN 安装包提供用于加载运行时的动态库搜索路径等基础环境set_sample_env.sh 是仓库内提供的样例级环境脚本用于自动识别SOC_VERSION与ASCENDC_CMAKE_DIR两个关键编译变量。步骤 3执行运行脚本。bash run.sh2.3 运行脚本做了什么run.sh 是整个样例的自动构建与校验入口其核心流程如下解析 CANN 环境通过source ${EXAMPLE_DIR}/common/resolve_cann_env.sh加载 resolve_cann_env.sh解析出ASCEND_INSTALL_PATH等 CANN 安装信息自动识别样例环境调用detect_sample_env()函数若SOC_VERSION与ASCENDC_CMAKE_DIR尚未设置则尝试 source 仓库根目录下的set_sample_env.sh自动探测两者任一缺失时给出明确报错并退出前置校验脚本显式检查${ASCENDC_CMAKE_DIR}/ascendc.cmake是否存在确保昇腾编译器AscendC的 CMake 工具链可用构建使用 CMake 配置并编译-j$(nproc)并行编译运行与校验执行编译产物./main将其标准输出同时落盘到output_msg.txt随后使用grep -c Hello World统计Hello World出现次数断言若Hello World出现次数不等于 8则判定为失败并退出非零否则输出[SUCCESS]: Launch kernels under resource limits successfully.。由此可见8是本样例的关键预期数值它既是进程中设置的 Cube Core 资源限制值blockDim也是打印 Kernel 实际执行后应产生的Hello World日志条数。2.4 CMake 构建配置要点CMakeLists.txt 揭示了该样例的编译模型include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) ascendc_library(kernels STATIC ../../../kernel_func/kernel_print.cpp)通过include(${ASCENDC_CMAKE_DIR}/ascendc.cmake)引入 AscendC 算子编译工具链使用ascendc_library将设备侧 Kernel 源码 kernel_print.cpp 编译为静态库kernels宿主侧main.cpp编译为可执行文件main并通过-D_GLIBCXX_USE_CXX11_ABI0保持与 CANN 动态库一致的 ABI链接${ASCEND_CANN_PACKAGE_PATH}/lib64/libacl_rt.so即 ACL Runtime 动态库RUN_MODE固定为npuCMAKE_BUILD_TYPE默认Debug。3. 核心 API 与底层实现3.1 资源限制类型枚举资源限制的类型由枚举aclrtDevResLimitType定义声明位于 acl_rt.htypedef enum { ACL_RT_DEV_RES_CUBE_CORE 0, ACL_RT_DEV_RES_VECTOR_CORE, } aclrtDevResLimitType;枚举值含义ACL_RT_DEV_RES_CUBE_CORE 0Cube Core 资源限制本样例使用该类型ACL_RT_DEV_RES_VECTOR_CORE 1Vector Core 资源限制由枚举递增关系可推断其值为 13.2 进程级 Device 资源限制 API与进程级 Device 资源限制相关的接口全部声明于 acl_rt.h 中除本样例使用的两个接口外还包含 Stream 级与线程级的能力API声明位置作用aclrtSetDeviceResLimit(int32_t deviceId, aclrtDevResLimitType type, uint32_t value)acl_rt.h#L4188设置当前进程在指定 Device 上的某类资源限制aclrtGetDeviceResLimit(int32_t deviceId, aclrtDevResLimitType type, uint32_t* value)acl_rt.h#L4177查询当前进程在指定 Device 上的某类资源限制aclrtSetStreamResLimit(aclrtStream stream, aclrtDevResLimitType type, uint32_t value)acl_rt.h#L4219设置指定 Stream 的资源限制aclrtGetStreamResLimit(aclrtStream stream, aclrtDevResLimitType type, uint32_t* value)acl_rt.h#L4208查询指定 Stream 的资源限制aclrtGetResInCurrentThread(aclrtDevResLimitType type, uint32_t* value)acl_rt.h#L4256查询当前线程可用的资源量从接口集合可以看出CANN Runtime 提供了Device 级进程维度、Stream 级与线程级三个层面的资源限制能力开发者可以根据任务的调度粒度选择最合适的配额维度。3.3 ACL 层的实现路径在仓库的 ACL 实现层aclrtSetDeviceResLimit与aclrtGetDeviceResLimit分别由 device.cpp 中的aclrtGetDeviceResLimitImpl/aclrtSetDeviceResLimitImpl承载aclError aclrtGetDeviceResLimitImpl(int32_t deviceId, aclrtDevResLimitType type, uint32_t* value) { ACL_PROFILING_REG(acl::AclProfType::AclrtGetDeviceResLimit); // 日志deviceId、type ACL_REQUIRES_RTS_OK(rtsGetDeviceResLimit(deviceId, static_castrtDevResLimitType_t(type), value)); // 日志成功 } aclError aclrtSetDeviceResLimitImpl(int32_t deviceId, aclrtDevResLimitType type, uint32_t value) { ACL_PROFILING_REG(acl::AclProfType::AclrtSetDeviceResLimit); // 日志deviceId、type、value ACL_REQUIRES_RTS_OK(rtsSetDeviceResLimit(deviceId, static_castrtDevResLimitType_t(type), value)); // 日志成功 }从源码结构看这两个 ACL 接口均先通过ACL_PROFILING_REG登记 profiling 采样点可通过ACL_PROFILING_REG宏与AclProfType::AclrtGetDeviceResLimit/AclrtSetDeviceResLimit关联到 prof 能力随后将aclrtDevResLimitType强转为 RTSRuntime System层的rtDevResLimitType_t枚举并下沉调用rtsSetDeviceResLimit/rtsGetDeviceResLimit完成实际的资源限制设置与查询最终由ACL_REQUIRES_RTS_OK统一将 RTS 层返回码转换为aclError返回给用户。这一ACL 接口 → RTS 接口的两层封装与仓库中 pkg_inc/runtime 下的 RTS 头文件体系保持一致。4. 样例源码逐段解析4.1 主流程代码main.cpp 是宿主侧核心逻辑采用典型的初始化 → 设 Device → 建 Stream → 设资源限制 → 下发 Kernel → 清理结构int32_t main() { const int32_t deviceId 0; const aclrtDevResLimitType resLimitType ACL_RT_DEV_RES_CUBE_CORE; const uint32_t blockDim 8; aclrtStream stream nullptr; bool aclInitialized false; bool deviceSet false; bool streamCreated false; const int32_t result []() - int32_t { // 1. 初始化 ACL CHECK_ERROR(aclInit(nullptr)); aclInitialized true; // 2. 指定运算 Device CHECK_ERROR(aclrtSetDevice(deviceId)); deviceSet true; // 3. 创建 Stream CHECK_ERROR(aclrtCreateStream(stream)); streamCreated true; // 4. 下发 Kernel 前设置进程级 Device 资源限制 CHECK_ERROR(aclrtSetDeviceResLimit(deviceId, resLimitType, blockDim)); INFO_LOG(Device resource limit type %d set to %u., ...); // 5. 查询限制值并以此为 blockDim 下发打印 Kernel if (LaunchPrintKernelWithCurrentResLimit(deviceId, resLimitType, stream) ! 0) { return -1; } return 0; }(); // ...清理与结果汇总 }各步骤的 CANN RUNTIME API 映射关系如下功能分类关键接口在样例中的作用初始化aclInit/aclFinalize初始化 / 去初始化 ACL 运行环境Device 管理aclrtSetDevice指定用于运算的 Device本样例为 device 0Device 管理aclrtResetDeviceForce强制复位 Device回收 Device 上的资源Stream 管理aclrtCreateStream创建 StreamStream 管理aclrtSynchronizeStream阻塞等待 Stream 上任务执行完成Stream 管理aclrtDestroyStreamForce强制销毁 Stream运行时配置aclrtSetDeviceResLimit设置当前进程的 Device 资源限制运行时配置aclrtGetDeviceResLimit获取当前进程的 Device 资源限制4.2 查询限制值并下发 Kernel资源限制查询与 Kernel 下发被封装在LaunchPrintKernelWithCurrentResLimit函数中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而是先调用aclrtGetDeviceResLimit查询当前进程的资源限制值coreDim再将该值作为 blockDim 传给PrintDo。也就是说打印 Kernel 的实际执行规模8 个块完全由此前aclrtSetDeviceResLimit(deviceId, ACL_RT_DEV_RES_CUBE_CORE, 8)设置的配额决定从而验证资源限制真正作用于 Kernel 下发这一核心语义。4.3 设备侧打印 Kernel打印 Kernel 定义于 kernel_print.cpp宿主侧的封装函数PrintDo声明在 kernel_ops.h 中#include kernel_operator.h #include kernel_ops.h extern C __global__ __aicore__ void KernelPrint() { AscendC::printf(Hello World\n); } void PrintDo(uint32_t blockDim, void* stream) { KernelPrintblockDim, nullptr, stream(); }要点说明KernelPrint是__global__ __aicore__修饰的 Device 侧核函数内部调用AscendC::printf输出Hello WorldPrintDo以类 CUDA 的blockDim, nullptr, stream语法将 Kernel 下发到指定 Stream其中 blockDim 即来自资源限制查询值本样例为 8每个块都会执行一次AscendC::printf(Hello World\n)因此 8 个块产生 8 行Hello World这与run.sh中的grep -c校验目标严格对应。4.4 错误处理与资源清理样例采用了精细的错误处理策略主流程使用 utils.h 中定义的CHECK_ERROR宏任一 ACL 调用返回非ACL_SUCCESS即打印错误日志并直接返回-1清理阶段使用UpdateFinalResultOnError辅助函数即便主流程已经失败也会继续依次执行aclrtDestroyStreamForce、aclrtResetDeviceForce、aclFinalize尽力回收 Stream、Device 与 ACL 环境资源并合并最终的返回结果aclrtDestroyStreamForce与aclrtResetDeviceForce均为强制语义接口用于在异常路径下确保资源释放。5. 运行结果解读5.1 预期输出编译运行后样例的标准输出如下示例[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.5.2 输出如何印证资源限制生效输出中的关键信息链可以这样解读[INFO] Device resource limit type 0 set to 8.——type 0即ACL_RT_DEV_RES_CUBE_CORE进程级 Cube Core 资源限制被设置为 8[INFO] Current device resource limit type 0 is 8.——aclrtGetDeviceResLimit查询结果与设置值一致说明限制设置成功且可被查询恰好8 行Hello World——打印 Kernel 以查询到的限制值 8 作为 blockDim 执行每个块打印一行实际执行规模与资源配额吻合[SUCCESS]: Launch kernels under resource limits successfully.——run.sh中grep -c Hello World计数等于 8 的断言通过样例整体成功。6. 延伸资源限制的典型应用场景结合本样例与仓库中的 API 全貌进程级 Device 资源限制在工程实践中通常用于以下场景多进程共享 Device 时的资源隔离多个进程跑在同一物理 Device 上时通过aclrtSetDeviceResLimit为各自进程声明 Cube/Vector Core 配额避免单个进程独占计算核心任务级资源预算控制结合 Stream 级接口aclrtSetStreamResLimit/aclrtGetStreamResLimit可以在更细粒度上约束特定 Stream 的可用资源动态资源感知借助aclrtGetDeviceResLimit/aclrtGetResInCurrentThread在运行时读取当前可用资源据此动态调整 Kernel 的 blockDim 等调度参数实现资源自适应的下发策略本样例正是这一思路的最小演示。需要注意的是资源限制的具体生效范围、枚举取值与硬件能力的对应关系以当前 CANN 版本的接口声明与产品文档为准本仓库中 acl_rt.h 的枚举定义与 device.cpp 的实现可作为接口层事实依据。7. 小结本样例用最小的代码量展示了 CANN Runtime 设置 Device 资源限制 → 查询限制 → 在限制下下发 Kernel 的完整闭环进程级资源限制通过aclrtSetDeviceResLimit设置、aclrtGetDeviceResLimit查询资源类型由aclrtDevResLimitType枚举Cube Core / Vector Core界定样例将查询到的限制值直接用作打印 Kernel 的 blockDim并用 8 行Hello World与脚本断言验证了限制确实作用于 Kernel 下发ACL 层实现最终下沉到rtsSetDeviceResLimit/rtsGetDeviceResLimit并带有 profiling 采样点为上层开发者提供了可靠、可观测的资源配额控制能力。对于希望在多进程、多任务场景下精细管理 AI Core 资源的开发者本样例是一个可以直接复用的最小范式将aclrtSetDeviceResLimit调用置于初始化之后、Kernel 下发之前即可为整个进程划定计算资源边界再配合aclrtGetDeviceResLimit动态适配 Kernel 调度规模。【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CANN Runtime 资源限制下 Kernel 下发实战：基于 aclrtSetDeviceResLimit/aclrtGetDeviceResLimit 的进程级 Device 资源配额示例解析
CANN Runtime 资源限制下 Kernel 下发实战基于 aclrtSetDeviceResLimit/aclrtGetDeviceResLimit 的进程级 Device 资源配额示例解析【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime导读本文以 CANN runtime 开源仓库中的1_launch_kernel_with_reslimit样例为线索完整讲解如何在进程内设置 Device 级资源限制Resource Limit并在此限制下下发 Ascend 算子Kernel的完整流程。读者将掌握aclrtSetDeviceResLimit与aclrtGetDeviceResLimit两个核心 API 的用法与枚举含义理解设置 Cube Core 资源限制 → 查询当前限制值 → 按限制值设置核函数 blockDim → 下发打印 Kernel的实战链路并能够独立编译、运行和校验该样例从而在多任务、多进程共享 Device 的场景下合理规划 AI Core 资源配额。1. 样例概述与适用产品1.1 样例要解决的问题现代异构计算中一个物理 Device 可能同时被多个进程、多个任务使用。若不加以限制单个进程可能抢占过多 AI Core 计算资源导致其他进程无法按时获得执行机会。CANN Runtime 提供了进程级Process-levelDevice 资源限制能力进程可以在初始化后、下发 Kernel 前通过aclrtSetDeviceResLimit接口为当前进程声明对指定类型计算核心如 Cube Core的资源配额后续 Kernel 下发即可依据该配额运行。本样例演示的正是这一基础流程在当前进程设置Cube Core资源限制通过aclrtGetDeviceResLimit查询当前限制值将查询到的限制值作为打印 Kernel 的 blockDim块维度下发出打印 Kernel运行后观察 Kernel 输出的Hello World日志并由脚本校验输出数量是否符合预期。1.2 样例在产品支持情况根据样例文档 README.md本样例支持以下产品产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品/Atlas A3 推理系列产品√Atlas A2 训练系列产品/Atlas A2 推理系列产品√注意不同产品对资源限制类型例如 Cube Core / Vector Core的支持能力可能不同实际使用前请确认当前环境的 SOC 版本支持情况。2. 编译运行全流程2.1 环境准备编译运行该样例需要满足以下前提已安装 CANN 软件包默认安装根目录为/usr/local/Ascend下文以${install_root}代替已下载 CANN runtime 开源仓库代码下文以${git_clone_path}代替仓库根目录目标环境为上述支持的昇腾产品。2.2 编译运行步骤步骤 1切换到样例目录。cd ${git_clone_path}/example/2_advanced_features/kernel/1_launch_kernel_with_reslimit步骤 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其中set_env.sh由 CANN 安装包提供用于加载运行时的动态库搜索路径等基础环境set_sample_env.sh 是仓库内提供的样例级环境脚本用于自动识别SOC_VERSION与ASCENDC_CMAKE_DIR两个关键编译变量。步骤 3执行运行脚本。bash run.sh2.3 运行脚本做了什么run.sh 是整个样例的自动构建与校验入口其核心流程如下解析 CANN 环境通过source ${EXAMPLE_DIR}/common/resolve_cann_env.sh加载 resolve_cann_env.sh解析出ASCEND_INSTALL_PATH等 CANN 安装信息自动识别样例环境调用detect_sample_env()函数若SOC_VERSION与ASCENDC_CMAKE_DIR尚未设置则尝试 source 仓库根目录下的set_sample_env.sh自动探测两者任一缺失时给出明确报错并退出前置校验脚本显式检查${ASCENDC_CMAKE_DIR}/ascendc.cmake是否存在确保昇腾编译器AscendC的 CMake 工具链可用构建使用 CMake 配置并编译-j$(nproc)并行编译运行与校验执行编译产物./main将其标准输出同时落盘到output_msg.txt随后使用grep -c Hello World统计Hello World出现次数断言若Hello World出现次数不等于 8则判定为失败并退出非零否则输出[SUCCESS]: Launch kernels under resource limits successfully.。由此可见8是本样例的关键预期数值它既是进程中设置的 Cube Core 资源限制值blockDim也是打印 Kernel 实际执行后应产生的Hello World日志条数。2.4 CMake 构建配置要点CMakeLists.txt 揭示了该样例的编译模型include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) ascendc_library(kernels STATIC ../../../kernel_func/kernel_print.cpp)通过include(${ASCENDC_CMAKE_DIR}/ascendc.cmake)引入 AscendC 算子编译工具链使用ascendc_library将设备侧 Kernel 源码 kernel_print.cpp 编译为静态库kernels宿主侧main.cpp编译为可执行文件main并通过-D_GLIBCXX_USE_CXX11_ABI0保持与 CANN 动态库一致的 ABI链接${ASCEND_CANN_PACKAGE_PATH}/lib64/libacl_rt.so即 ACL Runtime 动态库RUN_MODE固定为npuCMAKE_BUILD_TYPE默认Debug。3. 核心 API 与底层实现3.1 资源限制类型枚举资源限制的类型由枚举aclrtDevResLimitType定义声明位于 acl_rt.htypedef enum { ACL_RT_DEV_RES_CUBE_CORE 0, ACL_RT_DEV_RES_VECTOR_CORE, } aclrtDevResLimitType;枚举值含义ACL_RT_DEV_RES_CUBE_CORE 0Cube Core 资源限制本样例使用该类型ACL_RT_DEV_RES_VECTOR_CORE 1Vector Core 资源限制由枚举递增关系可推断其值为 13.2 进程级 Device 资源限制 API与进程级 Device 资源限制相关的接口全部声明于 acl_rt.h 中除本样例使用的两个接口外还包含 Stream 级与线程级的能力API声明位置作用aclrtSetDeviceResLimit(int32_t deviceId, aclrtDevResLimitType type, uint32_t value)acl_rt.h#L4188设置当前进程在指定 Device 上的某类资源限制aclrtGetDeviceResLimit(int32_t deviceId, aclrtDevResLimitType type, uint32_t* value)acl_rt.h#L4177查询当前进程在指定 Device 上的某类资源限制aclrtSetStreamResLimit(aclrtStream stream, aclrtDevResLimitType type, uint32_t value)acl_rt.h#L4219设置指定 Stream 的资源限制aclrtGetStreamResLimit(aclrtStream stream, aclrtDevResLimitType type, uint32_t* value)acl_rt.h#L4208查询指定 Stream 的资源限制aclrtGetResInCurrentThread(aclrtDevResLimitType type, uint32_t* value)acl_rt.h#L4256查询当前线程可用的资源量从接口集合可以看出CANN Runtime 提供了Device 级进程维度、Stream 级与线程级三个层面的资源限制能力开发者可以根据任务的调度粒度选择最合适的配额维度。3.3 ACL 层的实现路径在仓库的 ACL 实现层aclrtSetDeviceResLimit与aclrtGetDeviceResLimit分别由 device.cpp 中的aclrtGetDeviceResLimitImpl/aclrtSetDeviceResLimitImpl承载aclError aclrtGetDeviceResLimitImpl(int32_t deviceId, aclrtDevResLimitType type, uint32_t* value) { ACL_PROFILING_REG(acl::AclProfType::AclrtGetDeviceResLimit); // 日志deviceId、type ACL_REQUIRES_RTS_OK(rtsGetDeviceResLimit(deviceId, static_castrtDevResLimitType_t(type), value)); // 日志成功 } aclError aclrtSetDeviceResLimitImpl(int32_t deviceId, aclrtDevResLimitType type, uint32_t value) { ACL_PROFILING_REG(acl::AclProfType::AclrtSetDeviceResLimit); // 日志deviceId、type、value ACL_REQUIRES_RTS_OK(rtsSetDeviceResLimit(deviceId, static_castrtDevResLimitType_t(type), value)); // 日志成功 }从源码结构看这两个 ACL 接口均先通过ACL_PROFILING_REG登记 profiling 采样点可通过ACL_PROFILING_REG宏与AclProfType::AclrtGetDeviceResLimit/AclrtSetDeviceResLimit关联到 prof 能力随后将aclrtDevResLimitType强转为 RTSRuntime System层的rtDevResLimitType_t枚举并下沉调用rtsSetDeviceResLimit/rtsGetDeviceResLimit完成实际的资源限制设置与查询最终由ACL_REQUIRES_RTS_OK统一将 RTS 层返回码转换为aclError返回给用户。这一ACL 接口 → RTS 接口的两层封装与仓库中 pkg_inc/runtime 下的 RTS 头文件体系保持一致。4. 样例源码逐段解析4.1 主流程代码main.cpp 是宿主侧核心逻辑采用典型的初始化 → 设 Device → 建 Stream → 设资源限制 → 下发 Kernel → 清理结构int32_t main() { const int32_t deviceId 0; const aclrtDevResLimitType resLimitType ACL_RT_DEV_RES_CUBE_CORE; const uint32_t blockDim 8; aclrtStream stream nullptr; bool aclInitialized false; bool deviceSet false; bool streamCreated false; const int32_t result []() - int32_t { // 1. 初始化 ACL CHECK_ERROR(aclInit(nullptr)); aclInitialized true; // 2. 指定运算 Device CHECK_ERROR(aclrtSetDevice(deviceId)); deviceSet true; // 3. 创建 Stream CHECK_ERROR(aclrtCreateStream(stream)); streamCreated true; // 4. 下发 Kernel 前设置进程级 Device 资源限制 CHECK_ERROR(aclrtSetDeviceResLimit(deviceId, resLimitType, blockDim)); INFO_LOG(Device resource limit type %d set to %u., ...); // 5. 查询限制值并以此为 blockDim 下发打印 Kernel if (LaunchPrintKernelWithCurrentResLimit(deviceId, resLimitType, stream) ! 0) { return -1; } return 0; }(); // ...清理与结果汇总 }各步骤的 CANN RUNTIME API 映射关系如下功能分类关键接口在样例中的作用初始化aclInit/aclFinalize初始化 / 去初始化 ACL 运行环境Device 管理aclrtSetDevice指定用于运算的 Device本样例为 device 0Device 管理aclrtResetDeviceForce强制复位 Device回收 Device 上的资源Stream 管理aclrtCreateStream创建 StreamStream 管理aclrtSynchronizeStream阻塞等待 Stream 上任务执行完成Stream 管理aclrtDestroyStreamForce强制销毁 Stream运行时配置aclrtSetDeviceResLimit设置当前进程的 Device 资源限制运行时配置aclrtGetDeviceResLimit获取当前进程的 Device 资源限制4.2 查询限制值并下发 Kernel资源限制查询与 Kernel 下发被封装在LaunchPrintKernelWithCurrentResLimit函数中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而是先调用aclrtGetDeviceResLimit查询当前进程的资源限制值coreDim再将该值作为 blockDim 传给PrintDo。也就是说打印 Kernel 的实际执行规模8 个块完全由此前aclrtSetDeviceResLimit(deviceId, ACL_RT_DEV_RES_CUBE_CORE, 8)设置的配额决定从而验证资源限制真正作用于 Kernel 下发这一核心语义。4.3 设备侧打印 Kernel打印 Kernel 定义于 kernel_print.cpp宿主侧的封装函数PrintDo声明在 kernel_ops.h 中#include kernel_operator.h #include kernel_ops.h extern C __global__ __aicore__ void KernelPrint() { AscendC::printf(Hello World\n); } void PrintDo(uint32_t blockDim, void* stream) { KernelPrintblockDim, nullptr, stream(); }要点说明KernelPrint是__global__ __aicore__修饰的 Device 侧核函数内部调用AscendC::printf输出Hello WorldPrintDo以类 CUDA 的blockDim, nullptr, stream语法将 Kernel 下发到指定 Stream其中 blockDim 即来自资源限制查询值本样例为 8每个块都会执行一次AscendC::printf(Hello World\n)因此 8 个块产生 8 行Hello World这与run.sh中的grep -c校验目标严格对应。4.4 错误处理与资源清理样例采用了精细的错误处理策略主流程使用 utils.h 中定义的CHECK_ERROR宏任一 ACL 调用返回非ACL_SUCCESS即打印错误日志并直接返回-1清理阶段使用UpdateFinalResultOnError辅助函数即便主流程已经失败也会继续依次执行aclrtDestroyStreamForce、aclrtResetDeviceForce、aclFinalize尽力回收 Stream、Device 与 ACL 环境资源并合并最终的返回结果aclrtDestroyStreamForce与aclrtResetDeviceForce均为强制语义接口用于在异常路径下确保资源释放。5. 运行结果解读5.1 预期输出编译运行后样例的标准输出如下示例[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.5.2 输出如何印证资源限制生效输出中的关键信息链可以这样解读[INFO] Device resource limit type 0 set to 8.——type 0即ACL_RT_DEV_RES_CUBE_CORE进程级 Cube Core 资源限制被设置为 8[INFO] Current device resource limit type 0 is 8.——aclrtGetDeviceResLimit查询结果与设置值一致说明限制设置成功且可被查询恰好8 行Hello World——打印 Kernel 以查询到的限制值 8 作为 blockDim 执行每个块打印一行实际执行规模与资源配额吻合[SUCCESS]: Launch kernels under resource limits successfully.——run.sh中grep -c Hello World计数等于 8 的断言通过样例整体成功。6. 延伸资源限制的典型应用场景结合本样例与仓库中的 API 全貌进程级 Device 资源限制在工程实践中通常用于以下场景多进程共享 Device 时的资源隔离多个进程跑在同一物理 Device 上时通过aclrtSetDeviceResLimit为各自进程声明 Cube/Vector Core 配额避免单个进程独占计算核心任务级资源预算控制结合 Stream 级接口aclrtSetStreamResLimit/aclrtGetStreamResLimit可以在更细粒度上约束特定 Stream 的可用资源动态资源感知借助aclrtGetDeviceResLimit/aclrtGetResInCurrentThread在运行时读取当前可用资源据此动态调整 Kernel 的 blockDim 等调度参数实现资源自适应的下发策略本样例正是这一思路的最小演示。需要注意的是资源限制的具体生效范围、枚举取值与硬件能力的对应关系以当前 CANN 版本的接口声明与产品文档为准本仓库中 acl_rt.h 的枚举定义与 device.cpp 的实现可作为接口层事实依据。7. 小结本样例用最小的代码量展示了 CANN Runtime 设置 Device 资源限制 → 查询限制 → 在限制下下发 Kernel 的完整闭环进程级资源限制通过aclrtSetDeviceResLimit设置、aclrtGetDeviceResLimit查询资源类型由aclrtDevResLimitType枚举Cube Core / Vector Core界定样例将查询到的限制值直接用作打印 Kernel 的 blockDim并用 8 行Hello World与脚本断言验证了限制确实作用于 Kernel 下发ACL 层实现最终下沉到rtsSetDeviceResLimit/rtsGetDeviceResLimit并带有 profiling 采样点为上层开发者提供了可靠、可观测的资源配额控制能力。对于希望在多进程、多任务场景下精细管理 AI Core 资源的开发者本样例是一个可以直接复用的最小范式将aclrtSetDeviceResLimit调用置于初始化之后、Kernel 下发之前即可为整个进程划定计算资源边界再配合aclrtGetDeviceResLimit动态适配 Kernel 调度规模。【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考