企业官网建设流程全解析

CANN Runtime 三维 FDTD Stencil 样例深度解析：Kernel 属性校验、Device 变量与向量核下发实践

【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime

导读

本文围绕 CANN Runtime 开源仓库中的5_fdtd_stencil样例，完整讲解如何在单 Device 上实现一个三维有限差分时域（FDTD）Stencil 更新程序：从 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 上完成了一整条闭环链路：

1. 确认 FDTD Kernel 是兼容的 Vector Core Kernel（Kernel 类型为向量核）；
2. 将本轮计算所需的中心点系数和相邻点系数写入 Device 变量（Symbol）；
3. 对一块带 Halo 单元的确定性三维网格执行一次 Stencil 更新；
4. 将 Device 上全部 64 个内部点与 Host 参考实现逐点比较。

程序的失败判定条件非常明确，只要出现以下任一情况即返回非零退出码：

• Kernel 类型不兼容（不是向量核）；
• Device 变量（Symbol）的大小与预期不符；
• 最大误差超过1e-5；
• 任一 Runtime 操作或资源清理失败。

该样例的工程文件组成如下：

1.1 网格布局：6×6×6 外网格与 4×4×4 内部点

网格维度的定义集中在 fdtd_stencil.h：

constexpr 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_add（Host 内存加载 Kernel）、9_device_symbol_io（Device 变量读写）共同构成了 Kernel 加载、符号与执行方式的完整示例族。

2. 编译与运行：从克隆仓库到输出校验结果

2.1 产品支持范围

根据样例文档，该样例支持以下产品：

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.sh

run.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 910（Ascend910_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_VECTOR（AI VECTOR CORE，见 acl_rt.h）；
• Copied 2 FDTD coefficients to the Device variable.——两个 float 系数（0.5 与 1/12）已写入 Device 符号；
• Verified 64 interior points; max error is 0.00000000.——4×4×4=64 个内部点全部通过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.cpp）：

CHECK_ERROR(aclrtGetFuncBySymbol(reinterpret_cast<const 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.h）；
• aclrtGetFunctionAttribute：查询句柄的指定属性。属性枚举aclrtFuncAttribute定义于 acl_rt.h：ACL_FUNC_ATTR_KERNEL_TYPE = 1表示 Kernel 类型，此外还有ACL_FUNC_ATTR_KERNEL_RATIO = 2（AICore 占比）与ACL_FUNC_ATTR_KERNEL_SCHED_MODE = 3（调度模式）；
• ACL_KERNEL_TYPE_VECTOR = 2即"AI VECTOR CORE"，对应类型枚举aclrtKernelType（acl_rt.h）中的向量核。同枚举还包括ACL_KERNEL_TYPE_AICORE = 0（MIX）、ACL_KERNEL_TYPE_CUBE = 1（AI CUBE CORE）、ACL_KERNEL_TYPE_MIX = 3与ACL_KERNEL_TYPE_AICPU = 100。

这种"先查类型再下发"的做法保证了：如果误把 Cube 核或 AICPU 类型的函数句柄交给向量核执行路径，程序会在启动前以显式报错拦截，而不是在 Device 侧产生难以定位的非法指令或错误结果。

4.3 系数注入：Symbol 大小校验 +aclrtMemcpyToSymbol

CopyCoefficients演示了 Device 变量（Symbol）使用的完整姿势（main.cpp）：

size_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 缓冲区准备：确定性输入 + 双向拷贝

PrepareBuffers（fdtd_stencil.cpp）负责申请 Device 内存并生成确定性输入数据：

constexpr size_t bufferBytes = kVolumeSize * sizeof(float); CHECK_ERROR(aclrtMalloc(reinterpret_cast<void**>(&resources.inputDevice), bufferBytes, ACL_MEM_MALLOC_HUGE_FIRST)); CHECK_ERROR(aclrtMalloc(reinterpret_cast<void**>(&resources.outputDevice), bufferBytes, ACL_MEM_MALLOC_HUGE_FIRST)); for (uint32_t i = 0; i < kVolumeSize; ++i) { input[i] = static_cast<float>((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×4=864 字节；
• 输入数据由公式(i * 7) % 29 / 29.0生成。它是一个确定性伪随机序列：不依赖随机数种子，任何一次运行都会产生完全相同的输入，从而保证 Host 参考计算与 Device Kernel 计算面对同一份数据，这是"可复现校验"的基础；
• 随后aclrtMemcpy以ACL_MEMCPY_HOST_TO_DEVICE方向把整块网格送上 Device。

4.5 Kernel 下发与同步

ExecuteStencil（fdtd_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；第四个参数为额外扩展参数（此处传nullptr）；
• aclrtSynchronizeStream：阻塞等待 Stream 上的 Kernel 执行完成。这是异步执行模型下读取结果前必须的同步点；
• 同步完成后，aclrtMemcpy以ACL_MEMCPY_DEVICE_TO_HOST方向把输出网格拷回 Host 栈上的output数组。

需要说明的是，样例为单核小网格演示选择了 blockDim=1；在真实大规模 FDTD 场景中，可结合ACL_FUNC_ATTR_KERNEL_RATIO等属性评估算力占比，并按 Device 能力调整 blockDim 与网格切分策略。

4.6 Host 参考实现与逐点校验

VerifyResult（fdtd_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; }

校验只遍历z=1..4, y=1..4, x=1..4的内部点（共 64 个），Halo 层不参与比较。它采用最大绝对误差作为度量：只有当全部 64 点的误差都不超过1e-5时才判定通过。由于输入数据是分母为 29 的小数，float 运算误差被压得很低，实测最大误差通常为0.00000000。

4.7 资源释放：带失败记录的逆序清理

ReleaseResources（fdtd_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），适用于销毁尚有未完成任务但确认不再使用的 Stream；
• aclrtResetDeviceForce用于复位 Device 并回收相关资源（acl_rt.h），其注释明确说明无需在复位前手动销毁 Stream，只需调用一次复位即可。

5. CANN Runtime API 全景对照

将样例涉及的接口按功能域归类如下，便于在阅读源码时对照查阅：

初始化与去初始化

• aclInit(nullptr)：初始化 CANN Runtime；
• aclFinalize()：去初始化，回收全局运行期资源。

Device 管理

• aclrtSetDevice(kDeviceId)：指定执行 FDTD 计算的 Device（样例固定为 0 号设备）；
• aclrtResetDeviceForce(kDeviceId)：复位 Device 并回收相关资源。

Stream 管理

• aclrtCreateStream(&stream)：创建 Kernel 执行所用的 Stream；
• aclrtSynchronizeStream(stream)：阻塞等待 Kernel 执行完成；
• aclrtDestroyStreamForce(stream)：销毁 Stream。

内存与数据传输

• aclrtMalloc/aclrtFree：申请/释放输入输出 Device 内存（ACL_MEM_MALLOC_HUGE_FIRST大页优先）；
• aclrtMemcpy：Host 与 Device 间双向传输网格数据（H2D / D2H）；
• aclrtGetSymbolSize：确认 FDTD 系数 Device 变量大小；
• aclrtMemcpyToSymbol：将本轮 FDTD 系数写入 Device 变量。

Kernel 配置与执行

• aclrtGetFuncBySymbol：根据 Kernel 符号获取函数句柄；
• aclrtGetFunctionAttribute：查询 Kernel 类型等属性，确认 Vector Core 兼容性；
• aclrtLaunchKernelWithArgsArray：以参数数组形式下发 FDTD Stencil 任务。

6. 从样例到实践的扩展思考

1. 把 Kernel 类型校验沉淀为通用工具：ConfigureKernel中的"取句柄 → 查类型 → 比对ACL_KERNEL_TYPE_VECTOR"三段式可以直接抽象为工具函数，在所有向量核样例（如本仓库 kernel 目录下的其他样例）中复用，在启动早期拦截类型不匹配错误。

2. Symbol 机制 vs Kernel 参数：本样例刻意展示了第三条参数通道——g_fdtdCoefficientsDevice 变量。当参数需要频繁按轮次更新（如 FDTD 每时间步更换系数）、或参数是全局常量不希望进入 Kernel 形参列表时，aclrtMemcpyToSymbol+ Kernel 内直接读符号的写法比每次都重新组装参数数组更贴合语义。配套的可参考样例是 9_device_symbol_io，后者进一步演示了由 Kernel 写回 Device 变量并在 Host 读回的完整闭环。

3. 确定性数据是可复现校验的前提：(i * 7) % 29 / 29.0这种无随机数的数据生成方式保证了多次运行输入完全一致，配合1e-5的最大误差容差，使"通过/失败"判定完全可复现。在实际数值类样例设计中，这是值得沿用的模式。

4. Cache 一致性按架构分条件编译：Kernel 末尾的#if __NPU_ARCH__ == 3510说明不同指令集架构对 DMA 读回结果的 Cache 策略有差异，移植 Kernel 到新架构时需关注此类刷 Cache 指令的适配。

参考资料

• 样例主文档：example/2_advanced_features/kernel/5_fdtd_stencil/README_en.md（中文版见同目录 README.md）
• Kernel 与流程实现：main.cpp、fdtd_stencil.cpp
• 常量与资源结构定义：fdtd_stencil.h
• 构建与运行脚本：run.sh、CMakeLists.txt
• 公共工具头文件：example/utils.h
• 环境探测脚本：example/set_sample_env.sh
• ACL Runtime 接口声明：include/external/acl/acl_rt.h（属性与类型枚举见 L907-L919）

【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime