asc-devkit SIMT-API 详解：half 类型数学库函数（htanh/hexp/hlog 等 16 个接口）全面解析

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

本文聚焦 CANN asc-devkit 开源仓库中 SIMT-API 面向half（FP16）类型提供的数学库函数，覆盖htanh、hexp、hexp2、hexp10、hlog、hlog2、hlog10、hcos、hsin、hsqrt、hrsqrt、hrcp、hrint、hfloor、hceil、htrunc共 16 个接口。文章以 half_math_functions.md 文档为主体，结合 asc_fp16.h 接口声明、asc_fp16_impl.h 底层实现与 math_half.asc 测试用例，逐一讲解每个函数的原型、参数、返回值（含饱和模式特殊值语义）、约束、调用示例与实现原理，帮助读者在昇腾算子开发中准确选用与调用这些 FP16 数学函数。

一、函数总览与产品支持情况

half 类型数学库函数是 SIMT-API 中面向单元素 half 标量（以及配套的 half2 向量版本）提供的内联数学函数集合，完整清单与逐函数文档索引如下：

需要特别注意的是，上述16 个接口的产品支持情况完全一致，均仅在以下产品上受支持：

• Ascend 950PR & 950DT 系列产品：支持
• Atlas A3 系列产品：不支持
• Atlas A2 系列产品：不支持
• Atlas 200I/500 A2 推理产品：不支持
• Atlas 推理系列产品 AI Core：不支持
• Atlas 推理系列产品 Vector Core：不支持
• Atlas 训练系列产品：不支持

从源码结构看，这一"950 系列专用"特性与 half 数学函数族的定位一致：它们直接映射到昇腾 950 系列提供的标量数学指令/库函数，因此在选用前务必确认目标部署 SoC 属于 Ascend 950 系列。

二、公共约定：头文件、饱和模式与特殊值语义

2.1 头文件路径

所有 half 类型数学库函数均声明于同一个头文件：

#include "simt_api/asc_fp16.h"

该头文件位于 include/simt_api/asc_fp16.h，其中第 44～104 行依次声明了hcos、hsin、htanh、hexp、hexp2、hexp10、hlog、hlog2、hlog10、hsqrt、hrsqrt、hrcp、hfloor、hrint、hceil、htrunc的函数原型。这些接口均为inline内联函数，参数与返回值类型为half，便于在 Kernel 内部逐元素调用。

2.2 全局饱和寄存器与饱和模式

文档中每个接口的"返回值说明"都强调：本接口受全局饱和寄存器的影响。这意味着计算结果在写回前会依据饱和开关做一次收敛处理。各函数文档给出的特殊值表中，统一存在两列语义：

• 非饱和模式返回值：遵循 IEEE 754 语义，溢出为 ±inf、非法输入为 nan，符号位按数学规则保留；
• 饱和模式返回值：将 inf/nan 等"非有限值"收敛到 half 能表示的最大有限正常数，即ASCRT_MAX_NORMAL_FP16（FP16 最大正规数，约 65504），nan 则收敛为 0。

典型收敛规则归纳如下：

例如hexp在非饱和模式下inf输入返回inf，饱和模式下返回ASCRT_MAX_NORMAL_FP16；hlog在饱和模式下-inf、x<0均返回 0。hrsqrt是例外：其 nan 输入在饱和模式下返回ASCRT_MAX_NORMAL_FP16而非 0（详见第四节），使用时需注意这一差异。

2.3 统一的输入输出示例

各函数文档采用同一组输入数据x: -1.25, 0.5, 2.5, -3.75给出实测输出（half 精度下的近似值），便于横向对比函数行为，下文各节将逐一列出。

三、指数与对数族：hexp / hexp2 / hexp10 / hlog / hlog2 / hlog10

3.1 hexp：e 的 x 次方

• 函数原型：inline half hexp(half x)
• 参数：x（输入，源操作数）
• 功能：获取 e 的 x 次方，对应源码实现为__expf(x)（见 asc_fp16_impl.h）。

