Merge blocked
Hi @Mocchibird, welcome to submitting your first PR to CANN!
PR Merge Steps
1. CLA Signing
If the current PR label includes cann-cla/yes, it means you have signed the CLA and can proceed to the next step. If the label includes cann-cla/no, please sign the CLA first. If you have any questions, please refer to the FAQ.
2. CI Check
Please comment /compile to trigger the CI pipeline check. If the CI run is successful, the PR will be tagged with ci-pipeline-passed and you can proceed to the next step. If the CI run fails, the PR will be tagged with ci-pipeline-failed, please check the CI logs to fix the issues in the PR. If you have any questions, please refer to the FAQ.
3. Code Review
After CI passes, please refer to the PR Approval Progress and proactively @ the committers in the table to review the code. After approval, committers will comment /lgtm and /approve. Once the lgtm and approved labels are successfully added, the PR will be merged automatically.


Thanks for your pull-request.
The full list of commands accepted by me can be found at here.
You can get sig-info at here.
You can self-configure the PR merge rules for this repository. For more details, please refer to Here.
For more, you also can visit HICANN.
PR Approval Progress
⚠️ This PR does not yet meet the following requirements:lgtm (requires ≥ 2 person(s) per module)、approve (requires ≥ 1 person(s) per module)
Module Approval Details
| module | lgtm status | approve status |
|---|---|---|
| repo-cann/ops-tilelang | ❌ (0/2)(You can also ask: 俞郑中, jiawei_zhuang, kevin365, ChaoyangJi, 马兵) | ❌ (0/1)(You can also ask: 黄俊健, kevin365, 蒋礼锐, 刘丹, jiawei_zhuang) |
💡 Tip:
- Committer can comment
/approveor/lgtm- Commenting
/approveimplies both code review (lgtm) and intent to merge (approve)
CLA Signature Pass
Mocchibird, thanks for your pull request. All authors of the commits have signed the CLA. 👍


流水线分析报告 — ops-tilelang run#73
流水线: https://gitcode.com/CANN/ops-tilelang/actions/runs/559c41c676f24d12beee3dcdae67e2dd | PR#44 | 2026-11-06 23:00:33 ~ 2026-11-07 05:03:13 | 失败 Job 1 个
定界结论: 环境问题 / Runner机器异常(排队超时,可重试)
- 失败类型: 环境问题 / Runner机器异常
- 根因: 唯一失败 Job
ST_test始终处于 Queued 状态,排队 300s 超时(CP.COMP1105),未进入执行阶段;与 PR#44 代码变更无关。 - 处置建议: 直接重新触发流水线即可(平台调度/执行机资源临时紧张,可重试)。
- 排除项: 其余 11 个 Job(Test_image / Test_get-pr / Test_pre_pipeline / Test_get-shell / Antipoison / CodeCheck_style / SCA / CodeCheck_check-pr / CodeCheck / CodeCheck_precommit / Test_finish)全部 COMPLETED,无代码、编译或测试类失败。
失败根因分析
根因 #1: 环境问题 / Runner机器异常 — ST_test 排队 300s 超时未调度
- 判据:
(?i)queued state timeout|任务排队超时(平台调度类,可重试) - 定位步骤:
0_初始化步骤.log - 影响 Job:
ST_test - 根因证据(日志原文):
[2026-11-07 05:02:52.994+08:00] [ERROR] [server] - Queued state timeout, job_run_id: 5922092ee1cc4530b3af72ae46b4df55_1568399727533895717, state: Queued, timeout_seconds: 300
[2026-11-07 05:02:53.223+08:00] [ERROR] [server] - [CP.COMP1105][Description] 任务排队超时[Description_en] Queued timeout.(addition_en) Queued state timeout ... [RootCause] 可能是执行机配置或者网络问题。... [Solution] 请检查执行机配置或者网络。
- 说明: 该 Job 的
1_Checkout.log/2_download.log/3_run_st_test.log/4_后处理步骤.log均为空,确认任务尚未开始执行即因排队超时被平台终止,非 PR 代码导致。
结构化定界摘要
[{"repo": "ops-tilelang", "run_number": 73, "Failed_Type": "环境问题", "Failed_Sub_Type": "Runner机器异常", "Failed_Reason": "ST_test 任务排队超时 300s(CP.COMP1105)未调度执行,与 PR 代码无关", "evidence": "[CP.COMP1105][Description] 任务排队超时 Queued state timeout, state: Queued, timeout_seconds: 300", "suggestion": "重新触发流水线重试;若持续复现则联系 CI 维护者检查执行机资源", "jobs": "ST_test"}]
由门禁失败定界分析自动生成 | repo: ops-tilelang | run: 559c41c676f24d12beee3dcdae67e2dd | PR#44


