Triton Operator Precision Eval

Triton算子精度评估技能

1. 功能概述

该技能用于自动化评估Triton算子实现的精度，通过与PyTorch（CPU或NPU）的对应算子实现进行比对，生成详细的精度验证报告。

核心功能

• 自动接收Triton算子实现
• 支持与CPU或NPU上的Torch小算子进行比对
• 支持多种数据类型（float16、float32、int8、uint8等）
• 自动生成精度验证报告
• 支持批量测试不同参数配置

2. 调用时机

在以下情况下调用此技能：

• 需要验证Triton算子实现的精度正确性
• 需要与PyTorch原生算子进行精度比对
• 需要生成标准化的精度验证报告
• 需要批量测试Triton算子在不同数据类型和参数下的精度表现

3. 实现原理

3.1 工作流程

┌─────────────────┐     ┌─────────────────┐     ┌─────────────────┐
│ Triton算子实现  │────▶│ 生成测试数据    │────▶│ 执行Torch对比实现 │
└─────────────────┘     └─────────────────┘     └─────────────────┘
          ▲                     │                     │
          │                     ▼                     ▼
          │              ┌─────────────────┐     ┌─────────────────┐
          │              │ 执行Triton实现  │     │ 计算误差指标    │
          │              └─────────────────┘     └─────────────────┘
          │                     │                     │
          └─────────────────────┼─────────────────────┘
                                │
                                ▼
                        ┌─────────────────┐
                        │ 生成精度报告    │
                        └─────────────────┘

3.2 核心组件

1. 测试数据生成：使用 test_common.generate_numpy() 生成随机测试数据
2. Torch对比实现：用户提供的Torch算子实现（如示例中的 torch_pointwise()）
3. Triton算子执行：使用Triton JIT编译并执行用户提供的Triton kernel
4. 精度验证：使用 test_common.validate_cmp() 进行精度比对，支持不同数据类型的误差阈值
5. 报告生成：生成包含误差指标的精度验证报告

4. 使用方法

4.1 前置条件

• 已安装Triton和PyTorch环境
• 已安装 torch_npu（如果使用NPU进行测试）
• 已准备Triton算子实现代码

4.2 编写测试用例

创建测试文件（如 test_abs.py），包含以下内容：

1. 导入必要模块：

   import triton
   import triton.language as tl
   import numpy as np
   import torch
   import pytest
   import test_common
   

2. 实现Torch对比函数：

   def torch_pointwise(x0):
       # 实现与Triton算子对应的Torch功能
       return torch.abs(x0)
   

3. 实现Triton算子：

   @triton.jit
   def triton_abs(in_ptr0, out_ptr0, XBLOCK: tl.constexpr, XBLOCK_SUB: tl.constexpr):
       # Triton kernel实现
       offset = tl.program_id(0) * XBLOCK
       base1 = tl.arange(0, XBLOCK_SUB)
       loops1: tl.constexpr = (XBLOCK + XBLOCK_SUB - 1) // XBLOCK_SUB
       for loop1 in range(loops1):
           x0_prime = offset + (loop1 * XBLOCK_SUB) + base1
           x0 = offset + (loop1 * XBLOCK_SUB) + base1
           tmp0 = tl.load(in_ptr0 + (x0), None)
           tmp2 = tl.abs(tmp0)
           tl.store(out_ptr0 + (x0), tmp2, None)
   

4. 编写测试用例：

   @pytest.mark.parametrize('param_list',
                          [
                              ['float16', (2, 4096, 8), 32, 2048, 64],
                              ['float32', (2, 4096, 8), 32, 2048, 64],
                              ['int8', (2, 4096, 8), 32, 2048, 64],
                              ['uint8', (2, 4096, 8), 32, 2048, 64],
                          ]
                          )
   
   def test_case(param_list):
       dtype, shape, ncore, xblock, xblock_sub = param_list
       np_x0 = test_common.generate_numpy(shape, dtype)
       x0 = torch.from_numpy(np_x0).to(eval('torch.' + dtype)).npu()
       y_ref = torch_pointwise(x0)
       y_cal = torch.zeros(shape, dtype = eval('torch.' + dtype)).npu()
       triton_abs[ncore, 1, 1](x0, y_cal, xblock, xblock_sub)
       test_common.validate_cmp(dtype, y_cal, y_ref)
   

4.3 运行测试

# 运行单个测试文件
pytest test_abs.py -v

# 运行所有测试文件
pytest ./examples/ -v

5. 精度验证规则

5.1 不同数据类型的验证规则

