数据类型与精度保护【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills触发条件Agent 处理数据类型转换或精度保护时查阅本文档1. 910_95 支持的数据类型速查1.1 完全支持的数据类型数据类型Triton 名称字节数512B 对齐元素数典型用途int8tl.int81512量化推理、索引int16tl.int162256中等范围整数int32tl.int324128索引、计数器int64tl.int64864大偏移量Vector ADD/CMP 退化为 scalarfp16tl.float162256矩阵乘法训练、高精度推理bf16tl.bfloat162256NPU 推荐推理精度、矩阵乘法fp32tl.float324128累加器、归约、高精度计算booltl.int10.125-条件判断内部转为 int81.2 部分支持的数据类型数据类型Triton 名称支持范围910_95 额外支持fp8e4nvtl.float8e4nv类型转换tl.dot_scaled的 FP8 输入fp8e4b15tl.float8e4b15类型转换tl.dot_scaled的 FP8 输入fp8e5tl.float8e5类型转换tl.dot_scaled的 FP8 输入fp8e4b8tl.float8e4b8类型转换tl.dot_scaled的 FP8 输入fp8e5b16tl.float8e5b16类型转换tl.dot_scaled的 FP8 输入uint8tl.uint8Block Pointer 场景不支持-1.3 不支持的数据类型数据类型Triton 名称替代方案fp64tl.float64使用 fp32uint16tl.uint16使用 int16uint32tl.uint32使用 int32uint64tl.uint64使用 int64源码参考core.py 中 dtype 定义compiler.py 中 FP8 支持列表2. 归约操作升精度规则2.1 核心规则所有归约必须在 FP32 下进行NPU 上归约操作tl.sum、tl.max、tl.min、tl.argmax、tl.argmin的精度行为与 GPU 不同必须手动确保在 FP32 精度下执行归约。输入类型GPU 行为NPU (910_95) 行为正确做法fp16自动提升为 fp32 归约tl.sum直接 fp16 归约tl.max/tl.min自动提升为 fp32tl.sum需手动.to(tl.float32)后归约bf16直接 bf16 归约自动提升为 fp32 归约无需额外处理编译器自动提升int8提升为 int32 归约直接 int8 归约手动.to(tl.int32)后归约int16提升为 int32 归约直接 int16 归约手动.to(tl.int32)后归约2.2 标准写法triton.jit def layernorm_kernel(X, Out, Mean, Rstd, M, N, eps, BLOCK_N: tl.constexpr): cols tl.arange(0, BLOCK_N) mask cols N # 加载后立即升精度到 FP32 x tl.load(X cols, maskmask, other0.0).to(tl.float32) # FP32 归约 mean tl.sum(x, axis0) / N xbar tl.where(cols.to(tl.float32) N, x - mean, 0.0) var tl.sum(xbar * xbar, axis0) / N rstd 1 / tl.sqrt(var eps) out (x - mean) * rstd # 写回时降精度 tl.store(Out cols, out, maskmask)2.3 归约升精度速查表操作输入类型是否需要手动升精度原因tl.sumfp16是fp16 直接归约精度不足需手动.to(tl.float32)tl.sumbf16否编译器自动提升为 fp32tl.sumint8/int16是可能溢出需提升为 int32tl.max/tl.minfp16否编译器自动提升为 fp32bitwidth32 且 is_floatingtl.max/tl.minbf16否编译器自动提升为 fp32tl.dot累加fp16/bf16否硬件默认 fp32 累加源码参考semantic.py 中 bf16 自动提升说明3. 矩阵乘法混合精度模式3.1 标准模式低精度输入 - FP32 累加 - 低精度写回存储格式 (fp16/bf16) ──load── 计算格式 (fp16/bf16) ──dot── 累加器 (fp32) ──store── 存储格式 (fp16/bf16)3.2 标准写法triton.jit def matmul_kernel(A, B, C, M, N, K, stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr): pid_m tl.program_id(0) pid_n tl.program_id(1) offs_m pid_m * BLOCK_M tl.arange(0, BLOCK_M) offs_n pid_n * BLOCK_N tl.arange(0, BLOCK_N) offs_k tl.arange(0, BLOCK_K) # 累加器使用 FP32 acc tl.zeros((BLOCK_M, BLOCK_N), dtypetl.float32) for k in range(0, K, BLOCK_K): # 输入使用 fp16/bf16 a tl.load(A offs_m[:, None] * stride_am (offs_k[None, :] k) * stride_ak, mask(offs_m[:, None] M) (offs_k[None, :] k K), other0.0) b tl.load(B (offs_k[:, None] k) * stride_bk offs_n[None, :] * stride_bn, mask(offs_k[:, None] k K) (offs_n[None, :] N), other0.0) # tl.dot 自动将 fp16/bf16 输入累加到 fp32 acc tl.dot(a, b) # 写回时降精度 tl.store(C offs_m[:, None] * stride_cm offs_n[None, :] * stride_cn, acc.to(tl.float16), mask(offs_m[:, None] M) (offs_n[None, :] N))3.3 tl.dot 精度控制参数参数默认值可选值说明allow_tf32FalseTrue/FalseNPU 上 tf32 自动映射为 hf32input_precisionieeeieee/hf32仅 fp32 输入时 hf32 生效非 fp32 输入回退 ieee3.4 tl.dot 支持的输入类型组合输入 A输入 B累加器out_dtype说明int8int8fp32int32量化推理fp16fp16fp32fp32标准训练bf16bf16fp32fp32NPU 推荐推理fp32fp32fp32fp32高精度计算注意acc不支持 FP16硬件默认使用 FP32 累加。max_num_imprecise_acc暂不支持。源码参考04-linear-algebra-ops.md4. BF16 精度保护注意事项4.1 BF16 vs FP16 格式对比特性float16 (fp16)bfloat16 (bf16)总位数1616尾数位107指数位58指数偏移15127表示范围~5.96e-8 ~ 65504~1.18e-38 ~ 3.39e38十进制精度~3.3 位~2.1 位与 fp32 转换需调整指数和尾数仅截断/扩展尾数4.2 BF16 精度保护关键点归约自动提升bf16 输入的归约操作sum/max/min会自动提升为 fp32 执行因为 NPU 不支持 bf16 的 FMAX/FMIN/FCMP 操作。这是编译器行为无需手动处理但会产生额外的类型转换开销。比较操作自动提升bf16 张量参与比较操作、、、maximum、minimum时会自动提升为 fp32 执行。除法/取模自动提升bf16 的除法/和取模%运算会自动提升为 fp32。混合运算提升bf16 与 fp16 混合运算会自动提升为 fp32产生额外转换开销。建议同一 kernel 内统一使用一种 16-bit 浮点类型。libdevice 函数限制libdevice 路径下acos、asin、sinh、cosh、acosh、asinh、atanh、atan2、hypot等函数不支持 bf16 输入需先手动.to(tl.float32)转换。bf16 尾数仅 7 位对于需要高精度的累加或迭代计算bf16 的精度可能不足。关键中间结果应在 fp32 下计算。4.3 BF16 选择建议场景推荐类型原因矩阵乘法输入bf16范围与 fp32 一致不易溢出训练前向传播bf16范围大梯度稳定与 fp32 混合运算bf16bf16 与 fp32 转换开销更低需要高精度的累加fp32累加器默认使用 fp32训练反向传播fp32 或 bf16梯度计算需要足够精度对精度要求极高fp16 或 fp32fp16 尾数精度更高源码参考08-comparison-logical-ops.md 中 bf16 比较操作说明11-libdevice.md 中 bf16 限制5. care_padding 参数详解5.1 参数定义care_padding是tl.load的 NPU 专属扩展参数控制maskFalse时 padding 区域的填充行为tl.load(pointer, maskmask, otherNone, care_paddingTrue)源码参考core.py5.2 行为差异care_paddingother 参数padding 区域行为性能影响True默认None填充 0浮点填 0.0整数填 0MTE2 等待 Vector 初始化降低并行度True指定值填充指定值MTE2 等待 Vector 初始化降低并行度FalseNone随机值未定义MTE2 与 Vector 无依赖提升并行度任意非 None使用other指定的值care_padding不生效5.3 执行时序对比care_paddingTrue默认 Vector: |初始化全0| |计算| MTE2: |搬运有效数据| ↑ 必须等待初始化完成 care_paddingFalse MTE2: |搬运1|搬运2|搬运3| Vector: |计算1|计算2|计算3| ↑ 无需等待直接并行5.4 安全使用 care_paddingFalse 的场景场景 1Element-wise 操作padding 不被 storetriton.jit def add_kernel(x_ptr, y_ptr, out_ptr, n, BLOCK_SIZE: tl.constexpr): pid tl.program_id(0) offsets pid * BLOCK_SIZE tl.arange(0, BLOCK_SIZE) mask offsets n x tl.load(x_ptr offsets, maskmask, care_paddingFalse) y tl.load(y_ptr offsets, maskmask, care_paddingFalse) output x y tl.store(out_ptr offsets, output, maskmask)场景 2padding 被后续 where/select 覆盖triton.jit def masked_fill_kernel(inp, expand_mask, value, out, N, BLOCK_SIZE: tl.constexpr): pid tl.program_id(0) offsets pid * BLOCK_SIZE tl.arange(0, BLOCK_SIZE) mask offsets N input_vals tl.load(inp offsets, maskmask, care_paddingFalse) fill_mask_vals tl.load(expand_mask offsets, maskmask).to(tl.int1) result tl.where(fill_mask_vals, value, input_vals) tl.store(out offsets, result, maskmask)场景 3矩阵乘法 K 维度尾部for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)): x tl.load(x_ptrs, maskoffset_k[None, :] K - k * BLOCK_SIZE_K, other0.0, care_paddingFalse) w tl.load(w_ptrs, maskoffset_k[:, None] K - k * BLOCK_SIZE_K, other0.0, care_paddingFalse) accumulator tl.dot(x, w)5.5 不安全场景场景原因正确做法Reduction (sum)padding 随机值被累加使用other0.0Reduction (max)padding 随机值可能成为最大值使用other-float(inf)中间结果依赖 padding后续计算使用 padding 值保持care_paddingTrueStore 无 maskpadding 值被写出确保 store 有正确 mask5.6 性能影响访存密集型算子care_paddingFalse可提升 10%-30% 性能计算密集型算子提升较小瓶颈不在搬运结合 for 循环 Tiling care_paddingFalse效果最佳源码参考04-care-padding.md6. 精度验证方法6.1 torch.testing.assert_close 推荐参数场景rtolatol说明fp32 通用计算1e-51e-5标准浮点验证fp16/bf16 向量运算1e-31e-3半精度通用验证fp16/bf16 矩阵乘法2^-6 (~0.016)2^-6 (~0.016)matmul 精度差异在 1e-2 范围内属正常Softmax1e-51e-7NPU 上 tl.exp 是近似计算微小差异正常int8 量化推理00整数运算应精确匹配6.2 矩阵乘法分段验证对于 bf16/fp16 输入的 matmul绝对值较小和较大的区域应使用不同的验证策略mask golden.abs() 1.0 torch.testing.assert_close(result[mask], golden[mask], atol2**-6, rtol0) torch.testing.assert_close(result[~mask], golden[~mask], atol0, rtol2**-6)6.3 通用验证模板def verify_kernel(kernel_fn, *args, dtypetorch.float16, **kwargs): for size in [127, 128, 255, 256, 1023, 1024, 4096]: x torch.randn(size, dtypedtype, devicenpu) y_triton kernel_fn(x) y_torch torch_reference(x) if dtype in (torch.float16, torch.bfloat16): torch.testing.assert_close(y_triton, y_torch, rtol1e-3, atol1e-3) else: torch.testing.assert_close(y_triton, y_torch, rtol1e-5, atol1e-5)6.4 精度差异排查流程结果与参考不一致 ├── 差异在 1e-6 量级 → 浮点计算顺序差异属正常 ├── 差异在 1e-3 量级 → 检查是否缺少 .to(tl.float32) 升精度 ├── 差异在 1e-2 量级 → matmul 场景属正常其他场景检查累加器精度 ├── 差异较大 → │ ├── 检查 int64/int32 是否导致 scalar 退化 │ ├── 检查 bf16/fp16 精度损失 │ ├── 检查 mask/boundary_check 逻辑 │ └── 使用 TRITON_INTERPRET1 在 CPU 上运行作为基准 └── Softmax 差异 → 检查减最大值和 exp 近似源码参考05-faq.md7. 类型转换的性能影响7.1 类型转换开销分级转换类型开销说明bf16 - fp32低仅截断/扩展尾数指数位相同fp16 - fp32中需调整指数和尾数bf16 - fp16中需经过 fp32 中间转换fp8 - fp16/bf16中910_95 支持A2/A3 不支持int8 - fp32低整数到浮点转换int64 - fp32中Vector CMP 不支持 int64需转换避免 scalar 退化fp32 - fp16/bf16低降精度可控制舍入模式7.2 隐式类型提升规则computation_type二元运算中操作数类型会自动提升。了解这些规则有助于避免不必要的性能开销操作数 A操作数 B运算提升结果性能影响fp16fp16,-,*fp16无额外开销fp16fp16/,%fp32除法/取模自动升精度bf16bf16,-,*bf16无额外开销bf16bf16/,%fp32除法/取模自动升精度bf16fp16任意fp32混合类型额外转换开销fp32fp16/bf16任意fp32有 fp32 则提升fp8e4nvfp8e5任意fp16不同 fp8 提升为 fp16int8int32任意int32整数类型提升7.3 类型转换优化建议避免 bf16 与 fp16 混合运算混合运算自动提升为 fp32产生额外转换开销和内存占用。同一 kernel 内统一使用一种 16-bit 浮点类型。Vector CMP 类型转换NPU 的 Vector CMP 不支持 int64/int32会导致 scalar 退化。需要手动转换为 fp32# 优化前cols 是 int32CMP 退化为 scalar xbar tl.where(cols N, x - mean, 0.0) # 优化后转为 fp32 使用 Vector CMP cols_cmp cols.to(tl.float32) xbar tl.where(cols_cmp N, x - mean, 0.0)减少不必要的精度切换在 kernel 内部统一使用 fp32 计算仅在 load/store 时进行类型转换避免反复转换。fp64 替换为 fp32Ascend NPU 不支持 fp64所有 fp64 使用需替换为 fp32。uint 类型替换uint8/16/32/64 不支持需在 host 端转换为对应的 int 类型。降精度舍入模式控制# 默认 rtneRound To Nearest Even y tl.cast(x, tl.bfloat16) # 指定 rtzRound Toward Zero y tl.cast(x, tl.bfloat16, fp_downcast_roundingrtz)整数溢出保护# 默认 trunc截断溢出 y tl.cast(x, tl.int8) # 指定 saturate饱和截断Ascend 扩展 y tl.cast(x, tl.int8, overflow_modesaturate)源码参考semantic.py 中 integer_promote 和 computation_type 规则8. 910_95 特别注意8.1 FP8 支持910_95 系列对 FP8 的支持比 A2/A3 系列有显著增强功能A2/A3910_95FP8 类型转换不支持支持FP8 dot_scaled不支持支持FP8 tl.dot 输入不支持不支持需先转换为 fp16/bf168.2 dot_scaled 支持910_95 支持 FP8 格式的tl.dot_scaled缩放张量值为 int8GPU 为 uint8缩放张量类型fp4fp8bf16fp16910_95不支持支持支持支持8.3 其他 910_95 差异UB 空间256KBA2/A3 为 192KB开启 double buffer 时不超过 128KBfixpipe 直通支持 L0C - UB 直通路径详见 11-fixpipe-and-bias-fusion.mdMultiBuffer 默认关闭需显式设置multibufferTrue详见 07-compile-params.md完整硬件规格详见 00-hardware-quick-ref.md9. 精度保护标准写法速查9.1 归约操作x_fp32 x.to(tl.float32) result tl.sum(x_fp32, axis-1)9.2 矩阵乘法acc tl.zeros([BLOCK_M, BLOCK_N], dtypetl.float32) for k in range(0, K, BLOCK_K): a tl.load(a_ptr ...) b tl.load(b_ptr ...) acc tl.dot(a, b) tl.store(c_ptr ..., acc.to(tl.float16))9.3 LayerNormx tl.load(X cols, maskmask, other0.0).to(tl.float32) mean tl.sum(x, axis0) / N xbar tl.where(cols.to(tl.float32) N, x - mean, 0.0) var tl.sum(xbar * xbar, axis0) / N rstd 1 / tl.sqrt(var eps) out (x - mean) * rstd tl.store(Out cols, out, maskmask)9.4 Softmaxx tl.load(x_ptr offsets, maskmask).to(tl.float32) x_max tl.max(x, axis-1) x_shifted x - x_max[:, None] numerator tl.exp(x_shifted) denominator tl.sum(numerator, axis-1)[:, None] result numerator / denominator tl.store(out_ptr offsets, result.to(tl.float16), maskmask)相关文档链接数据类型支持矩阵 - 完整的数据类型支持矩阵数据类型详解 - 数据类型详细规格与对齐要求care_padding 优化 - care_padding 参数完整说明归约操作 API - 归约操作精度行为线性代数 API - tl.dot 精度控制比较与逻辑操作 API - bf16 比较操作说明FAQ - 精度相关常见问题API 差异 - GPU/NPU API 行为差异【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考