Draft
[WIP][New operators] MACAQ W4A4 dense and grouped matmul with CCE #44
[WIP][New operators] MACAQ W4A4 dense and grouped matmul with CCE #44
Draft
Mocchibird Add on 20 days ago
Mocchibird
Mocchibird
20 days ago

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) @ H and 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:

Model reference Perplexity
BF16 21.3111
Ordinary MXFP4 W4A4, scaling algorithm 0 33.4286
MACAQ W4A4 25.4962

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.

MACAQ perplexity compared with ordinary W4A4 and BF16

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.

Actual NPU projection error for MACAQ and both native scaling presets

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, or npu_grouped_matmul for 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.

Workload Regular native/MACAQ Regular wins Prepared native/MACAQ Prepared wins
Dense M=16 1.106× 20/24 1.771× 22/24
Dense M=128 1.008× 18/24 1.441× 19/24
Dense M=4096 0.626× 1/24 0.643× 1/24
Grouped E=4 0.618× 4/24 0.747× 5/24
All 96, equal-case GM 0.810× 43/96 1.052× 47/96

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.

MACAQ latency by workload, including all-case geometric mean

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.

  • 132 correctness/contract cases pass; three level-2 cases are skipped and 96 benchmark cases deselected in the validation run.
  • The full 96-case elapsed sweep passes 576 complete output checks across regular MACAQ, prepared MACAQ and native CANN.
  • Stream/update tests change X, D and routes in place, including alternating-16 and mixed routes, output guards and alias rejection.
  • The global schedule quantizes X once and shares it through one fresh per-call workspace, capped at 64 MiB, inside the same fused kernel. Other shapes retain UB/L1 reuse or streaming.
  • Eligible default-16 grouped tiles inspect routes on device and use 128-row matrix tiles only when uniform, retaining a 16-row fallback inside the same kernel.
  • Additional structured checks cover M=65,536. Historical all-shape maximum-row and >2 GiB expert-bank tests have not all been repeated for this schedule.
  • Ruff check/format pass. All six MACAQ source files match the timing and accuracy metadata.
  • Dense accuracy checks: 36 exact CPU-reference matches, plus nine native algorithm-2 decoding checks.
  • Grouped arithmetic: 24 exact checks on three real down projections with four replicated expert banks. This is not a MoE-model quality evaluation or runtime measurement.
  • Supported-stack CI remains pending; earlier validation reported a driver-version mismatch.

API, calibration contract and reproduction commands · Related implementation: ops-nn PR #12568.

AI/Agent disclosure

likedislike
Merge blocked
Mocchibird Mocchibird
20 days ago Created pull request, commit 8cf51cc6
CANN-robot
CANN-robotCollaborators
20 days ago Comments:

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.

likedislike
CANN-robotCANN-robotCollaborators
20 days ago Added label: cann-cla/yes
CANN-robot
CANN-robotCollaborators
20 days ago Comments:

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 /approve or /lgtm
  • Commenting /approve implies 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. 👍

likedislike
CANN-robotCANN-robotCollaborators
20 days ago Set mywaaagh_admin,songchangxia,tanghaoran7,ChaoyangJi,kevin365,shentukangjun,monologue815,jiang-lirui,Allan_Yu,jiawei_zhuang,liudan12,Larry_liuyi,mabing1118 as approver
127 messages are collapsed here。 See more
CANN-robotCANN-robotCollaborators
5 days ago Added label: ci-pipeline-running
CANN-robotCANN-robotCollaborators
5 days ago Removed label: ci-pipeline-running
CANN-robotCANN-robotCollaborators
5 days ago Added label: ci-pipeline-failed
CANN-robot
CANN-robotCollaborators
5 days ago Comments:

流水线分析报告 — 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

likedislike
Mocchibird Mocchibird
5 hours ago Modify description of pull request