企业官网建设流程全解析

• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend

导读

本文以 asc-devkit 开源仓库中的 Log 样例（examples/01_simd_cpp_api/04_advanced_api/10_math/log）为主线，完整讲解基于 Ascend C 高阶 APILog实现按元素以 e、2、10 为底做对数运算的算子开发全流程。你将掌握：Log/Log2/Log10接口的声明与底层推导原理、Tiling 参数的设置方法、核函数的编写与调用方式，以及从数据生成、三种编译运行模式到精度校验的完整实操闭环，可直接照搬到其他数学类高阶 API 样例开发中。

样例概述：基于 Log 高阶 API 的对数运算

本样例实现的功能是按元素（element-wise）以 e、2、10 为底做对数运算。它使用 Ascend C 的Log高阶 API 完成核心计算，核函数名为log_custom，固定输入 shape 为x[8, 2048]、输出 shape 为y[8, 2048]，数据类型为float、format 为 ND。样例默认以 e 为底演示（即自然对数），但通过替换接口即可扩展为以 2、10 为底。

接口提示：除本样例使用的Log接口外，Ascend C 还提供了以下对数运算相关高阶 API 接口：

• Lgamma：伽马函数的对数（自然对数）计算，声明位于 include/adv_api/math/lgamma.h，实现lgamma(n) = log(|Γ(n)|)，例如lgamma(1) = 0。

对数运算属于逐元素（Unary）类运算，在 AI Vector Core 上执行，接口定义在 include/adv_api/math/log.h 中，以#pragma begin_pipe(V)标记，说明其计算流水（Pipe）类型为向量计算流水PIPE_V。

支持的产品与 CANN 软件版本

从源码看，Log接口的编译期架构分支支持__NPU_ARCH__为 2201、2002、3510、9201、5102、5162、3003、3113 的架构（见 include/adv_api/math/log.h 中的条件编译），并在 AI Vector Core（ASCEND_IS_AIC）上直接返回，明确限定该接口只用于向量核。

目录结构介绍

├── log │ ├── scripts │ │ └── gen_data.py // 输入数据和真值数据生成脚本 │ ├── CMakeLists.txt // 编译工程文件 │ ├── data_utils.h // 数据读入写出函数 │ ├── log.asc // Ascend C算子实现 & 调用样例 │ └── README.md // 样例说明文档

• log.asc：算子实现与调用样例的主文件，包含 Tiling 数据定义、核函数类、核函数入口与 host 端main；
• gen_data.py：用 NumPy 随机生成输入数据并计算真值（golden）；
• data_utils.h：二进制文件的读写工具，供 host 端读取输入、写出结果；
• CMakeLists.txt：编译工程文件，定义运行模式与 NPU 架构选项。

样例规格

样例实现拆解：从 Tiling 到 Kernel 再到调用

原文档将实现划分为 Kernel 实现、Tiling 实现、调用实现三部分，下面结合 log.asc 的源码逐层深入。

1. Tiling 实现：totalLength 与 tileNum

Tiling（数据切分）通过自定义的LogCustomTilingData结构体承载两个参数：

struct LogCustomTilingData { uint32_t totalLength; // 参与计算的元素总数 uint32_t tileNum; // 每个核内分块（tile）数量 }; void GenerateTilingData(uint8_t* tilingBuf, const uint32_t totalLength, const uint32_t coreNum) { LogCustomTilingData* tiling = reinterpret_cast<LogCustomTilingData*>(tilingBuf); tiling->totalLength = totalLength; tiling->tileNum = coreNum; }

样例按"输入总长度 + 所用核数量"确定 Tiling 参数：totalLength = 8 * 2048 = 16384，核数量USED_CORE_NUM = 8（定义在log.asc的匿名 namespace 中）。核内再按tileNum * BUFFER_NUM分块，其中BUFFER_NUM = 2，即每个核使用双缓冲流水。切分关系在KernelLog::Init中计算：

this->blockLength = tilingData.totalLength / AscendC::GetBlockNum(); // 每个核处理的元素数 this->tileLength = this->blockLength / tileNum / BUFFER_NUM; // 单次搬运/计算的元素数

• GetBlockNum()返回实际使用的核数（与log_custom<<<USED_CORE_NUM, ...>>>的 block dim 对应）；
• GetBlockIdx()用于定位当前核在全局内存中的起始偏移，见xGm.SetGlobalBuffer((__gm__ T*)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength)；
• 输入输出各初始化 2 个缓冲：pipe->InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(T))，实现"搬运-计算-搬出"三级流水。