特殊值：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: 0.286621, 1.64844, 12.1797, 0.0235138

3.2 hexp2：2 的 x 次方

• 函数原型：inline half hexp2(half x)
• 参数：x（输入，源操作数）
• 功能：获取 2 的 x 次方，实现见 asc_fp16_impl.h。

特殊值：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: 0.42041, 1.41406, 5.65625, 0.0743408

3.3 hexp10：10 的 x 次方

• 函数原型：inline half hexp10(half x)
• 参数：x（输入，源操作数）
• 功能：获取 10 的 x 次方，实现见 asc_fp16_impl.h。

特殊值：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: 0.0562439, 3.16211, 316.25, 0.00017786

3.4 hlog：以 e 为底的对数

• 函数原型：inline half hlog(half x)
• 参数：x（输入，源操作数）
• 功能：获取以 e 为底、x 的对数，实现为__logf(x)（见 asc_fp16_impl.h）。

特殊值：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: nan, -0.693359, 0.916504, nan

3.5 hlog2：以 2 为底的对数

• 函数原型：inline half hlog2(half x)
• 参数：x（输入，源操作数）
• 功能：获取以 2 为底、x 的对数。实现上先将 half 升精度为 float，计算logf(x) / logf(2.0f)后再转回 half（见 asc_fp16_impl.h），通过中间浮点计算换取精度。

特殊值：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: nan, -1, 1.32227, nan

3.6 hlog10：以 10 为底的对数

• 函数原型：inline half hlog10(half x)
• 参数：x（输入，源操作数）
• 功能：获取以 10 为底、x 的对数，实现见 asc_fp16_impl.h。

特殊值：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: nan, -0.301025, 0.397949, nan

四、根号与倒数族：hsqrt / hrsqrt / hrcp

4.1 hsqrt：平方根

• 函数原型：inline half hsqrt(half x)
• 参数：x（输入，源操作数）
• 功能：获取输入数据 x 的平方根，实现为__sqrtf(x)（见 asc_fp16_impl.h）。
• 约束：输入数据范围要求x ≥ 0，否则非饱和模式返回 nan、饱和模式返回 0。

特殊值：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: nan, 0.707031, 1.58105, nan

4.2 hrsqrt：平方根的倒数

• 函数原型：inline half hrsqrt(half x)
• 参数：x（输入，源操作数）
• 功能：获取输入数据 x 的平方根的倒数，实现为(half)1.0 / hsqrt(x)（见 asc_fp16_impl.h）。
• 约束：输入数据范围要求x ≥ 0，否则返回 nan。

特殊值（注意饱和模式下 nan 类输入收敛为 ASCRT_MAX_NORMAL_FP16，与多数接口"nan→0"不同）：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: nan, 1.41406, 0.632324, nan

4.3 hrcp：倒数

• 函数原型：inline half hrcp(half x)
• 参数：x（输入，源操作数）
• 功能：获取输入数据 x 的倒数，实现为static_cast<half>(1.0) / x（见 asc_fp16_impl.h），即标准的除法求倒。

特殊值：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: -0.799805, 2, 0.399902, -0.266602

五、三角与双曲族：hsin / hcos / htanh

5.1 hsin：正弦

• 函数原型：inline half hsin(half x)
• 参数：x（输入，源操作数）
• 功能：获取输入数据的三角函数正弦值，实现见 asc_fp16_impl.h。

特殊值（注意 ±0 的符号位被保留）：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: -0.949219, 0.479492, 0.598633, 0.571777

5.2 hcos：余弦

• 函数原型：inline half hcos(half x)
• 参数：x（输入，源操作数）
• 功能：获取输入数据的三角函数余弦值，实现见 asc_fp16_impl.h。

特殊值：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: 0.31543, 0.877441, -0.80127, -0.820801

5.3 htanh：双曲正切

• 函数原型：inline half htanh(half x)
• 参数：x（输入，源操作数）
• 功能：获取输入数据的双曲正切值。实现先将 half 升为 float，调用tanhf(tmp)后转回 half（见 asc_fp16_impl.h）。

