external-cannbot-ops-triton-op-verifier

Kernel Verifier Skill

Verification Workflow

输入：generated_code.py + task_file.py
    ↓
[0. Triton 退化预检查] → scripts/validate_triton_impl.py (AST 静态分析)
    ↓ (通过)
[1. 创建验证项目] → 两个文件
    ↓
[2. 执行验证脚本] → scripts/verify.py --op_name ...
    ↓
[3. 收集验证结果]
    ↓
[验证通过] → [4. 执行性能测试] → scripts/benchmark.py --op_name ...
    ↓
[5. 收集性能结果]
    ↓
输出：验证结果 + 性能数据

---

Step 0: Triton Degradation Pre-check (AST Static Analysis)

Before creating the verification project, first use validate_triton_impl.py to perform degradation detection on the generated code. This check is pure AST static analysis, requires no NPU/torch runtime, and completes in milliseconds.

Command Template：

python3 <本skill所在目录的绝对路径>/scripts/validate_triton_impl.py \
    <生成代码文件路径> --json

Detect three degradation types：

Type	Meaning	Detection Method
Type 1	Completely absent @triton.jit kernel	No triton.jit decorated function definition in the AST
Type 2	Has kernel but forward() is not called	Kernel definition exists but ModelNew.forward() is not referenced (including wrapper function tracing)
Type 3	Part of the computation uses PyTorch	forward() Prohibited torch.* / F.* computation operations exist in (accurate to the line number)

Result Determination：

• exit code == 0 → passed, continue to Step 1
• exit code != 0 → degradation detected, parse the regression_type and suggestionin the JSON, and return failure directly

JSON Output Format：

{
  "valid": false,
  "regression_type": 3,
  "checks": {
    "triton_kernel_exists": {"passed": true, "kernels": [...]},
    "kernel_called_from_forward": {"passed": true, "called": [...]},
    "no_forbidden_torch_ops": {"passed": false, "violations": [{"line": 45, "call": "F.softmax", "reason": "..."}]}
  },
  "suggestion": "..."
}

---

Step 1: Create Verification Project

In the verification directory of the current iteration (e.g. {output-path}/iter_{iteration}/verify/), create two files:

File 1: {op_name}_torch.py

Directly copy the complete contents of the task file. This file contains Model、get_inputs()、get_init_inputs()。

File 2: {op_name}_triton_ascend_impl.py

Directly copy the complete contents of the generated code. This file contains the ModelNew class.

---

Step 2: Execute Verification (⚠️ You must use this script; creating your own test method is prohibited)

You must use bash tool to call the scripts/verify.py script included with this skill.

Command Template：

python3 <本skill所在目录的绝对路径>/scripts/verify.py \
    --op_name <算子名> \
    --verify_dir <验证目录> \
    --triton_impl_name <triton实现模块名> \
    --timeout 900