从源码结构看，Log运算本身不需要额外临时空间：Tiling 侧接口GetLogMaxMinTmpSize/GetLog10MaxMinTmpSize（impl/adv_api/tiling/math/log_tiling.cpp）对float输入直接返回maxValue = 0, minValue = 0。而Log2对half类型需要 256 字节临时缓冲（LOG2_TMP_SIZE = 256），且half输入时最大临时空间为inputSize * sizeof(float)与 256 的较大者；GetLog2TmpBufferFactorSize对half返回maxLiveNodeCount = 2、对float返回 0。这些接口的声明见 include/adv_api/math/log_tiling.h，可作为上层算子 Tiling 设计（如GetLog2MaxMinTmpSize的maxTmpBuffer = maxLiveNodeCount * inputSize * typeSize + extraBuf）的参考。

2. Kernel 实现：Log 接口的调用与底层推导

核函数类KernelLog<T>采用经典的 CopyIn → Compute → CopyOut 三段式处理：

__aicore__ inline void Process() { int32_t loopCount = this->tileNum * BUFFER_NUM; for (int32_t i = 0; i < loopCount; i++) { CopyIn(i); // 通过 DataCopy 将 xGm 分块数据搬入 VECIN 队列 Compute(i); // 调用 Log 接口完成对数计算 CopyOut(i); // 将 VECOUT 队列结果搬回 yGm } }

其中CopyIn使用AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength)将全局内存数据搬入LocalTensor，Compute中即为核心计算语句：

// 使用Log接口计算对数 // 模板参数： // - T: 输入输出数据类型 // - false: 是否复用源操作数 // 参数说明： // - yLocal: 输出Tensor，存储计算结果 // - xLocal: 输入Tensor // - this->tileLength: 计算元素个数 AscendC::Log<T, false>(yLocal, xLocal, this->tileLength);

Log 接口族声明（include/adv_api/math/log.h）：

template <typename T, bool isReuseSource = false> __aicore__ inline void Log(const LocalTensor<T>& dstTensor, const LocalTensor<T>& srcTensor, uint32_t calCount); // 语义：dst = log(src) template <typename T, bool isReuseSource = false> __aicore__ inline void Log2(const LocalTensor<T>& dstTensor, const LocalTensor<T>& srcTensor, uint32_t calCount); // 语义：dst = log2(src) template <typename T, bool isReuseSource = false> __aicore__ inline void Log10(const LocalTensor<T>& dstTensor, const LocalTensor<T>& srcTensor, uint32_t calCount); // 语义：dst = log10(src)

三者均有省略calCount的重载，缺省时以srcTensor.GetSize()作为计算元素个数。模板参数isReuseSource表示是否复用源操作数的输入空间。

底层实现原理（impl/adv_api/detail/math/log/log_common_impl.h）：三个接口并非各自独立实现，而是统一归结到基础向量指令Ln（自然对数）上，再通过一次标量乘法换算底数：

• Log：Logx = Lnx，直接调用Ln<T, false>(dstTensor, srcTensor, MASK_PLACEHOLDER, 1, unaryParams)，并在调用前用SetVectorMask<T, MaskMode::COUNTER>(0, calCount)设置按计数模式的向量掩码（LogImpl，L88-L104）；
• Log2：Log2x = Lnx / Ln2，即Lnx * 1.4426950408889634（Ln2Reciprocal = 1.0/Ln2），实现为先Ln再Muls（Log2Compute，L45-L53）；
• Log10：Log10x = Lnx / Ln10，即Lnx * 0.43429448190325176（Ln10Reciprocal = 1.0/Ln10），实现同样为先Ln再Muls，中间以PipeBarrier<PIPE_V>()保证向量流水同步（Log10Impl，L144-L164）。

对于__NPU_ARCH__为 3510、5102、5162、3003、3113 的架构，Log会切入寄存器级实现 impl/adv_api/detail/math/log/log_3510_impl.h：通过static_assert限定仅支持half与float，以Reg::Log(vreg1, vreg0, mask)逐向量寄存器（vector register）完成计算，LogXImpl对half额外做 b16/b32 的 Cast 转换后再计算。

注意事项：Log相关接口声明在#if defined(__NPU_COMPILER_INTERNAL_PURE_SIMT__)下直接#error，即该接口族不能与--enable-simt编译选项（SIMT 模式）同时使用。此外 include/adv_api/math/kernel_operator_log_intf.h 已被标记为 deprecated，官方提示改用log.h。

3. 调用实现：内核调用符与 ACL 流程

核函数入口通过内核调用符<<<>>>启动（log.asc中__global__ __vector__声明）：

__global__ __vector__ void log_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, LogCustomTilingData tiling) { AscendC::TPipe pipe; MyCustomKernel::KernelLog<float> op; op.Init(x, y, tiling, &pipe); op.Process(); }

