CANN Runtime 三维 FDTD Stencil 样例深度解析Kernel 属性校验、Device 变量与向量核下发实践【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime导读本文围绕 CANN Runtime 开源仓库中的5_fdtd_stencil样例完整讲解如何在单 Device 上实现一个三维有限差分时域FDTDStencil 更新程序从 Kernel 类型校验确认其为 Vector Core Kernel、通过 Device 变量Symbol注入本轮差分系数到基于aclrtLaunchKernelWithArgsArray下发向量核任务再到 Host 侧逐点比对验证结果。读完本文你将掌握aclrtGetFuncBySymbol、aclrtGetFunctionAttribute、aclrtGetSymbolSize、aclrtMemcpyToSymbol等一组 Kernel 配置与执行接口的组合用法并能独立复现和改造这一确定性网格 Halo 单元 误差校验的数值样例。1. 样例概述用最小三维网格演示完整的 Stencil 计算闭环5_fdtd_stencil是一个面向开发者的教学型样例位于仓库的 example/2_advanced_features/kernel/5_fdtd_stencil 目录。它演示的是一类在电磁场仿真FDTD 方法、流体力学、图像处理等领域广泛存在的计算模式当前点的更新值由自身及相邻点的旧值加权求和得到。该样例在单 Device 上完成了一整条闭环链路确认 FDTD Kernel 是兼容的 Vector Core KernelKernel 类型为向量核将本轮计算所需的中心点系数和相邻点系数写入 Device 变量Symbol对一块带 Halo 单元的确定性三维网格执行一次 Stencil 更新将 Device 上全部 64 个内部点与 Host 参考实现逐点比较。程序的失败判定条件非常明确只要出现以下任一情况即返回非零退出码Kernel 类型不兼容不是向量核Device 变量Symbol的大小与预期不符最大误差超过1e-5任一 Runtime 操作或资源清理失败。该样例的工程文件组成如下文件作用fdtd_stencil.h定义网格维度、系数个数等常量与运行期资源结构体RuntimeResourcesfdtd_stencil.cpp实现 Runtime 初始化、缓冲区准备、Kernel 执行与结果校验、资源释放main.cpp声明 Device 系数变量编写FdtdStencilKernel核函数并组织完整流程run.sh一键完成环境检查、CMake 编译与运行CMakeLists.txt按 SOC 版本选择昇腾指令集架构ASC并链接libacl_rt.so1.1 网格布局6×6×6 外网格与 4×4×4 内部点网格维度的定义集中在 fdtd_stencil.hconstexpr int32_t kDeviceId 0; constexpr uint32_t kInnerDim 4; // 内部有效计算维度 constexpr uint32_t kOuterDim kInnerDim 2; // 外层含 Halo 的维度 constexpr uint32_t kVolumeSize kOuterDim * kOuterDim * kOuterDim; // 216 constexpr uint32_t kCoefficientCount 2; // 中心点系数 相邻点系数这里的带 Halo 单元是 Stencil 计算的标准手法物理上只需要计算 4×4×4 的内部区域但每个内部点在三维的六个方向±x、±y、±z上都要读取邻居值因此在实际存储中把每个维度向外扩 2 个点形成 6×6×6 的网格。这样 Kernel 内部索引center±1、center±strideY、center±strideZ始终落在合法内存范围内无需做边界分支判断。从仓库目录 example/2_advanced_features/kernel 看该样例与6_memory_loaded_vector_addHost 内存加载 Kernel、9_device_symbol_ioDevice 变量读写共同构成了 Kernel 加载、符号与执行方式的完整示例族。2. 编译与运行从克隆仓库到输出校验结果2.1 产品支持范围根据样例文档该样例支持以下产品产品是否支持Ascend 950PR / Ascend 950DT是Atlas A3 训练系列产品 / Atlas A3 推理系列产品是Atlas A2 训练系列产品 / Atlas A2 推理系列产品是2.2 三步运行流程第 1 步进入样例目录。将代码下载到已安装 CANN 软件的环境后切换目录cd ${git_clone_path}/example/2_advanced_features/kernel/5_fdtd_stencil第 2 步设置环境变量。${install_root}替换为 CANN 安装根目录source ${install_root}/set_env.sh source ${git_clone_path}/example/set_sample_env.sh其中example/set_sample_env.sh会完成三件关键事可参考 set_sample_env.sh 源码通过一个基于aclrtGetSocName的小工具探测当前 SOC 版本并导出SOC_VERSION在 CANN 安装目录下定位ascendc_kernel_cmake目录并导出ASCENDC_CMAKE_DIR同时导出ASCEND_INSTALL_PATH与ASCEND_HOME_PATH。若探测失败脚本会打印明确的[ERROR]信息并终止。第 3 步一键编译并运行bash run.shrun.sh 内部会依次执行检查环境变量、打印当前 SOC 版本例如Ascend910_9362、用 CMake 构建、安装产物最后运行./build/main并把输出同时写入output_msg.txt。脚本启用了set -euo pipefail任何一步失败都会立即终止并返回非零状态。2.3 CMake 中的 SOC 到指令集架构映射样例的 CMakeLists.txt 展示了 SOC 版本到昇腾指令集架构的映射关系这是让 Kernel 能编出正确向量指令的关键set(SOC_VERSION $ENV{SOC_VERSION}) string(TOLOWER ${SOC_VERSION} SOC_VERSION_LOWER) if(SOC_VERSION_LOWER MATCHES ^ascend(950|350)) set(CMAKE_ASC_ARCHITECTURES dav-3510) elseif(SOC_VERSION_LOWER MATCHES ^ascend910(b|_93)) set(CMAKE_ASC_ARCHITECTURES dav-2201) else() message(FATAL_ERROR Unsupported SOC_VERSION: ${SOC_VERSION}) endif()dav-3510对应 Ascend 950 / 350 系列即支持列表中的 Ascend 950PR/950DT、Atlas A3 系列dav-2201对应 Ascend 910B / 910_93 系列Atlas A2 系列。随后工程通过find_package(ASC REQUIRED)引入昇腾编译器工具链将 main.cpp 以 ASC 语言参与编译并以libacl_rt.so链接运行期库。可以看到 Kernel 定义__global__ __vector__与 Host 主逻辑fdtd_stencil.cpp在同一可执行文件中完成混合编译这是单文件内嵌 Kernel样例的典型工程形态。3. 样例输出与判定标准在 Ascend 910Ascend910_9362上运行后典型输出如下[INFO]: Current compile soc version is Ascend910_9362 ... [INFO] Start to run the 5_fdtd_stencil sample. [INFO] Kernel type 2 confirms Vector Core compatibility. [INFO] Copied 2 FDTD coefficients to the Device variable. [INFO] Verified 64 interior points; max error is 0.00000000. [INFO] Run the 5_fdtd_stencil sample successfully.日志行含义逐条对应前文流程Kernel type 2 confirms Vector Core compatibility.——aclrtGetFunctionAttribute返回的 Kernel 类型为 2即ACL_KERNEL_TYPE_VECTORAI VECTOR CORE见 acl_rt.hCopied 2 FDTD coefficients to the Device variable.——两个 float 系数0.5 与 1/12已写入 Device 符号Verified 64 interior points; max error is 0.00000000.——4×4×464 个内部点全部通过1e-5容差校验最大误差为 0。日志宏INFO_LOG/ERROR_LOG/CHECK_ERROR来自公共头文件 example/utils.h其中CHECK_ERROR会在任何 ACL 调用返回非ACL_SUCCESS时打印出具体失败调用与错误码并立即返回-1是样例统一的错误传播机制。4. 核心实现逐段拆解4.1 Kernel 定义__gm__全局变量 __vector__向量核main.cpp 中的 Kernel 是本样例的灵魂__gm__ float g_fdtdCoefficients[kCoefficientCount]; extern C __global__ __vector__ void FdtdStencilKernel(__gm__ float* output, __gm__ float* input) { const uint32_t strideY kOuterDim; const uint32_t strideZ kOuterDim * kOuterDim; for (uint32_t z 1; z kInnerDim; z) { for (uint32_t y 1; y kInnerDim; y) { for (uint32_t x 1; x kInnerDim; x) { const uint32_t center z * strideZ y * strideY x; const float neighbors input[center - 1] input[center 1] input[center - strideY] input[center strideY] input[center - strideZ] input[center strideZ]; output[center] g_fdtdCoefficients[0] * input[center] g_fdtdCoefficients[1] * neighbors; } } } #if __NPU_ARCH__ 3510 dcci(reinterpret_cast__gm__ int64_t*(output), cache_line_t::ENTIRE_DATA_CACHE, dcci_dst_t::CACHELINE_OUT); #endif }要点说明__gm__表示全局内存Global Memory指针或变量用于 Device 侧访问 Host 分配的 Device 内存__vector__声明该 Kernel 运行在 AI 向量核Vector Core上这与后续aclrtGetFunctionAttribute查询到的ACL_KERNEL_TYPE_VECTOR相互印证g_fdtdCoefficients是一个由__gm__修饰的全局符号Kernel 直接以g_fdtdCoefficients[0]、g_fdtdCoefficients[1]读取系数——它不经过 Kernel 参数列表而是通过Device 变量机制从 Host 侧注入Cache 一致性处理当__NPU_ARCH__ 3510对应 dav-3510 架构时Kernel 在写完后调用dcci数据缓存失效指令将输出数据从缓存行刷出保证 Host 侧通过 DMA 读回时拿到的是最新值。计算式本身是七点 Stencil 的经典形式output c_center * input c_neighbors * (六个邻居之和)本样例取c_center 0.5、c_neighbors 1/12。4.2 Kernel 配置取句柄、查类型、验兼容ConfigureKernel完成三步操作main.cppCHECK_ERROR(aclrtGetFuncBySymbol(reinterpret_castconst void*(FdtdStencilKernel), funcHandle)); int64_t kernelType -1; CHECK_ERROR(aclrtGetFunctionAttribute(funcHandle, ACL_FUNC_ATTR_KERNEL_TYPE, kernelType)); if (kernelType ! ACL_KERNEL_TYPE_VECTOR) { ERROR_LOG(FDTD stencil requires a Vector Core kernel, but kernel type is %ld., kernelType); return -1; }aclrtGetFuncBySymbol根据 Kernel 符号地址获取运行期函数句柄aclrtFuncHandle该句柄是后续一切 Kernel 属性查询与下发动作的输入声明见 acl_rt.haclrtGetFunctionAttribute查询句柄的指定属性。属性枚举aclrtFuncAttribute定义于 acl_rt.hACL_FUNC_ATTR_KERNEL_TYPE 1表示 Kernel 类型此外还有ACL_FUNC_ATTR_KERNEL_RATIO 2AICore 占比与ACL_FUNC_ATTR_KERNEL_SCHED_MODE 3调度模式ACL_KERNEL_TYPE_VECTOR 2即AI VECTOR CORE对应类型枚举aclrtKernelTypeacl_rt.h中的向量核。同枚举还包括ACL_KERNEL_TYPE_AICORE 0MIX、ACL_KERNEL_TYPE_CUBE 1AI CUBE CORE、ACL_KERNEL_TYPE_MIX 3与ACL_KERNEL_TYPE_AICPU 100。这种先查类型再下发的做法保证了如果误把 Cube 核或 AICPU 类型的函数句柄交给向量核执行路径程序会在启动前以显式报错拦截而不是在 Device 侧产生难以定位的非法指令或错误结果。4.3 系数注入Symbol 大小校验 aclrtMemcpyToSymbolCopyCoefficients演示了 Device 变量Symbol使用的完整姿势main.cppsize_t symbolSize 0; CHECK_ERROR(aclrtGetSymbolSize(g_fdtdCoefficients, symbolSize)); if (symbolSize ! sizeof(coefficients)) { ERROR_LOG(Coefficient symbol size is %zu, expected %zu., symbolSize, sizeof(coefficients)); return -1; } CHECK_ERROR( aclrtMemcpyToSymbol(g_fdtdCoefficients, coefficients, sizeof(coefficients), 0, ACL_MEMCPY_HOST_TO_DEVICE));aclrtGetSymbolSize查询 Device 符号对应的实际字节数。样例刻意先校验symbolSize sizeof(coefficients)即 2×sizeof(float)8 字节防止 Kernel 编译产物与 Host 侧对符号大小的假设不一致aclrtMemcpyToSymbol将 Host 缓冲区coefficients拷贝到符号指向的 Device 内存count为拷贝字节数offset 0表示从符号起始地址写入kind ACL_MEMCPY_HOST_TO_DEVICE指明传输方向。两个 API 的 C 声明均位于 acl_rt.h 与 acl_rt.h。在头文件 acl_rt_api.h 中还提供了模板化的 C 重载可以直接传符号名数组本身而免去手动reinterpret_cast本项目样例选用的是底层 C 接口。系数{0.5F, 1.0F / 12.0F}在运行时才被写入这正体现了 Symbol 机制的典型用途Kernel 二进制与数值参数解耦同一份 Kernel 可以通过改写 Device 变量复用于不同物理场景。4.4 缓冲区准备确定性输入 双向拷贝PrepareBuffersfdtd_stencil.cpp负责申请 Device 内存并生成确定性输入数据constexpr size_t bufferBytes kVolumeSize * sizeof(float); CHECK_ERROR(aclrtMalloc(reinterpret_castvoid**(resources.inputDevice), bufferBytes, ACL_MEM_MALLOC_HUGE_FIRST)); CHECK_ERROR(aclrtMalloc(reinterpret_castvoid**(resources.outputDevice), bufferBytes, ACL_MEM_MALLOC_HUGE_FIRST)); for (uint32_t i 0; i kVolumeSize; i) { input[i] static_castfloat((i * 7U) % 29U) / 29.0F; } CHECK_ERROR(aclrtMemcpy(resources.inputDevice, bufferBytes, input, bufferBytes, ACL_MEMCPY_HOST_TO_DEVICE));aclrtMalloc使用ACL_MEM_MALLOC_HUGE_FIRST分配策略优先大页内存输入输出各 216×4864 字节输入数据由公式(i * 7) % 29 / 29.0生成。它是一个确定性伪随机序列不依赖随机数种子任何一次运行都会产生完全相同的输入从而保证 Host 参考计算与 Device Kernel 计算面对同一份数据这是可复现校验的基础随后aclrtMemcpy以ACL_MEMCPY_HOST_TO_DEVICE方向把整块网格送上 Device。4.5 Kernel 下发与同步ExecuteStencilfdtd_stencil.cpp完成下发、同步与结果回传void* args[] {resources.outputDevice, resources.inputDevice}; CHECK_ERROR(aclrtLaunchKernelWithArgsArray(funcHandle, 1, resources.stream, nullptr, args)); CHECK_ERROR(aclrtSynchronizeStream(resources.stream)); CHECK_ERROR(aclrtMemcpy(output, bufferBytes, resources.outputDevice, bufferBytes, ACL_MEMCPY_DEVICE_TO_HOST));aclrtLaunchKernelWithArgsArray以参数数组形式下发 Kernel声明见 acl_rt.h。args中依次是outputDevice、inputDevice两个 Device 指针的地址与 Kernel 签名的参数顺序一一对应第二个参数1表示 blockDim网格维度为 1即单核执行第三个参数指定执行 Stream第四个参数为额外扩展参数此处传nullptraclrtSynchronizeStream阻塞等待 Stream 上的 Kernel 执行完成。这是异步执行模型下读取结果前必须的同步点同步完成后aclrtMemcpy以ACL_MEMCPY_DEVICE_TO_HOST方向把输出网格拷回 Host 栈上的output数组。需要说明的是样例为单核小网格演示选择了 blockDim1在真实大规模 FDTD 场景中可结合ACL_FUNC_ATTR_KERNEL_RATIO等属性评估算力占比并按 Device 能力调整 blockDim 与网格切分策略。4.6 Host 参考实现与逐点校验VerifyResultfdtd_stencil.cpp用与 Kernel 完全相同的索引公式在 Host 侧重算一遍参考值再逐点比较const float expected coefficients[0] * input[center] coefficients[1] * neighbors; const float error std::fabs(output[center] - expected); maxError error maxError ? error : maxError; ... if (maxError kTolerance) { // kTolerance 1e-5F ERROR_LOG(FDTD result mismatch: max error %.8f exceeds %.8f., maxError, kTolerance); return -1; }校验只遍历z1..4, y1..4, x1..4的内部点共 64 个Halo 层不参与比较。它采用最大绝对误差作为度量只有当全部 64 点的误差都不超过1e-5时才判定通过。由于输入数据是分母为 29 的小数float 运算误差被压得很低实测最大误差通常为0.00000000。4.7 资源释放带失败记录的逆序清理ReleaseResourcesfdtd_stencil.cpp按照后创建先释放的原则逆序清理并且在清理过程中遇到失败不会中断后续清理而是通过RecordCleanupError记录错误并把最终结果置为失败if (resources.outputDevice ! nullptr) { RecordCleanupError(aclrtFree(output), aclrtFree(resources.outputDevice), result); } if (resources.inputDevice ! nullptr) { RecordCleanupError(aclrtFree(input), aclrtFree(resources.inputDevice), result); } if (resources.streamCreated) { RecordCleanupError(aclrtDestroyStreamForce, aclrtDestroyStreamForce(resources.stream), result); } if (resources.deviceSet) { RecordCleanupError(aclrtResetDeviceForce, aclrtResetDeviceForce(kDeviceId), result); } if (resources.initialized) { RecordCleanupError(aclFinalize, aclFinalize(), result); }这一设计与 fdtd_stencil.h 中RuntimeResources的布尔标记配合initialized、deviceSet、streamCreated各自独立记录初始化进度从而保证即使中途失败也只清理真正完成初始化的资源不会对未创建的对象做无效释放。其中aclrtDestroyStreamForce为强制销毁 Stream 接口acl_rt.h适用于销毁尚有未完成任务但确认不再使用的 StreamaclrtResetDeviceForce用于复位 Device 并回收相关资源acl_rt.h其注释明确说明无需在复位前手动销毁 Stream只需调用一次复位即可。5. CANN Runtime API 全景对照将样例涉及的接口按功能域归类如下便于在阅读源码时对照查阅初始化与去初始化aclInit(nullptr)初始化 CANN RuntimeaclFinalize()去初始化回收全局运行期资源。Device 管理aclrtSetDevice(kDeviceId)指定执行 FDTD 计算的 Device样例固定为 0 号设备aclrtResetDeviceForce(kDeviceId)复位 Device 并回收相关资源。Stream 管理aclrtCreateStream(stream)创建 Kernel 执行所用的 StreamaclrtSynchronizeStream(stream)阻塞等待 Kernel 执行完成aclrtDestroyStreamForce(stream)销毁 Stream。内存与数据传输aclrtMalloc/aclrtFree申请/释放输入输出 Device 内存ACL_MEM_MALLOC_HUGE_FIRST大页优先aclrtMemcpyHost 与 Device 间双向传输网格数据H2D / D2HaclrtGetSymbolSize确认 FDTD 系数 Device 变量大小aclrtMemcpyToSymbol将本轮 FDTD 系数写入 Device 变量。Kernel 配置与执行aclrtGetFuncBySymbol根据 Kernel 符号获取函数句柄aclrtGetFunctionAttribute查询 Kernel 类型等属性确认 Vector Core 兼容性aclrtLaunchKernelWithArgsArray以参数数组形式下发 FDTD Stencil 任务。6. 从样例到实践的扩展思考把 Kernel 类型校验沉淀为通用工具ConfigureKernel中的取句柄 → 查类型 → 比对ACL_KERNEL_TYPE_VECTOR三段式可以直接抽象为工具函数在所有向量核样例如本仓库 kernel 目录下的其他样例中复用在启动早期拦截类型不匹配错误。Symbol 机制 vs Kernel 参数本样例刻意展示了第三条参数通道——g_fdtdCoefficientsDevice 变量。当参数需要频繁按轮次更新如 FDTD 每时间步更换系数、或参数是全局常量不希望进入 Kernel 形参列表时aclrtMemcpyToSymbol Kernel 内直接读符号的写法比每次都重新组装参数数组更贴合语义。配套的可参考样例是 9_device_symbol_io后者进一步演示了由 Kernel 写回 Device 变量并在 Host 读回的完整闭环。确定性数据是可复现校验的前提(i * 7) % 29 / 29.0这种无随机数的数据生成方式保证了多次运行输入完全一致配合1e-5的最大误差容差使通过/失败判定完全可复现。在实际数值类样例设计中这是值得沿用的模式。Cache 一致性按架构分条件编译Kernel 末尾的#if __NPU_ARCH__ 3510说明不同指令集架构对 DMA 读回结果的 Cache 策略有差异移植 Kernel 到新架构时需关注此类刷 Cache 指令的适配。参考资料样例主文档example/2_advanced_features/kernel/5_fdtd_stencil/README_en.md中文版见同目录 README.mdKernel 与流程实现main.cpp、fdtd_stencil.cpp常量与资源结构定义fdtd_stencil.h构建与运行脚本run.sh、CMakeLists.txt公共工具头文件example/utils.h环境探测脚本example/set_sample_env.shACL Runtime 接口声明include/external/acl/acl_rt.h属性与类型枚举见 L907-L919【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考