MACAQ: better W4A4 accuracy with a fused inference path
Project repository: huawei-csl/MACaQ.
Add dense and grouped MACAQ W4A4 matmul for Ascend 950 using exported TileLang Ascend APIs and the native CCE backend.
MACAQ (Model Agnostic Calibrated Quantization) applies channel scaling,
block-32 BF16 Hadamard rotation, MXFP4 activation quantization and matrix
multiplication in one mixed Vector/Cube kernel, producing BF16 output.
Calibration, weight transformation
(W / D) @ Hand Flip-Refinement are offline.The key result is better quantization accuracy with competitive small-batch
kernel latency. The online accuracy transform is included in every MACAQ
timing below. Native CANN quantization and matmul are unchanged independent
baselines. Large-prefill and grouped losses remain visible.
Accuracy results
Model perplexity: 23.7% lower than ordinary W4A4
Qwen3-0.6B on held-out WikiText-2; lower is better:
This is a CPU model-reference pilot, using 16,384 held-out tokens
(16,368 scored predictions, context 1,024). All 196 decoder linear projections
are quantized; embeddings and the LM head remain BF16. Calibration uses
1,024 training tokens, with no held-out hyperparameter tuning. It is not
full-dataset perplexity or full-model NPU execution. Algorithm 2 model
perplexity has not been measured.
Actual NPU outputs: lower error on all nine fixed projections
MACAQ lowers normalized mean squared error against both native FP4 scaling
presets on q/up/down projections at layers 0, 14 and 27. Native error is
1.58–4.61× MACAQ error for algorithm 0 and
1.68–4.14× for algorithm 2.
The measured ops-nn, TileLang regular and TileLang prepared outputs agree
exactly with the MACAQ CPU reference on all nine fixtures. Evaluation uses
1,024 separate held-out test tokens. Unsupported ops-nn dimensions are
zero-padded for accuracy coverage only, then sliced; those calls make no latency claim.
Accuracy protocol, all projection rows, grouped arithmetic checks and raw data.
Runtime and accuracy are separate experiments; this is not a joint
model-level speed-and-accuracy result.
Latency against existing native CANN operators
Baseline:
npu_dynamic_mx_quant(block32 E2M1, round-to-even,scale_alg=0)followed by
npu_quant_matmul, ornpu_grouped_matmulfor grouped workloads.MACAQ additionally performs scaling and Hadamard rotation. Each numerical
function is checked against its own CPU reference. Our split MACAQ paths
are ablations, not existing external baselines.
Native latency / MACAQ latency; above 1 favors MACAQ. Geometric means
give every configuration equal weight: all 24 (K,N) pairs at dense
M=16/128/4096 plus four-expert grouped workloads, for 96 cases in total.
Win counts compare medians and do not establish statistical significance.
Prepared plans cache static validation, layout views and kernel selection; dynamic guards, fresh outputs and workspace allocations remain inside each call. Prepared grouped calls explicitly promise
expert_block_rows=128; the regular API retains its 16-row contract. Both use the same padded fixtures and rows as native CANN.One call per configuration has total native/MACAQ time ratios 0.672× regular and 0.747× prepared; both are slower in that mixture. The prepared equal-case mean is near parity, with substantial small-batch gains and remaining large-case losses. An earlier same-computation sweep measured 1.036× prepared versus 1.052× here; that variation is not an algorithm gain.
The accuracy gain comes without a latency penalty on the faster measured
workloads. It is not free for every shape: large prefill and many grouped
cases still lose. No end-to-end NPU throughput gain is claimed.
Complete protocol, all 96 rows, heatmap and raw samples.
Regular MACAQ API results.
Timing protocol and current-source validation
Ascend 950, CANN 9.3 weekly, PyTorch 2.9.0 / torch_npu 2.9.0.post2,
native TileLang 0.1.15. Each case uses 20 warmups and 15 rotated trials of
20 eager calls. Two distinct activation buffers alternate. Outputs and all
intermediate allocations count inside timed calls; compilation, calibration
and static weight/layout preparation are excluded. No explicit cache flush.
These are repeated API timings, not single-request latency or model throughput.
Within-run bootstrap intervals and an earlier equivalent TileLang sweep are
available in the reports; they do not cover systematic or across-run uncertainty.
API, calibration contract and reproduction commands · Related implementation: ops-nn PR #12568.
AI/Agent disclosure