host 端main函数则完成标准的 ACL（Ascend CANN Lite）流程：

aclInit(nullptr); aclrtSetDevice(deviceId); aclrtCreateStream(&stream); // 申请 host/device 内存，读取 input_x.bin 并 H2D 拷贝 log_custom<<<USED_CORE_NUM, 0, stream>>>(xDevice, yDevice, tiling); // 8 个核，stream 指定流 aclrtSynchronizeStream(stream); // 同步等待核执行完成 // D2H 拷贝结果，写出 output.bin，与 golden 对比，最后释放资源

Tiling 数据在 host 端直接以结构体实参传入：LogCustomTilingData tiling = {TOTAL_LENGTH, 8}，与核函数签名(x, y, tiling)一一对应。

数据生成与真值校验

gen_data.py 用 NumPy 生成测试数据：

input_shape = [8, 2048] input_x = np.random.uniform(0, 9999, input_shape).astype(dtype) # 正数区间随机采样 golden = np.log(input_x).astype(dtype) # np.log 即自然对数真值

• 输入 x 在(0, 9999)区间随机采样，保证真数恒为正，规避log定义域问题；
• 输入写出为./input/input_x.bin，真值写出为./output/golden.bin，均为二进制文件。

精度对比在log.asc的CompareResult中完成：以EPS = 1e-4为阈值，对每个元素同时比较绝对误差ae = |a - b|与相对误差re = ae / |b|，仅当两者都超过EPS才判为错误元素；只要存在任一错误元素即返回失败。因此最终输出test pass!代表全量 16384 个元素全部通过精度校验。

编译运行：完整步骤与三种运行模式

配置环境变量

请根据当前环境上 CANN 开发套件包的安装方式配置环境变量：

source ${install_path}/cann/set_env.sh

说明：${install_path}为 CANN 包安装目录，未指定安装目录时默认安装至/usr/local/Ascend下。

样例执行

在本样例目录（examples/01_simd_cpp_api/04_advanced_api/10_math/log）下执行如下命令：

mkdir -p build && cd build; # 创建并进入build目录 cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程，默认npu模式 python3 ../scripts/gen_data.py # 生成测试输入数据 ./demo # 执行编译生成的可执行程序，执行样例

CPU 调试与 NPU 仿真模式

使用 CPU 调试或 NPU 仿真模式时，添加-DCMAKE_ASC_RUN_MODE=cpu或-DCMAKE_ASC_RUN_MODE=sim参数即可，示例如：

cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # cpu调试模式 cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # NPU仿真模式

注意：切换编译模式前需清理 cmake 缓存，可在 build 目录下执行rm CMakeCache.txt后重新 cmake。

编译选项说明

这两个选项的默认值定义在样例的 CMakeLists.txt 中，并通过find_package(ASC REQUIRED)引入 Ascend C 编译框架：

set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim") set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture: dav-2201, dav-3510") find_package(ASC REQUIRED) project(kernel_samples LANGUAGES ASC CXX) add_executable(demo log.asc) target_link_libraries(demo PRIVATE tiling_api register platform m dl) target_compile_options(demo PRIVATE $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>)

其中链接库tiling_api、register、platform对应 Ascend C 的 Tiling 接口、算子注册与平台支撑库；编译选项以--npu-arch将目标架构传给 ASC 语言编译器。

执行结果

执行./demo后，若精度对比成功，输出如下：

test pass!

此前还会依次打印ReadFile golden.bin success!与CompareResult golden.bin success!（来自 log.asc 的CompareResult函数）。

延伸阅读

• 若需了解 CANN 开发套件的环境准备与安装细节，可阅读 docs/zh/quick_start.md；
• 完整的Log/Log2/Log10接口声明与注释见 include/adv_api/math/log.h，Tiling 临时空间估算接口见 include/adv_api/math/log_tiling.h；
• 向量核与 3510 系列架构下的底层实现分别见 impl/adv_api/detail/math/log/log_common_impl.h 与 impl/adv_api/detail/math/log/log_3510_impl.h；
• 本样例隶属于高阶 API 数学类算子集，同目录（examples/01_simd_cpp_api/04_advanced_api/10_math）还包含其他数学运算样例，可作为对照参考。

• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend