★ FEATURED ARTICLE 人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载导读本文聚焦 asc-devkitCANN 昇腾 AI 处理器算子程序开发语言SIMT-API 中面向half2向量类型一次承载两个 fp16 半精度元素的数学库函数系统讲解h2exp、h2log、h2sin、h2cos、h2tanh、h2sqrt、h2rcp及取整族等 16 个接口的调用方式、特殊值语义与饱和模式行为。读完本文你将掌握在 SIMT 纯线程场景与 SIMDSIMT 混合编程场景中正确使用这批 half2 数学函数并能结合源码理解其先拓宽到 float 计算、再收窄回 half2的实现原理写出数值行为可预期的算子代码。一、half2 数学函数库概览在 SIMT 编程模型中half2是一个包含两个halffp16分量的向量类型一次指令可同时处理两个数据元素是提升半精度计算吞吐的常用载体。asc-devkit 在 SIMT-API 的数学库中为half2提供了一组与单元素h*函数一一对应的批量版本函数名前缀h2索引文档位于 half2_math_functions.md共 16 个接口按功能可划分为四族功能族函数数学含义指数/对数族h2exp、h2exp2、h2exp10、h2log、h2log2、h2log10e^x、2^x、10^x、ln(x)、log₂(x)、log₁₀(x)三角函数族h2sin、h2cos、h2tanhsin、cos、tanh根号/倒数族h2sqrt、h2rsqrt、h2rcp√x、1/√x、1/x取整族h2rint、h2floor、h2ceil、h2trunc就近偶数取整、向下取整、向上取整、截断取整所有接口均声明于头文件simt_api/asc_fp16.h仓库内对应 include/simt_api/asc_fp16.h并以inline内联函数形式实现于 impl/simt_api/asc_fp16_impl.h调用时仅需包含头文件即可无需额外链接数学库。二、产品支持情况与通用语义2.1 支持的产品形态16 个half2数学函数的产品支持情况完全一致Ascend 950PR 950DT 系列产品支持Atlas A3 系列产品不支持Atlas A2 系列产品不支持Atlas 200I/500 A2 推理产品不支持Atlas 推理系列产品 AI Core不支持Atlas 推理系列产品 Vector Core不支持Atlas 训练系列产品不支持。也就是说这批接口当前仅面向 Ascend 950 系列产品开放在其它形态上编译或运行会受限移植代码时需先核对目标 SoC 的系列归属。2.2 全局饱和寄存器与非饱和/饱和模式文档中返回值说明反复出现本接口受全局饱和寄存器影响。这是 SIMT 执行模型中的关键语义当全局饱和寄存器使能饱和模式时任何会越界的计算结果都会被钳制saturate到 fp16 的合法范围或特殊值而不是溢出为 inf/nan。对照特殊值表可以直观看出两种模式的差异例如h2exp在输入为inf时非饱和模式返回inf饱和模式返回ASCRT_MAX_NORMAL_FP16fp16 最大正规数h2log在输入为±0时非饱和模式返回-inf饱和模式返回-ASCRT_MAX_NORMAL_FP16h2sin/h2cos/h2tanh等接口在输入为nan时非饱和模式透传nan饱和模式返回0。编写依赖特殊值判断如用nan检测非法输入的代码时务必确认饱和开关状态否则结果语义会与预期不符。ASCRT_MAX_NORMAL_FP16即 fp16 的最大正规有限值ASCRT_MIN_DENORM_FP16为最小次正规值它们作为饱和钳制边界反复出现在各接口的返回值表中。2.3 通用参数与返回值形态所有 16 个接口的函数原型均为统一形态inline half2 h2xxx(half2 x)参数名输入/输出描述x输入源操作数half2类型两个元素独立参与运算返回值同样是half2其中每个分量对应输入同位置分量的运算结果互不干扰。绝大多数接口的约束说明为无仅h2sqrt、h2rsqrt要求输入元素x 0否则非饱和模式下返回nan详见下文对应小节。三、指数与对数族函数3.1 h2exp / h2exp2 / h2exp10三种底数的幂运算三个接口分别计算输入各元素的 e 次方、2 次方与 10 次方inline half2 h2exp(half2 x); // 获取 e 的 x 元素次方 inline half2 h2exp2(half2 x); // 获取 2 的 x 元素次方 inline half2 h2exp10(half2 x); // 获取 10 的 x 元素次方h2exp的特殊值语义非饱和/饱和模式x 值非饱和模式返回值饱和模式返回值±011nannan0infinfASCRT_MAX_NORMAL_FP16-inf00ASCRT_MAX_NORMAL_FP16infASCRT_MAX_NORMAL_FP16-ASCRT_MAX_NORMAL_FP1600h2exp2与h2exp10的特殊值语义一致x 值非饱和模式返回值饱和模式返回值±011nannan0infinfASCRT_MAX_NORMAL_FP16-inf00x ASCRT_MAX_NORMAL_FP16infASCRT_MAX_NORMAL_FP16三者约束说明均为无。以h2exp的输入输出示例验证数值行为x: (-1.25, 0.5), (2.5, -3.75), (-2, 1.25), (0.75, -0.25) result: (0.286621, 1.64844), (12.1797, 0.0235138), (0.135376, 3.49023), (2.11719, 0.778809)3.2 h2log / h2log2 / h2log10三种底数的对数运算三个接口分别计算以 e、2、10 为底的各元素对数inline half2 h2log(half2 x); // 获取以 e 为底的对数 inline half2 h2log2(half2 x); // 获取以 2 为底的对数 inline half2 h2log10(half2 x); // 获取以 10 为底的对数三者的特殊值语义相同仅h2log在±0与-inf的处理上同样遵循下表h2log2、h2log10表中±0行合并-inf行亦返回 nanx 值非饱和模式返回值饱和模式返回值±0-inf-ASCRT_MAX_NORMAL_FP16nannan0x 0nan0-infnan0infinfASCRT_MAX_NORMAL_FP16约束说明均为无。值得注意的是对数族对非正输入在非饱和模式下返回nan这与后续h2sqrt的约束逻辑一致即非法定义域在非饱和模式下以 nan 透传在饱和模式下钳制为 0。以h2log2的输入输出示例验证x: (-1.25, 0.5), (2.5, -3.75), (-2, 1.25), (0.75, -0.25) result: (nan, -1), (1.32227, nan), (nan, 0.322021), (-0.415039, nan)可见负数元素确实返回nan。四、三角函数族h2sin / h2cos / h2tanhinline half2 h2sin(half2 x); // 各元素的正弦值 inline half2 h2cos(half2 x); // 各元素的余弦值 inline half2 h2tanh(half2 x); // 各元素的双曲正切值h2sin 特殊值x 值非饱和模式返回值饱和模式返回值000-0-0-0nannan0±infnan0h2cos 特殊值x 值非饱和模式返回值饱和模式返回值±011nannan0±infnan0h2tanh 特殊值x 值非饱和模式返回值饱和模式返回值±000nannan0inf11-inf-1-1ASCRT_MAX_NORMAL_FP1611-ASCRT_MAX_NORMAL_FP16-1-1三者约束说明均为无。从语义上看h2tanh天然有界值域 (-1, 1)因此即使输入为最大正规数也收敛到 ±1不受饱和钳制影响这是它区别于指数族的关键特征。以h2tanh的输入输出示例验证x: (-1.25, 0.5), (2.5, -3.75), (-2, 1.25), (0.75, -0.25) result: (-0.848145, 0.462158), (0.986816, -0.999023), (-0.963867, 0.848145), (0.635254, -0.244873)五、根号与倒数族h2sqrt / h2rsqrt / h2rcpinline half2 h2sqrt(half2 x); // 各元素的平方根 inline half2 h2rsqrt(half2 x); // 各元素的平方根倒数 inline half2 h2rcp(half2 x); // 各元素的倒数h2sqrt 特殊值x 值非饱和模式返回值饱和模式返回值000-0-0-0nannan0infinfASCRT_MAX_NORMAL_FP16-infnan0x 0nan0h2sqrt 约束说明本接口支持的输入数据各元素范围为x大于等于 0否则在非饱和模式下返回值为nan。h2rsqrt 特殊值x 值非饱和模式返回值饱和模式返回值0infASCRT_MAX_NORMAL_FP16-0-inf-ASCRT_MAX_NORMAL_FP16nannanASCRT_MAX_NORMAL_FP16x 0nanASCRT_MAX_NORMAL_FP16-infnanASCRT_MAX_NORMAL_FP16inf01.52588e-05h2rsqrt 约束说明本接口支持的输入数据范围为 x 大于等于 0否则非饱和模式下返回值为 nan饱和模式下的返回值参见上表。h2rcp 特殊值x 值非饱和模式返回值饱和模式返回值0infASCRT_MAX_NORMAL_FP16ASCRT_MAX_NORMAL_FP16infASCRT_MAX_NORMAL_FP16-0-inf-ASCRT_MAX_NORMAL_FP16nannan0inf00-inf-0-0h2rcp约束说明为无。注意h2rsqrt在饱和模式下输入inf返回1.52588e-05即 2^-16 量级的最小次正规数的倒数下界这是倒数类指令在饱和钳制下的特有表现与指数族的ASCRT_MAX_NORMAL_FP16钳制策略不同编程时需区分对待。六、取整族h2rint / h2floor / h2ceil / h2truncinline half2 h2rint(half2 x); // 就近取整相等距离取偶数银行家舍入 inline half2 h2floor(half2 x); // 小于或等于各元素的最大整数 inline half2 h2ceil(half2 x); // 大于或等于各元素的最小整数 inline half2 h2trunc(half2 x); // 向零截断的整数h2rint 特殊值注意其在 0.5 边界上的取偶数行为x 值非饱和模式返回值饱和模式返回值±000nannan0infinfASCRT_MAX_NORMAL_FP16-inf-inf-ASCRT_MAX_NORMAL_FP16ASCRT_MAX_NORMAL_FP16ASCRT_MAX_NORMAL_FP16ASCRT_MAX_NORMAL_FP16-ASCRT_MAX_NORMAL_FP16-ASCRT_MAX_NORMAL_FP16-ASCRT_MAX_NORMAL_FP16±0.5001.522-1.5-2-22.522h2floor 特殊值x 值非饱和模式返回值饱和模式返回值±000nannan0infinfASCRT_MAX_NORMAL_FP16-inf-inf-ASCRT_MAX_NORMAL_FP16ASCRT_MIN_DENORM_FP1600h2ceil 特殊值x 值非饱和模式返回值饱和模式返回值±000ASCRT_MIN_DENORM_FP1611nannan0infinfASCRT_MAX_NORMAL_FP16-inf-inf-ASCRT_MAX_NORMAL_FP16h2trunc 特殊值x 值非饱和模式返回值饱和模式返回值±000nannan0infinfASCRT_MAX_NORMAL_FP16-inf-inf-ASCRT_MAX_NORMAL_FP16四者约束说明均为无。h2rint是取整族中语义最特殊的一个——它采用就近取偶bankers rounding例如±0.5 → 0、1.5 → 2、-1.5 → -2、2.5 → 2与h2trunc的向零截断-1.25 → -1、0.75 → 0、h2floor的向下取整-0.25 → -1形成三套不同规则按需选用。七、SIMT 与 SIMDSIMT 两种编程场景的调用范式每个接口的文档都给出了两种编程场景的调用示例范式完全一致可归纳为如下模板以h2tanh为例。场景一纯 SIMT 编程#include simt_api/asc_fp16.h __global__ __launch_bounds__(1024) void kernel_tanh(half2* dst, half2* x, uint32_t total_length) { int idx threadIdx.x blockIdx.x * blockDim.x; if (idx total_length) { return; } dst[idx] h2tanh(x[idx]); }场景二SIMD 与 SIMT 混合编程#include simt_api/asc_fp16.h __simt_vf__ __launch_bounds__(1024) inline void kernel_tanh(__gm__ half2* dst, __gm__ half2* x, uint32_t total_length) { int idx threadIdx.x blockIdx.x * blockDim.x; if (idx total_length) { return; } dst[idx] h2tanh(x[idx]); }两种场景的差异仅在于混合场景的 kernel 函数需以__simt_vf__修饰、并显式使用__gm__标注全局内存指针核心调用dst[idx] h2xxx(x[idx])完全一致。其余 15 个接口只需将h2tanh替换为对应函数名即可套用部分示例如h2exp、h2rcp、h2rint、h2ceil、h2trunc在无边界检查的简化写法中直接计算dst[idx] h2xxx(x[idx])实战中建议保留total_length边界检查以避免越界访问。八、源码级实现原理half2 数学函数的底层实现剖析打开 impl/simt_api/asc_fp16_impl.h 可以看到这 16 个接口并非直接调用单条 fp16 硬件指令而是遵循half2 → float2 拓宽计算 → half2 收窄或逐分量标量复用两类实现策略策略一拓宽到 float2 后调用单精度库函数指数/对数/三角函数族的主体__SIMT_DEVICE_FUNCTIONS_DECL__ inline half2 h2tanh(half2 x) { float2 x_fp32 __cvt_float2__internal_get_round__RoundMode::CAST_RINT(), RoundingSaturation::RS_DISABLE_VALUE(x); float2 res_fp32; res_fp32.x tanhf(x_fp32.x); res_fp32.y tanhf(x_fp32.y); return __cvt_half2__internal_get_round__RoundMode::CAST_RINT(), RoundingSaturation::RS_DISABLE_VALUE(res_fp32); }h2log2、h2log10则利用单精度对数换底公式logf(x) / logf(2.0f)或logf(10.0f)先做浮点拓宽再逐分量换底最后收窄回half2h2sin、h2cos同样在 float2 域调用sinf/cosf。策略二直接复用标量/半精度内建函数逐分量组合根号、倒数、取整族__SIMT_DEVICE_FUNCTIONS_DECL__ inline half2 h2sqrt(half2 x) { return __sqrtf(x); } __SIMT_DEVICE_FUNCTIONS_DECL__ inline half2 h2rcp(half2 x) { half tmp1 static_casthalf(1.0) / x.x; half tmp2 static_casthalf(1.0) / x.y; return {tmp1, tmp2}; } __SIMT_DEVICE_FUNCTIONS_DECL__ inline half2 h2floor(half2 x) { half tmp1 __floorf(x.x); half tmp2 __floorf(x.y); return {tmp1, tmp2}; }其中h2sqrt直接透传__sqrtfh2rcp用1.0 / 分量构造h2rint/h2ceil分别透传__rintf/__ceilfh2trunc则按符号位分支正数走__floorf、负数走__ceilf从而精确实现向零截断语义__SIMT_DEVICE_FUNCTIONS_DECL__ inline half2 h2trunc(half2 x) { half tmp1 x.x; half tmp2 x.y; if (x.x static_casthalf(0)) { tmp1 __floorf(tmp1); } else { tmp1 __ceilf(tmp1); } // x.y 同理 return {tmp1, tmp2}; }h2rsqrt存在两条实现路径文件内出现两次定义属于不同代码分支/条件编译其一是static_casthalf(1.0) / __sqrtf(分量)其二是复用标量版hrsqrt(分量)逐分量组合从源码结构看后者在部分分支下通过封装好的标量半精度函数串联实现。这两类策略共同保证了用户视角下调用是透明的half2 - half2而底层计算精度与中间舍入行为由 float 中间变量和CAST_RINT舍入模式决定。此外include/simt_api/asc_fp16.h 中所有接口均以__SIMT_DEVICE_FUNCTIONS_DECL__ inline前缀声明如inline half2 h2tanh(half2 x);表明它们属于设备侧内联函数编译期直接内联到 kernel 中不产生额外函数调用开销。九、测试与验证如何在仓库中确认接口行为仓库的 SIMT API 测试目录 tests/api/simt_api 提供了可直接参考的验证用例其中 tests/api/simt_api/presmoke/half_math_functions.asc 覆盖了half类型数学函数的冒烟测试bfloat16_math_functions.asc则对应 bf16 变体tests/api/simt_api/ascendc_case_ascend950pr_9599_simt/test_operator_simt_bf16.cpp 是针对 Ascend 950PR/9599 的 SIMT 算子用例可作为在目标 SoC 上跑通h2*接口的最小工程模板。建议实践路径编写包含边界值±0、nan、±inf、ASCRT_MAX_NORMAL_FP16与常规值混合的half2输入分别在饱和/非饱和两种模式下运行对照本文各表核对返回值将结果与文档给出的输入输出示例比对示例中nan元素即负数输入在非饱和模式下的表现。十、实战要点总结头文件所有接口统一包含simt_api/asc_fp16.h即可声明见 include/simt_api/asc_fp16.h实现见 impl/simt_api/asc_fp16_impl.h平台约束16 个接口当前仅支持 Ascend 950PR 950DT 系列其余 Atlas 系列均不支持饱和语义nan在饱和模式下通常钳制为 0inf钳制为ASCRT_MAX_NORMAL_FP16h2rsqrt例外地钳为1.52588e-05±0的 log 钳为-ASCRT_MAX_NORMAL_FP16定义域约束h2sqrt/h2rsqrt要求x 0对数族对x 0返回nan非饱和模式取整规则差异h2rint就近取偶、h2floor向下、h2ceil向上、h2trunc向零截断边界行为各不相同性能路径底层多数接口先拓宽 float2 再计算精度与中间舍入由CAST_RINT舍入模式控制取整/倒数类直接逐分量组合内建函数均为编译期内联无调用开销。将以上要点与文档中的特殊值表结合使用即可在 Ascend 950 系列产品的 SIMT 与混合编程场景中安全、准确地完成 half2 半精度数学运算的开发与调试。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐asc-devkit 昇腾 SIMT 编程 half2 类型比较函数完全指南41 个 API 语义、特殊值与实战示例asc devkit 昇腾 SIMT 编程 half2 类型比较函数完全指南41 个 API 语义、特殊值与实战示例 本文围绕 asc devkit 仓库中人工智能深度学习算子库CANNAscendasc-devkit SIMT API 之 bfloat16 数学函数族类型体系、四类函数与 SIMT/SIMD 混合编程实践asc devkit SIMT API 之 bfloat16 数学函数族类型体系、四类函数与 SIMT/SIMD 混合编程实践 本篇指南系统梳理 asc de人工智能深度学习算子库CANNAscend昇腾SIMT Warp函数编程指南Vote、Shfl、Reduce与Lane ID 全解析昇腾SIMT Warp函数编程指南Vote、Shfl、Reduce与Lane ID 全解析 Warp 是 SIMT 线程执行与调度的基本单位Warp 函数则人工智能深度学习算子库CANNAscend上一篇Qwen3VL文本编码器与Boogu-Image协同工作提升AI绘图精度的秘诀下一篇Oink用户指南如何分析日志并找到内存消耗最高的控制器动作创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
阅读完成 · 觉得有帮助？