特殊值：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: -0.848145, 0.462158, 0.986816, -0.999023

六、舍入族：hrint / hfloor / hceil / htrunc

6.1 hrint：舍入到最接近的整数（ties-to-even）

• 函数原型：inline half hrint(half x)
• 参数：x（输入，源操作数）
• 功能：获取与输入数据最接近的整数；若存在两个同样接近的整数，则取其中的偶数（银行家舍入，round-to-nearest-even），实现为__rintf(x)（见 asc_fp16_impl.h）。

特殊值：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: -1, 0, 2, -4

上例中 0.5 舍入为 0、2.5 舍入为 2，正是 ties-to-even 的体现。

6.2 hfloor：向下取整

• 函数原型：inline half hfloor(half x)
• 参数：x（输入，源操作数）
• 功能：获取小于或等于输入数据的最大整数值（floor），实现为__floorf(x)（见 asc_fp16_impl.h）。

特殊值：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: -2, 0, 2, -4

6.3 hceil：向上取整

• 函数原型：inline half hceil(half x)
• 参数：x（输入，源操作数）
• 功能：获取大于或等于输入数据的最小整数值（ceil），实现为__ceilf(x)（见 asc_fp16_impl.h）。

特殊值：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: -1, 1, 3, -3

6.4 htrunc：向零截断

• 函数原型：inline half htrunc(half x)
• 参数：x（输入，源操作数）
• 功能：获取对输入数据浮点数截断后的整数（丢弃小数部分，向零收敛）。实现上对正数取__floorf、对负数取__ceilf（见 asc_fp16_impl.h），等价于标准 trunc 语义。

特殊值：

示例：

x: -1.25, 0.5, 2.5, -3.75 result: -1, 0, 2, -3

对比同一输入可清晰看出四种舍入语义的差异：hrint得-1, 0, 2, -4，hfloor得-2, 0, 2, -4，hceil得-1, 1, 3, -3，htrunc得-1, 0, 2, -3。

七、两种编程场景下的调用方式

每个函数文档均给出两套调用示例，分别对应两种 Kernel 编写场景。

7.1 SIMT 编程场景（标量线程模型）

在__global__Kernel 中按线程索引逐元素调用：

#include "simt_api/asc_fp16.h" __global__ __launch_bounds__(1024) void KernelTanh(half* dst, half* x) { int idx = threadIdx.x + blockIdx.x * blockDim.x; dst[idx] = htanh(x[idx]); }

以hsqrt为例，还可以加入边界保护，避免越界访问：

#include "simt_api/asc_fp16.h" __global__ __launch_bounds__(1024) void kernel_sqrt(half* dst, half* x, uint32_t total_length) { uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx >= total_length) { return; } dst[idx] = hsqrt(x[idx]); }

7.2 SIMD 与 SIMT 混合编程场景（向量函数）

在混合编程场景中，数学函数封装在__simt_vf__（向量函数）中，并通过asc_vf_call以dim3(1024)的网格配置发起调用。hsqrt、hfloor、hsin等函数文档给出了完整范例：

#include "simt_api/asc_fp16.h" __simt_vf__ __launch_bounds__(1024) inline void kernel_sqrt(__gm__ half* dst, __gm__ half* x, uint32_t total_length) { uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx >= total_length) { return; } dst[idx] = hsqrt(x[idx]); } __global__ __vector__ void run_sqrt(__gm__ half* dst, __gm__ half* x, uint32_t total_length) { asc_vf_call<kernel_sqrt>(dim3(1024), dst, x, total_length); }

要点说明：

• __simt_vf__标记的函数运行于 SIMT 向量线程，内部仍可使用threadIdx.x、blockIdx.x、blockDim.x计算全局索引；
• 指针需使用__gm__修饰以标明全局内存地址空间；
• asc_vf_call<kernel_sqrt>(dim3(1024), ...)以 1024 线程/块的配置实例化并调度该向量函数；
• 不含边界保护的简化形式（如KernelTanh、KernelRsqrt、KernelCeil等）在数据长度恰为线程数整数倍时也可直接使用。