5.2 误差指标

• 平均相对误差(MERE)：所有元素相对误差的平均值
• 最大相对误差(MARE)：所有元素相对误差的最大值
• 绝对误差：元素值之差的绝对值

6. 精度报告格式

生成的精度报告（如 eco_report.txt）包含以下内容：

================================================================================
                              Triton算子精度验证报告                               
--------------------------------------------------------------------------------
[验证配置]:
  数据类型: float32 (Single Precision)
  MERE阈值: 1.220703e-04
  MARE阈值: 1.220703e-03 (10×MERE阈值)
  小值域阈值: 1.000000e-07
--------------------------------------------------------------------------------
[精度标准]:
  float16: 相对误差 rtol=1e-03, atol=1e-03
  float32: 相对误差 rtol=1e-04, atol=1e-04
  bfloat16: 相对误差 rtol=1e-03, atol=1e-03
  int32/int64/int16/int8: 完全相等
  uint32/uint64/uint16/uint8: 完全相等
  bool: 完全相等
--------------------------------------------------------------------------------
[验证结果]:
  验证结果: FAIL
  样本总数: 4096
--------------------------------------------------------------------------------
[误差指标]:
  平均相对误差(MERE): 6.642197e-03
    阈值要求: MERE < 1.220703e-04
  最大相对误差(MARE): 3.458786e+00
    阈值要求: MARE < 1.220703e-03
--------------------------------------------------------------------------------
[判定条件]:
  ✓ MERE < 阈值: False
  ✓ MARE < 10×阈值: False
  ✓ 总体结果: False
================================================================================

6.1 报告内容要求

精度报告必须包含以下内容：

1. 验证配置：算子名称、测试形状、数据类型、NPU核心数等
2. 精度标准：每个数据类型的具体精度要求（误差阈值或完全相等）
3. 验证结果：测试总数、通过数量、失败数量、总体结果
4. 详细误差指标：每个数据类型的平均相对误差、最大相对误差、最大绝对误差
5. 判定条件：所有数据类型测试通过的状态

其中，精度标准部分必须列出所有支持的数据类型及其对应的精度要求，确保报告的可读性和可追溯性。

7. 示例

7.1 输入示例（Triton算子实现）

@triton.jit
def triton_abs(in_ptr0, out_ptr0, XBLOCK: tl.constexpr, XBLOCK_SUB: tl.constexpr):
    offset = tl.program_id(0) * XBLOCK
    base1 = tl.arange(0, XBLOCK_SUB)
    loops1: tl.constexpr = (XBLOCK + XBLOCK_SUB - 1) // XBLOCK_SUB
    for loop1 in range(loops1):
        x0_prime = offset + (loop1 * XBLOCK_SUB) + base1
        x0 = offset + (loop1 * XBLOCK_SUB) + base1
        tmp0 = tl.load(in_ptr0 + (x0), None)
        tmp2 = tl.abs(tmp0)
        tl.store(out_ptr0 + (x0), tmp2, None)

7.2 输出示例（精度报告）

================================================================================
                              Triton算子精度验证报告                               
--------------------------------------------------------------------------------
[验证配置]:
  数据类型: float32 (Single Precision)
  MERE阈值: 1.220703e-04
  MARE阈值: 1.220703e-03 (10×MERE阈值)
  小值域阈值: 1.000000e-07
--------------------------------------------------------------------------------
[验证结果]:
  验证结果: PASS
  样本总数: 4096
--------------------------------------------------------------------------------
[误差指标]:
  平均相对误差(MERE): 0.000000e+00
    阈值要求: MERE < 1.220703e-04
  最大相对误差(MARE): 0.000000e+00
    阈值要求: MARE < 1.220703e-03
--------------------------------------------------------------------------------
[判定条件]:
  ✓ MERE < 阈值: True
  ✓ MARE < 10×阈值: True
  ✓ 总体结果: True
================================================================================

8. 注意事项

1. 环境配置：确保已正确安装Triton、PyTorch和必要的依赖
2. 数据类型：不同数据类型有不同的精度验证规则，请根据算子类型选择合适的数据类型
3. 参数配置：根据算子复杂度和硬件资源调整测试参数（如ncore、xblock等）
4. 错误处理：如果测试失败，检查Triton算子实现和Torch对比实现是否功能一致
5. 报告解读：根据精度报告中的误差指标调整算子实现以提高精度

9. 故障处理

10. 参考资料

• Triton官方文档
• PyTorch官方文档
• 昇腾PyTorch插件文档