Actual Invocation Example(assuming the verification directory is /tmp/workspace/softmax/verify, and the operator name is softmax）：

python3 /path/to/triton-op-verifier/scripts/verify.py \
    --op_name softmax \
    --verify_dir /tmp/workspace/softmax/verify \
    --triton_impl_name triton_ascend_impl \
    --timeout 900

Parameter Description：

Parameter	Required	Description
--op_name	Yes	Operator name, corresponding to the file name prefix
--verify_dir	No	Verification directory path, defaults to the current directory
--triton_impl_name	No	Triton implementation module name (excluding the {op_name}_ prefix), defaults to triton_ascend_impl
--timeout	No	Timeout in seconds, default 900

Timeout Setting: default 900 seconds; can be appropriately increased for complex operators.

⛔ Prohibitions：

• Writing your own Python code to test the operator is prohibited (e.g. manually importing and comparing forward)
• Using torch.allclose or other self-created methods to replace scripts/verify.py
• Skipping this step and directly reporting the verification result is prohibited

---

Step 3: Collect Verification Results

verify.py will generate verify_dir under verify_result.json(or --output specified path), including:

{
  "op_name": "softmax",
  "total_cases": 5,
  "passed_cases": 4,
  "failed_cases": 1,
  "failures": [
    {
      "case_idx": 2,
      "input_desc": [
        {"type": "tensor", "shape": [128, 256], "dtype": "torch.float16"}
      ],
      "error_type": "CompilationError",
      "error_msg": "..."
    }
  ]
}

Multi-shape Behavior: Each shape has independent try/except; a failure does not stop subsequent shapes; only after all have finished running is it written to disk and exited.

Exit code semantics (Policy A: strict)：

• passed_cases == total_cases → exit 0，verifier_result = true
• passed_cases < total_cases → exit 1，verifier_result = false，verifier_error should read verify_result.json.failures ofall entries(not just the first), and after aggregation submit to the Conductor.

Timeout: the script outputs "验证超时" and the exit code is 1 → verifier_error = "验证超时（{timeout}秒）"。

---

Step 4: Execute performance testing (execute after verification passes)

Preconditions (enforced at the L1 script layer): when benchmark.py starts, it automatically derives the corresponding --triton_impl_name from verify_result file and validates passed_cases == total_cases; if it does not pass, directly exit 2, and running benchmark is prohibited. See the "L1 verify gate" section below for details.

Only executed when verify.py's passed_cases == total_cases (Policy A). If verify has any failure → executing benchmark.py is prohibited.

Use the bash tool to call the scripts/benchmark.py script bundled with this skill.

Command template：

python3 <本skill所在目录的绝对路径>/scripts/benchmark.py \
    --op_name <算子名> \
    --verify_dir <验证目录> \
    --triton_impl_name <triton实现模块名> \
    --warmup <warmup次数> \
    --repeats <测试次数> \
    --output <输出文件路径>

Actual invocation example：

python3 /path/to/triton-op-verifier/scripts/benchmark.py \
    --op_name softmax \
    --verify_dir /tmp/workspace/softmax/verify \
    --triton_impl_name triton_ascend_impl \
    --warmup 5 \
    --repeats 50 \
    --output /tmp/workspace/softmax/iter_0/perf_result.json

Note：--output The path is specified by the caller, and the performance report will be written to that path. It is usually specified by the kernelgen-workflow SubAgent as {output-path}/iter_{iteration}/perf_result.json。

Parameter description：

Parameter	Required	Description
--op_name	Yes	Operator name
--verify_dir	No	Verification directory path, defaults to the current directory
--triton_impl_name	No	Triton implementation module name (excluding the {op_name}_ prefix), default triton_ascend_impl
--warmup	No	Number of warmup iterations, default 5
--repeats	No	Number of formal test iterations, default 50
--output	No	Performance report output path (JSON format)
--verify_not_required	No	Skip the L1 verify gate (by default, all verify_result must pass)

---

L1 verify gate

When benchmark.py starts, it derives the corresponding verify_result filename from --triton_impl_name :

triton_impl_name	Corresponding verify json
triton_ascend_impl(default, Phase 3)	verify_result.json
triton_baseline（Phase 4 baseline）	verify_result_baseline.json
triton_optimized（Phase 4 optimized）	verify_result_optimized.json
Other triton_xxx	verify_result_xxx.json

Judgment rules(enabled by default; passing --verify_not_required can skip it):

Case	Exit code	Description
File does not exist	exit 2	verify.py must be run first
File read failure	exit 2	JSON is corrupted
total_cases == 0	exit 2	verify did not actually run any shape
passed_cases < total_cases	exit 2	Precision did not fully pass; benchmark is meaningless and will contaminate downstream
passed_cases == total_cases > 0	Continue executing benchmark	—

On exit 2, stderr will print: verify_json path / passed/total / first 5 failures, so the upstream agent can equivalently map the error to the verify failure handling path.

---

Step 5: Collect performance results

After performance testing is complete, read the results from the JSON file specified by --output .

Performance report format

{
  "op_name": "softmax",
  "warmup": 5,
  "repeats": 50,
  "total_cases": 3,
  "passed_cases": 3,
  "failed_cases": 0,
  "nan_indices": [],
  "inf_indices": [],
  "zero_indices": [],
  "negative_indices": [],
  "none_indices": [],
  "framework": {
    "avg_latency_ms": 1.2345,
    "peak_memory_mb": 256.00,
    "operators": {"...": 0.0}
  },
  "implementation": {
    "avg_latency_ms": 0.5678,
    "peak_memory_mb": 128.00,
    "operators": {"...": 0.0}
  },
  "speedup_vs_torch": 2.1746,
  "per_shape_results": [
    {
      "case_idx": 1,
      "input_desc": [{"type":"tensor","shape":[128,256],"dtype":"torch.float16"}],
      "status": "pass",
      "framework": {"avg_latency_ms": 1.23, "peak_memory_mb": 64.0},
      "implementation": {"avg_latency_ms": 0.56, "peak_memory_mb": 32.0},
      "speedup_vs_torch": 2.19,
      "error_type": null,
      "error_msg": null
    }
  ]
}

Field description：

Metric	Description
avg_latency_ms	Arithmetic mean of latency across shapes (compatible semantics)
peak_memory_mb	Peak memory usage (MB)
speedup_vs_torch	Geometric mean speedup = (∏ s_i)^(1/n), only for status==pass and s_i Take the geometric mean over the shapes that are finite positive numbers; when all are abnormal, it is null
passed_cases / failed_cases	Multiple shapes pass/fail counts (abnormal shapes are still counted in passed_cases, because the operator functions normally)
nan_indices / inf_indices / zero_indices / negative_indices / none_indices	case_idx list for each type of abnormal s_i (starting from 1), not included in the geometric mean; when there are no abnormalities, it is []
per_shape_results[].status	"pass" or "fail"
per_shape_results[].speedup_vs_torch	The speedup ratio of the shape; when fail or abnormal, it is null

Boundary value handling：

s_i = framework_latency_ms / impl_latency_ms Abnormal values may appear due to profiler failure, extremely small latency, etc.compute_overall For each s_i classify by the following priority:

Category	Determination	Persistence behavior
none	s_i is None	per_shape.speedup_vs_torch = null, case_idx goes into none_indices
nan	math.isnan(s_i)	Same as above, goes into nan_indices
inf	math.isinf(s_i)	Same as above, goes into inf_indices
negative	s_i < 0	Same as above, goes into negative_indices
zero	s_i == 0	Same as above, goes into zero_indices
valid	Finite positive number	Included in the geometric mean

Abnormal shape Still counted in passed_cases(the operator functions normally, only the measurement data is unreliable), but s_i does not participate in the overall geometric mean. When all shapes are abnormal, speedup_vs_torch = null。

Exit code：

• exit 0: benchmark completes normally (try/except inside per shape, pass/fail written to per_shape_results)
• exit 1: the script itself crashes
• exit 2: L1 verify gate rejects (precondition not satisfied, benchmark did not actually run)

The caller determines by reading the JSON passed_cases == total_cases; when exit 2 there is no JSON output, which should be treated as equivalent to "the corresponding verify failed".

Returns：

• perf_result: dict (complete performance data)
• perf_report_path: str (performance report file path)

---

Accuracy threshold description

Verification uses a data-type-based MERE/MARE dual-threshold relative errordetermination (NPU Benchmark standard), which is different from torch.allclose .

Determination formula(both must be satisfied):

MERE < threshold        且        MARE < 10 × threshold

Where:

• MERE = mean(|actual - golden| / max(|golden|, threshold)), mean relative error
• MARE = max(|actual - golden| / max(|golden|, threshold)), maximum relative error
• Before calculation, both sides are uniformly promoted to float32 to avoid contamination from the low-precision dtype's own error
• The denominator uses clamp(min=threshold) instead of +epsilon: when |golden| < threshold(the reference value is already as small as the dtype precision limit), rel_err degenerates to |diff| / threshold, which is equivalent to normalizing by absolute error, avoiding false positives near zero or extremely small values

dtype threshold table(powers of 2):

Data type	threshold	MERE upper limit	MARE upper limit (10×t)
float16	2⁻¹⁰ ≈ 9.77e-4	9.77e-4	9.77e-3
bfloat16	2⁻⁷ ≈ 7.81e-3	7.81e-3	7.81e-2
float32	2⁻¹³ ≈ 1.22e-4	1.22e-4	1.22e-3
hifloat32	2⁻¹¹ ≈ 4.88e-4	4.88e-4	4.88e-3
float8_e4m3	2⁻³ = 0.125	0.125	1.25
float8_e5m2	2⁻² = 0.25	0.25	2.5
Other dtype (fallback)	2⁻¹³	1.22e-4	1.22e-3

Pre-comparison checks(in order, any failure is judged fail):

1. Shapes must be consistent
2. NaN positions must be completely identical (mask bitwise equal)
3. Inf positions and signs must be completely identical
4. bool dtype: requires torch.equal to be completely equal, does not enter MERE/MARE determination
5. Only performs MERE/MARE calculation on finite_mask ; when the dtype is inconsistent, impl will be cast to golden's dtype

---

Script location

The validation scripts are located in the scripts/ directory of this skill:

Script	Purpose
scripts/validate_triton_impl.py	Degradation pre-check (AST static analysis)
scripts/verify.py	Verify correctness
scripts/benchmark.py	Test performance

CLI parameters：

• validate_triton_impl.py: <file_path>, [--json]
• verify.py: --op_name, --verify_dir, --triton_impl_name, --timeout, --output
• benchmark.py: --op_name, --verify_dir, --triton_impl_name, --warmup, --repeats, --output, --skip_framework, --framework_latency_ms, --verify_not_required