八、源码实现剖析

16 个接口的声明与实现分别位于接口层与实现层两个文件，形成"声明-实现分离"的层次：

接口声明（include/simt_api/asc_fp16.h）：统一使用__SIMT_DEVICE_FUNCTIONS_DECL__ inline half hxxx(half x)形式声明，随头文件一并提供half2向量版本（如h2tanh、h2exp、h2log2、h2sqrt、h2rcp、h2floor、h2rint、h2ceil、h2trunc等，声明位置紧邻标量版本）。

实现层（impl/simt_api/asc_fp16_impl.h），可归纳为三类实现策略：

1. 直接映射硬件内建函数：hsqrt → __sqrtf、hfloor → __floorf、hrint → __rintf、hceil → __ceilf、hlog → __logf、hexp → __expf，这些函数以一行内联实现直接映射昇腾硬件标量指令，开销最低；
2. 复合运算实现：hrcp用static_cast<half>(1.0) / x实现；hrsqrt用(half)1.0 / hsqrt(x)组合实现；htrunc通过"正数 floor、负数 ceil"的条件分支实现；
3. 升精度计算保证精度：htanh（half→float→tanhf→half）、hlog2（logf(x)/logf(2.0f)）等函数先将 half 升为 float 参与标准库计算再转回 half，以中间精度的代价换取更精确的结果（见 asc_fp16_impl.h）。

从源码结构可以推断，half2向量版本大多为标量版本的逐通道展开（如h2rcp分别计算x.x与x.y的倒数后组合返回），便于在需要同时处理两个 half 元素时减少线程开销。

九、测试验证

仓库提供了专门的 SIMT half 数学函数测试，可对照验证文档行为：

• 单元级测试：tests/api/simt_api/math_half/math_half.asc 中math_customKernel 在边界保护下通过hlog2、hlog10、hcos、hsin对输入逐元素计算，随后借助aclrtMallocHost/aclrtMalloc/aclrtMemcpy完成 host/device 数据传输与核函数调度（见 math_half.asc）；
• 预冒烟测试（presmoke）：tests/api/simt_api/presmoke/half_math_functions.asc 将 half 数学函数族纳入 CI 预冒烟用例，覆盖函数族的可编译性与基础运行正确性。

这些用例与本文档中的特殊值表、输入输出示例共同构成"文档—源码—测试"三层证据链，读者可按 tests/api/simt_api 目录下的组织方式扩展自己的验证用例。

十、使用注意事项小结

1. 产品兼容性：16 个接口目前仅 Ascend 950PR&950DT 系列支持，跨 SoC 移植前务必核对目标产品清单；
2. 饱和模式语义：默认受全局饱和寄存器影响，nan→0、inf→ASCRT_MAX_NORMAL_FP16；其中hrsqrt的 nan 在饱和模式下收敛为ASCRT_MAX_NORMAL_FP16而非 0，属于特殊情形；
3. 输入域约束：hsqrt、hrsqrt要求 x ≥ 0；hlog/hlog2/hlog10对负数输入返回 nan（非饱和）/0（饱和）；指数族在结果溢出 half 范围时返回 inf 或饱和值；
4. 舍入语义差异：hrint为 ties-to-even、hfloor向下、hceil向上、htrunc向零，四者在负数与 .5 边界处结果不同，选择时需匹配业务语义；
5. 编程场景：SIMT 场景直接在线程函数中调用；SIMD/SIMT 混合场景需配合__simt_vf__、__gm__与asc_vf_call使用，并建议在 Kernel 内加入idx >= total_length边界保护。

通过本文档与 include/simt_api/asc_fp16.h、impl/simt_api/asc_fp16_impl.h 源码对照阅读，开发者即可在 Ascend 950 系列上熟练使用这 16 个 half 类型数学函数，实现高效、可预期的 FP16 标量数学计算。

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