[WIP][New operators] MACAQ W4A4 dense and grouped matmul with CCE-ops-tilelang-AtomGit

Adocs/benchmarks/macaq_20261027/README.md+139-0

@@ -0,0 +1,139 @@
	1	+# TileLang MACAQ (regular API): performance against native CANN
	2	+
	3	+Accuracy at competitive small-batch latency: the separate Qwen3-0.6B CPU-reference pilot reduces perplexity from 33.4286 (ordinary W4A4) to 25.4962 (MACAQ), a 23.7% reduction. Actual NPU projection NMSE improves against both native scaling presets on all nine fixed projections. [Accuracy tables, plots and evaluation limits](accuracy.md). The model accuracy and kernel latency studies are separate; large-case latency losses remain below.
	4	+
	5	+These results measure the modified working tree whose source hashes are recorded in the metadata. The base-source study remains available in [the earlier report](submitted-source-before-optimization.md). Kernel execution remains fully fused in one mixed Vector/Cube launch.
	6	+
	7	+The [MACAQ prepared-plan report](tilelang_prepared-performance.md) separately measures static validation and layout views prepared once. It keeps output allocation inside each call and explicitly declares the stronger 128-row routing contract for grouped fixtures. The regular MACAQ API retains its 16-row contract.
	8	+
	9	+The regular TileLang MACAQ grouped API retains its original 16-row constant-routing contract. Eligible shapes inspect routes on device and use 128-row matrix tiles when uniform, otherwise retaining 16-row matrix tiles.
	10	+
	11	+Selected TileLang shapes share quantized activations through one per-call GM workspace, limited to 64 MiB, inside ONE mixed Vector/Cube kernel. Other shapes retain UB/L1 reuse or streaming. The workspace is freshly allocated inside each call; its allocation and all temporary memory count in these timings. The MACAQ prepared grouped path uses an explicit 128-row contract. Eligible MACAQ default-16 grouped shapes use device-side uniformity checks and a 16-row fallback inside the same kernel.
	12	+
	13	+Independent baseline: native `npu_dynamic_mx_quant` followed by `npu_quant_matmul`, or `npu_grouped_matmul` for grouped workloads. The activation quantizer uses block32 E2M1, round-to-even and `scale_alg=0`. The baseline omits the MACAQ diagonal and Hadamard transform. MACAQ is our contribution and its split implementations are ablations. Each algorithm is checked against its own CPU reference.
	14	+
	15	+The source commits in metadata identify base commits; file hashes pin the modified working tree used in this experiment. File hashes identify the measured MACAQ sources included in this PR.
	16	+
	17	+Ascend 950, CANN 9.3 weekly, torch 2.9.0 / torch_npu 2.9.0.post2, TileLang 0.1.15. Each case uses 20 warmups and 15 rotated trials of 20 eager calls, synchronized around each trial bracket. Two distinct activation buffers with identical deterministic values alternate between calls. All arms allocate outputs inside every timed call; intermediate allocations count. Static packed weights/layout views are prepared outside calls. ops-nn accepts raw D and normalizes it inside the fused kernel; TileLang accepts its documented prepared D and scale pairs. Compilation and calibration are excluded. No explicit cache flush is used. This measures repeated public-API execution, not single-request latency, graph replay or model throughput. The MACAQ and vendor arms for each case use the same NPU under the same scheduler lock.
	18	+
	19	+All 384 full-output checks passed at BF16 rtol=atol=1/128. The 96 cases cover all 24 supported (K,N) pairs at dense M=16,128,4096 and four-expert grouped workloads. Grouped rows use one ops-nn alignment unit per expert; alignment padding is included identically in both arms. Win counts compare sample medians and do not establish statistical significance.
	20	+
	21	+![All elapsed ratios](https://raw.gitcode.com/Mocchibird/ops-tilelang/blobs/6a429fc64bd8dd9391a9ee926294cad4e1e84163/docs/benchmarks/macaq_20261027/tilelang_pr-elapsed.png)
	22	+
	23	+\| Workload \| Vendor / MACAQ geomean \| Wins \| Ratio range \|
	24	+\|---\|---:\|---:\|---:\|
	25	+\| Dense M=16 \| 1.106x \| 20/24 \| 0.844–1.218x \|
	26	+\| Dense M=128 \| 1.008x \| 18/24 \| 0.473–1.209x \|
	27	+\| Dense M=4096 \| 0.626x \| 1/24 \| 0.417–1.182x \|
	28	+\| Grouped E=4 \| 0.618x \| 4/24 \| 0.407–1.248x \|
	29	+
	30	+Overall: 43/96 median wins; 0.810x geometric mean vendor/MACAQ. The paired-trial within-run bootstrap 95% interval is 0.809–0.816x; it does not include across-run or systematic uncertainty. An equal-call mixture (one call per configuration) has total vendor/MACAQ time ratio 0.672x. This weights time spent on larger cases more heavily than the equal-case geometric mean of ratios. Ratios below one are losses and remain in the table and plot. Aggregate parity does not mean that every shape matches the vendor. No claim of an all-shape speedup is made. Accuracy is measured in a separate [held-out study](accuracy.md); these kernel timings are not model-level speed-and-accuracy results.
	31	+
	32	+## Complete table
	33	+
	34	+\| Case \| Workload \| M \| K \| N \| MACAQ µs \| Vendor µs \| Vendor / MACAQ \|
	35	+\|---:\|---\|---:\|---:\|---:\|---:\|---:\|---:\|
	36	+\| 0 \| dense \| 16 \| 512 \| 2048 \| 28.906 \| 34.247 \| 1.185x \|
	37	+\| 1 \| dense \| 128 \| 512 \| 2048 \| 28.336 \| 34.272 \| 1.209x \|
	38	+\| 2 \| dense \| 4096 \| 512 \| 2048 \| 29.327 \| 34.669 \| 1.182x \|
	39	+\| 3 \| grouped \| 1536 \| 512 \| 2048 \| 30.027 \| 37.477 \| 1.248x \|
	40	+\| 4 \| dense \| 16 \| 768 \| 2560 \| 27.873 \| 33.840 \| 1.214x \|
	41	+\| 5 \| dense \| 128 \| 768 \| 2560 \| 27.752 \| 33.479 \| 1.206x \|
	42	+\| 6 \| dense \| 4096 \| 768 \| 2560 \| 35.956 \| 33.588 \| 0.934x \|
	43	+\| 7 \| grouped \| 512 \| 768 \| 2560 \| 30.206 \| 35.623 \| 1.179x \|
	44	+\| 8 \| dense \| 16 \| 1280 \| 4096 \| 28.674 \| 33.797 \| 1.179x \|
	45	+\| 9 \| dense \| 128 \| 1280 \| 4096 \| 28.050 \| 32.691 \| 1.165x \|
	46	+\| 10 \| dense \| 4096 \| 1280 \| 4096 \| 58.612 \| 37.493 \| 0.640x \|
	47	+\| 11 \| grouped \| 512 \| 1280 \| 4096 \| 50.861 \| 35.697 \| 0.702x \|
	48	+\| 12 \| dense \| 16 \| 1536 \| 2048 \| 28.033 \| 33.034 \| 1.178x \|
	49	+\| 13 \| dense \| 128 \| 1536 \| 2048 \| 28.199 \| 33.403 \| 1.185x \|
	50	+\| 14 \| dense \| 4096 \| 1536 \| 2048 \| 51.185 \| 34.096 \| 0.666x \|
	51	+\| 15 \| grouped \| 512 \| 1536 \| 2048 \| 31.901 \| 37.154 \| 1.165x \|
	52	+\| 16 \| dense \| 16 \| 2048 \| 1024 \| 29.244 \| 33.776 \| 1.155x \|
	53	+\| 17 \| dense \| 128 \| 2048 \| 1024 \| 28.702 \| 33.332 \| 1.161x \|
	54	+\| 18 \| dense \| 4096 \| 2048 \| 1024 \| 54.003 \| 35.226 \| 0.652x \|
	55	+\| 19 \| grouped \| 1536 \| 2048 \| 1024 \| 47.985 \| 36.129 \| 0.753x \|
	56	+\| 20 \| dense \| 16 \| 2048 \| 3072 \| 28.932 \| 33.616 \| 1.162x \|
	57	+\| 21 \| dense \| 128 \| 2048 \| 3072 \| 28.373 \| 33.101 \| 1.167x \|
	58	+\| 22 \| dense \| 4096 \| 2048 \| 3072 \| 79.102 \| 43.982 \| 0.556x \|
	59	+\| 23 \| grouped \| 1536 \| 2048 \| 3072 \| 53.876 \| 38.886 \| 0.722x \|
	60	+\| 24 \| dense \| 16 \| 2048 \| 4096 \| 28.267 \| 33.059 \| 1.170x \|
	61	+\| 25 \| dense \| 128 \| 2048 \| 4096 \| 32.984 \| 34.361 \| 1.042x \|
	62	+\| 26 \| dense \| 4096 \| 2048 \| 4096 \| 89.129 \| 54.672 \| 0.613x \|
	63	+\| 27 \| grouped \| 1536 \| 2048 \| 4096 \| 52.184 \| 37.263 \| 0.714x \|
	64	+\| 28 \| dense \| 16 \| 2048 \| 6144 \| 27.661 \| 33.346 \| 1.206x \|
	65	+\| 29 \| dense \| 128 \| 2048 \| 6144 \| 31.714 \| 33.414 \| 1.054x \|
	66	+\| 30 \| dense \| 4096 \| 2048 \| 6144 \| 118.544 \| 73.504 \| 0.620x \|
	67	+\| 31 \| grouped \| 1536 \| 2048 \| 6144 \| 67.392 \| 37.260 \| 0.553x \|
	68	+\| 32 \| dense \| 16 \| 2048 \| 7168 \| 28.268 \| 33.392 \| 1.181x \|
	69	+\| 33 \| dense \| 128 \| 2048 \| 7168 \| 32.252 \| 33.752 \| 1.047x \|
	70	+\| 34 \| dense \| 4096 \| 2048 \| 7168 \| 127.097 \| 83.047 \| 0.653x \|
	71	+\| 35 \| grouped \| 1536 \| 2048 \| 7168 \| 82.907 \| 39.215 \| 0.473x \|
	72	+\| 36 \| dense \| 16 \| 2048 \| 8192 \| 27.487 \| 32.968 \| 1.199x \|
	73	+\| 37 \| dense \| 128 \| 2048 \| 8192 \| 31.550 \| 33.281 \| 1.055x \|
	74	+\| 38 \| dense \| 4096 \| 2048 \| 8192 \| 143.204 \| 98.037 \| 0.685x \|
	75	+\| 39 \| grouped \| 1536 \| 2048 \| 8192 \| 86.732 \| 44.730 \| 0.516x \|
	76	+\| 40 \| dense \| 16 \| 2048 \| 12288 \| 27.889 \| 33.974 \| 1.218x \|
	77	+\| 41 \| dense \| 128 \| 2048 \| 12288 \| 31.932 \| 34.483 \| 1.080x \|
	78	+\| 42 \| dense \| 4096 \| 2048 \| 12288 \| 201.651 \| 132.114 \| 0.655x \|
	79	+\| 43 \| grouped \| 1536 \| 2048 \| 12288 \| 125.492 \| 60.836 \| 0.485x \|
	80	+\| 44 \| dense \| 16 \| 2560 \| 1280 \| 28.407 \| 33.938 \| 1.195x \|
	81	+\| 45 \| dense \| 128 \| 2560 \| 1280 \| 28.243 \| 33.877 \| 1.199x \|
	82	+\| 46 \| dense \| 4096 \| 2560 \| 1280 \| 70.410 \| 35.338 \| 0.502x \|
	83	+\| 47 \| grouped \| 512 \| 2560 \| 1280 \| 31.017 \| 37.829 \| 1.220x \|
	84	+\| 48 \| dense \| 16 \| 3072 \| 7168 \| 28.665 \| 34.154 \| 1.191x \|
	85	+\| 49 \| dense \| 128 \| 3072 \| 7168 \| 32.451 \| 34.687 \| 1.069x \|
	86	+\| 50 \| dense \| 4096 \| 3072 \| 7168 \| 181.774 \| 119.005 \| 0.655x \|
	87	+\| 51 \| grouped \| 512 \| 3072 \| 7168 \| 48.432 \| 38.047 \| 0.786x \|
	88	+\| 52 \| dense \| 16 \| 4096 \| 2560 \| 28.742 \| 34.096 \| 1.186x \|
	89	+\| 53 \| dense \| 128 \| 4096 \| 2560 \| 32.904 \| 34.450 \| 1.047x \|
	90	+\| 54 \| dense \| 4096 \| 4096 \| 2560 \| 150.233 \| 64.050 \| 0.426x \|
	91	+\| 55 \| grouped \| 1536 \| 4096 \| 2560 \| 92.537 \| 37.950 \| 0.410x \|
	92	+\| 56 \| dense \| 16 \| 4096 \| 4096 \| 28.686 \| 34.670 \| 1.209x \|
	93	+\| 57 \| dense \| 128 \| 4096 \| 4096 \| 32.420 \| 34.188 \| 1.055x \|
	94	+\| 58 \| dense \| 4096 \| 4096 \| 4096 \| 179.473 \| 99.019 \| 0.552x \|
	95	+\| 59 \| grouped \| 1536 \| 4096 \| 4096 \| 101.371 \| 48.002 \| 0.474x \|
	96	+\| 60 \| dense \| 16 \| 4096 \| 24576 \| 33.139 \| 35.736 \| 1.078x \|
	97	+\| 61 \| dense \| 128 \| 4096 \| 24576 \| 39.163 \| 35.952 \| 0.918x \|
	98	+\| 62 \| dense \| 4096 \| 4096 \| 24576 \| 764.154 \| 587.611 \| 0.769x \|
	99	+\| 63 \| grouped \| 1536 \| 4096 \| 24576 \| 515.314 \| 216.182 \| 0.420x \|
	100	+\| 64 \| dense \| 16 \| 5120 \| 34816 \| 50.669 \| 47.055 \| 0.929x \|
	101	+\| 65 \| dense \| 128 \| 5120 \| 34816 \| 65.696 \| 47.529 \| 0.723x \|
	102	+\| 66 \| dense \| 4096 \| 5120 \| 34816 \| 1296.152 \| 1081.125 \| 0.834x \|
	103	+\| 67 \| grouped \| 512 \| 5120 \| 34816 \| 348.464 \| 158.788 \| 0.456x \|
	104	+\| 68 \| dense \| 16 \| 6144 \| 2048 \| 31.901 \| 33.974 \| 1.065x \|
	105	+\| 69 \| dense \| 128 \| 6144 \| 2048 \| 32.381 \| 34.738 \| 1.073x \|
	106	+\| 70 \| dense \| 4096 \| 6144 \| 2048 \| 191.286 \| 79.774 \| 0.417x \|
	107	+\| 71 \| grouped \| 512 \| 6144 \| 2048 \| 42.998 \| 37.418 \| 0.870x \|
	108	+\| 72 \| dense \| 16 \| 6144 \| 4096 \| 31.200 \| 28.340 \| 0.908x \|
	109	+\| 73 \| dense \| 128 \| 6144 \| 4096 \| 32.796 \| 29.135 \| 0.888x \|
	110	+\| 74 \| dense \| 4096 \| 6144 \| 4096 \| 261.207 \| 145.547 \| 0.557x \|
	111	+\| 75 \| grouped \| 512 \| 6144 \| 4096 \| 54.634 \| 39.214 \| 0.718x \|
	112	+\| 76 \| dense \| 16 \| 7168 \| 4096 \| 32.138 \| 29.145 \| 0.907x \|
	113	+\| 77 \| dense \| 128 \| 7168 \| 4096 \| 31.967 \| 28.706 \| 0.898x \|
	114	+\| 78 \| dense \| 4096 \| 7168 \| 4096 \| 287.589 \| 166.680 \| 0.580x \|
	115	+\| 79 \| grouped \| 1536 \| 7168 \| 4096 \| 171.462 \| 73.326 \| 0.428x \|
	116	+\| 80 \| dense \| 16 \| 7168 \| 6144 \| 33.025 \| 34.613 \| 1.048x \|
	117	+\| 81 \| dense \| 128 \| 7168 \| 6144 \| 33.849 \| 34.336 \| 1.014x \|
	118	+\| 82 \| dense \| 4096 \| 7168 \| 6144 \| 384.464 \| 235.441 \| 0.612x \|
	119	+\| 83 \| grouped \| 1536 \| 7168 \| 6144 \| 242.461 \| 100.984 \| 0.416x \|
	120	+\| 84 \| dense \| 16 \| 8192 \| 4096 \| 32.772 \| 34.564 \| 1.055x \|
	121	+\| 85 \| dense \| 128 \| 8192 \| 4096 \| 32.404 \| 33.733 \| 1.041x \|
	122	+\| 86 \| dense \| 4096 \| 8192 \| 4096 \| 346.262 \| 188.090 \| 0.543x \|
	123	+\| 87 \| grouped \| 1536 \| 8192 \| 4096 \| 201.769 \| 82.204 \| 0.407x \|
	124	+\| 88 \| dense \| 16 \| 12288 \| 4096 \| 33.529 \| 34.635 \| 1.033x \|
	125	+\| 89 \| dense \| 128 \| 12288 \| 4096 \| 40.383 \| 34.131 \| 0.845x \|
	126	+\| 90 \| dense \| 4096 \| 12288 \| 4096 \| 543.457 \| 286.346 \| 0.527x \|
	127	+\| 91 \| grouped \| 512 \| 12288 \| 4096 \| 116.066 \| 54.986 \| 0.474x \|
	128	+\| 92 \| dense \| 16 \| 17408 \| 5120 \| 41.120 \| 34.704 \| 0.844x \|
	129	+\| 93 \| dense \| 128 \| 17408 \| 5120 \| 72.084 \| 34.126 \| 0.473x \|
	130	+\| 94 \| dense \| 4096 \| 17408 \| 5120 \| 887.334 \| 530.846 \| 0.598x \|
	131	+\| 95 \| grouped \| 512 \| 17408 \| 5120 \| 227.851 \| 93.099 \| 0.409x \|
	132	+
	133	+## Repeatability
	134	+
	135	+An earlier full sweep of the same kernel computations gave 0.777x geometric mean and 19/96 median wins; the current run gives 0.810x and 43/96. Only import ordering was cleaned up between these runs; kernel/test function ASTs were verified unchanged. These differences therefore are not attributed to an algorithm change. Host and device conditions vary across runs. Within-run bootstrap intervals do not capture that variation, and ratios close to one should be interpreted cautiously.
	136	+
	137	+[Earlier same-computation raw sweep](tilelang_pr-earlier-equivalent.json) · [Earlier source/runtime metadata](tilelang_pr-earlier-equivalent.metadata.json).
	138	+
	139	+Raw benchmark data is stored as standard JSON arrays. Files converted from CSV retain the header as the first row and preserve every cell as a string; files converted from JSON Lines retain each original record object. Chart images link to immutable, previously published Git blobs.

Adocs/benchmarks/macaq_20261027/accuracy-grouped-dynamic.json+188-0

@@ -0,0 +1,188 @@
	1	+[
	2	+ {
	3	+ "fixture": "accuracy-fixture-model_layers_0_mlp_down_proj.pt",
	4	+ "fixture_sha256": "4f73bdf7d8a927172fea0866e27cbe850c894deb329930bd08b99aab6cff444c",
	5	+ "m": 1024,
	6	+ "k": 3072,
	7	+ "n": 1024,
	8	+ "tilelang_padded_n": 2048,
	9	+ "opsnn_padded_n": 7168,
	10	+ "experts": 4,
	11	+ "outputs": {
	12	+ "native_grouped_alg0": {
	13	+ "nmse": 0.030516970236956067,
	14	+ "nrmse": 0.17469107085640087,
	15	+ "max_abs_error": 0.25
	16	+ },
	17	+ "native_grouped_alg2": {
	18	+ "nmse": 0.031153790077059946,
	19	+ "nrmse": 0.1765043627706124,
	20	+ "max_abs_error": 0.19921875
	21	+ },
	22	+ "runs_standard": {
	23	+ "nmse": 0.018508736055048453,
	24	+ "nrmse": 0.13604681567404822,
	25	+ "max_abs_error": 0.248046875
	26	+ },
	27	+ "runs_prepared16": {
	28	+ "nmse": 0.018508736055048453,
	29	+ "nrmse": 0.13604681567404822,
	30	+ "max_abs_error": 0.248046875
	31	+ },
	32	+ "runs_prepared128": {
	33	+ "nmse": 0.018508736055048453,
	34	+ "nrmse": 0.13604681567404822,
	35	+ "max_abs_error": 0.248046875
	36	+ },
	37	+ "alternating16_standard": {
	38	+ "nmse": 0.018508736055048453,
	39	+ "nrmse": 0.13604681567404822,
	40	+ "max_abs_error": 0.248046875
	41	+ },
	42	+ "alternating16_prepared16": {
	43	+ "nmse": 0.018508736055048453,
	44	+ "nrmse": 0.13604681567404822,
	45	+ "max_abs_error": 0.248046875
	46	+ },
	47	+ "runs_opsnn": {
	48	+ "nmse": 0.018508736055048453,
	49	+ "nrmse": 0.13604681567404822,
	50	+ "max_abs_error": 0.248046875
	51	+ }
	52	+ },
	53	+ "cpu_max_abs": {
	54	+ "native_grouped_alg0": 0.0,
	55	+ "native_grouped_alg2": 0.0,
	56	+ "runs_standard": 0.0,
	57	+ "runs_prepared16": 0.0,
	58	+ "runs_prepared128": 0.0,
	59	+ "alternating16_standard": 0.0,
	60	+ "alternating16_prepared16": 0.0,
	61	+ "runs_opsnn": 0.0
	62	+ }
	63	+ },
	64	+ {
	65	+ "fixture": "accuracy-fixture-model_layers_14_mlp_down_proj.pt",
	66	+ "fixture_sha256": "06327450f71e490385a6137915d300273b901ed344017c98b3cfaf8a7061b579",
	67	+ "m": 1024,
	68	+ "k": 3072,
	69	+ "n": 1024,
	70	+ "tilelang_padded_n": 2048,
	71	+ "opsnn_padded_n": 7168,
	72	+ "experts": 4,
	73	+ "outputs": {
	74	+ "native_grouped_alg0": {
	75	+ "nmse": 0.03410957884917448,
	76	+ "nrmse": 0.18468778749331122,
	77	+ "max_abs_error": 1.5625
	78	+ },
	79	+ "native_grouped_alg2": {
	80	+ "nmse": 0.03799854210203062,
	81	+ "nrmse": 0.19493214743092177,
	82	+ "max_abs_error": 1.8125
	83	+ },
	84	+ "runs_standard": {
	85	+ "nmse": 0.021642116309140973,
	86	+ "nrmse": 0.1471125973842518,
	87	+ "max_abs_error": 0.7265625
	88	+ },
	89	+ "runs_prepared16": {
	90	+ "nmse": 0.021642116309140973,
	91	+ "nrmse": 0.1471125973842518,
	92	+ "max_abs_error": 0.7265625
	93	+ },
	94	+ "runs_prepared128": {
	95	+ "nmse": 0.021642116309140973,
	96	+ "nrmse": 0.1471125973842518,
	97	+ "max_abs_error": 0.7265625
	98	+ },
	99	+ "alternating16_standard": {
	100	+ "nmse": 0.021642116309140973,
	101	+ "nrmse": 0.1471125973842518,
	102	+ "max_abs_error": 0.7265625
	103	+ },
	104	+ "alternating16_prepared16": {
	105	+ "nmse": 0.021642116309140973,
	106	+ "nrmse": 0.1471125973842518,
	107	+ "max_abs_error": 0.7265625
	108	+ },
	109	+ "runs_opsnn": {
	110	+ "nmse": 0.021642116309140973,
	111	+ "nrmse": 0.1471125973842518,
	112	+ "max_abs_error": 0.7265625
	113	+ }
	114	+ },
	115	+ "cpu_max_abs": {
	116	+ "native_grouped_alg0": 0.0,
	117	+ "native_grouped_alg2": 0.0,
	118	+ "runs_standard": 0.0,
	119	+ "runs_prepared16": 0.0,
	120	+ "runs_prepared128": 0.0,
	121	+ "alternating16_standard": 0.0,
	122	+ "alternating16_prepared16": 0.0,
	123	+ "runs_opsnn": 0.0
	124	+ }
	125	+ },
	126	+ {
	127	+ "fixture": "accuracy-fixture-model_layers_27_mlp_down_proj.pt",
	128	+ "fixture_sha256": "18f102d83016adb97aa542df1515be28ab7acea4059b91fcd8acb22427382c25",
	129	+ "m": 1024,
	130	+ "k": 3072,
	131	+ "n": 1024,
	132	+ "tilelang_padded_n": 2048,
	133	+ "opsnn_padded_n": 7168,
	134	+ "experts": 4,
	135	+ "outputs": {
	136	+ "native_grouped_alg0": {
	137	+ "nmse": 0.021753115704479336,
	138	+ "nrmse": 0.14748937488673322,
	139	+ "max_abs_error": 720.0
	140	+ },
	141	+ "native_grouped_alg2": {
	142	+ "nmse": 0.019562300238837857,
	143	+ "nrmse": 0.13986529318897473,
	144	+ "max_abs_error": 192.0
	145	+ },
	146	+ "runs_standard": {
	147	+ "nmse": 0.0047220473550336605,
	148	+ "nrmse": 0.06871715473616222,
	149	+ "max_abs_error": 288.0
	150	+ },
	151	+ "runs_prepared16": {
	152	+ "nmse": 0.0047220473550336605,
	153	+ "nrmse": 0.06871715473616222,
	154	+ "max_abs_error": 288.0
	155	+ },
	156	+ "runs_prepared128": {
	157	+ "nmse": 0.0047220473550336605,
	158	+ "nrmse": 0.06871715473616222,
	159	+ "max_abs_error": 288.0
	160	+ },
	161	+ "alternating16_standard": {
	162	+ "nmse": 0.0047220473550336605,
	163	+ "nrmse": 0.06871715473616222,
	164	+ "max_abs_error": 288.0
	165	+ },
	166	+ "alternating16_prepared16": {
	167	+ "nmse": 0.0047220473550336605,
	168	+ "nrmse": 0.06871715473616222,
	169	+ "max_abs_error": 288.0
	170	+ },
	171	+ "runs_opsnn": {
	172	+ "nmse": 0.0047220473550336605,
	173	+ "nrmse": 0.06871715473616222,
	174	+ "max_abs_error": 288.0
	175	+ }
	176	+ },
	177	+ "cpu_max_abs": {
	178	+ "native_grouped_alg0": 0.0,
	179	+ "native_grouped_alg2": 0.0,
	180	+ "runs_standard": 0.0,
	181	+ "runs_prepared16": 0.0,
	182	+ "runs_prepared128": 0.0,
	183	+ "alternating16_standard": 0.0,
	184	+ "alternating16_prepared16": 0.0,
	185	+ "runs_opsnn": 0.0
	186	+ }
	187	+ }
	188	+]

Adocs/benchmarks/macaq_20261027/accuracy-grouped-dynamic.metadata.json+24-0

@@ -0,0 +1,24 @@
	1	+{
	2	+ "checks": 24,
	3	+ "fixtures": 3,
	4	+ "source_sha256": {
	5	+ "src/cann_ops_tilelang/__init__.py": "a81e80bbf9fd868d5109d2ab725dec9afe16d0a346ccad7809e55fb6dddfd2a7",
	6	+ "src/cann_ops_tilelang/matmul/__init__.py": "13c8e8bdf2323d8787bcdf944464e1975343b35fe95a95e40d8d5612cec5a1eb",
	7	+ "src/cann_ops_tilelang/matmul/macaq_w4a4_matmul.py": "039334d011a6c5714be28066c1dc12a00868ac4a9fea1c640ed974a61bf830b8",
	8	+ "src/cann_ops_tilelang/matmul/macaq_w4a4_matmul_kernel.py": "fb33d479bb11db9ca3617c029d8c40dc7135af1baf0571346b6add11884bd777",
	9	+ "src/cann_ops_tilelang/matmul/macaq_w4a4_global_kernel.py": "128169df74198a73da97292edc8d9aefe9caee6cfe70c3c2b2b88fe024c0d65c",
	10	+ "src/cann_ops_tilelang/matmul/macaq_w4a4_dynamic_group_kernel.py": "c4a61194f87466de63e30b5f412146ce08f14eab7c5a1887cf83a2328b39212a",
	11	+ "tests/matmul/test_macaq_w4a4_matmul.py": "4af44fcb180beade6cbba3fe49c156603c69c7c8a4a652a52f134e01abf37293",
	12	+ "tests/test_macaq_contract.py": "fd971a2be888b2a2387a39e6a75a5f4765bb023635301a20bdc848a866eba978"
	13	+ },
	14	+ "fixture_sha256": {
	15	+ "accuracy-fixture-model_layers_0_mlp_down_proj.pt": "4f73bdf7d8a927172fea0866e27cbe850c894deb329930bd08b99aab6cff444c",
	16	+ "accuracy-fixture-model_layers_14_mlp_down_proj.pt": "06327450f71e490385a6137915d300273b901ed344017c98b3cfaf8a7061b579",
	17	+ "accuracy-fixture-model_layers_27_mlp_down_proj.pt": "18f102d83016adb97aa542df1515be28ab7acea4059b91fcd8acb22427382c25"
	18	+ },
	19	+ "harness_sha256": "dc688dc4f2786a3b9005fe2ea5cd2660d4100d732adeaa29d3c644e0353d955b",
	20	+ "scope": "Grouped arithmetic check on three preselected real down-projection fixtures. Four replicated expert banks; Nzero-padding only, sliced original columns. No MoE-model evaluation or runtime claim. Alternating16 routes valid only for standard/prepared16; prepared128 andopsnn usecontiguous256-row runs.",
	21	+ "torch": "2.9.0+cpu",
	22	+ "torch_npu": "2.9.0.post2",
	23	+ "device": "Ascend950DT_9592"
	24	+}

Adocs/benchmarks/macaq_20261027/accuracy-layer-study.json+200-0

@@ -0,0 +1,200 @@
	1	+[
	2	+ {
	3	+ "layer": "model.layers.0.self_attn.q_proj",
	4	+ "m": 1024,
	5	+ "k": 1024,
	6	+ "n": 2048,
	7	+ "plain_w4a4": {
	8	+ "nmse": 0.012961809523403645,
	9	+ "nrmse": 0.11384994301010275,
	10	+ "sqnr_db": 18.873343648873842,
	11	+ "cosine": 0.9924591183662415,
	12	+ "max_abs_error": 1.015625
	13	+ },
	14	+ "macaq_w4a4": {
	15	+ "nmse": 0.005795283708721399,
	16	+ "nrmse": 0.07612676079225622,
	17	+ "sqnr_db": 22.369252982480702,
	18	+ "cosine": 0.9948984980583191,
	19	+ "max_abs_error": 0.46875
	20	+ },
	21	+ "elapsed_cpu_seconds": 0.3079274800256826,
	22	+ "nmse_reduction_factor": 2.236613455851565
	23	+ },
	24	+ {
	25	+ "layer": "model.layers.0.mlp.up_proj",
	26	+ "m": 1024,
	27	+ "k": 1024,
	28	+ "n": 3072,
	29	+ "plain_w4a4": {
	30	+ "nmse": 0.03272208198904991,
	31	+ "nrmse": 0.18089245973519713,
	32	+ "sqnr_db": 14.8515907151692,
	33	+ "cosine": 0.9813424348831177,
	34	+ "max_abs_error": 0.640625
	35	+ },
	36	+ "macaq_w4a4": {
	37	+ "nmse": 0.018918711692094803,
	38	+ "nrmse": 0.1375453077792725,
	39	+ "sqnr_db": 17.231084410877195,
	40	+ "cosine": 0.9877645969390869,
	41	+ "max_abs_error": 0.3671875
	42	+ },
	43	+ "elapsed_cpu_seconds": 0.5210358299664222,
	44	+ "nmse_reduction_factor": 1.7296147074709562
	45	+ },
	46	+ {
	47	+ "layer": "model.layers.0.mlp.down_proj",
	48	+ "m": 1024,
	49	+ "k": 3072,
	50	+ "n": 1024,
	51	+ "plain_w4a4": {
	52	+ "nmse": 0.0305169690400362,
	53	+ "nrmse": 0.17469106743058216,
	54	+ "sqnr_db": 15.154586029109385,
	55	+ "cosine": 0.9849362373352051,
	56	+ "max_abs_error": 0.25
	57	+ },
	58	+ "macaq_w4a4": {
	59	+ "nmse": 0.018508734181523323,
	60	+ "nrmse": 0.13604680878845826,
	61	+ "sqnr_db": 17.32623281775322,
	62	+ "cosine": 0.9904028177261353,
	63	+ "max_abs_error": 0.248046875
	64	+ },
	65	+ "elapsed_cpu_seconds": 0.8360411199973896,
	66	+ "nmse_reduction_factor": 1.6487874719439386
	67	+ },
	68	+ {
	69	+ "layer": "model.layers.14.self_attn.q_proj",
	70	+ "m": 1024,
	71	+ "k": 1024,
	72	+ "n": 2048,
	73	+ "plain_w4a4": {
	74	+ "nmse": 0.014667452313005924,
	75	+ "nrmse": 0.12110925775103208,
	76	+ "sqnr_db": 18.336453150926808,
	77	+ "cosine": 0.9908798933029175,
	78	+ "max_abs_error": 2.640625
	79	+ },
	80	+ "macaq_w4a4": {
	81	+ "nmse": 0.007576686330139637,
	82	+ "nrmse": 0.08704416310206926,
	83	+ "sqnr_db": 21.20520691886017,
	84	+ "cosine": 0.9940814971923828,
	85	+ "max_abs_error": 2.1171875
	86	+ },
	87	+ "elapsed_cpu_seconds": 0.24619967996841297,
	88	+ "nmse_reduction_factor": 1.9358663766585684
	89	+ },
	90	+ {
	91	+ "layer": "model.layers.14.mlp.up_proj",
	92	+ "m": 1024,
	93	+ "k": 1024,
	94	+ "n": 3072,
	95	+ "plain_w4a4": {
	96	+ "nmse": 0.027274655178189278,
	97	+ "nrmse": 0.16515040168945783,
	98	+ "sqnr_db": 15.642407313044702,
	99	+ "cosine": 0.9838981628417969,
	100	+ "max_abs_error": 1.109375
	101	+ },
	102	+ "macaq_w4a4": {
	103	+ "nmse": 0.01696912758052349,
	104	+ "nrmse": 0.13026560398095688,
	105	+ "sqnr_db": 17.703404851251104,
	106	+ "cosine": 0.98872971534729,
	107	+ "max_abs_error": 0.671875
	108	+ },
	109	+ "elapsed_cpu_seconds": 0.4196922399569303,
	110	+ "nmse_reduction_factor": 1.6073103964103654
	111	+ },
	112	+ {
	113	+ "layer": "model.layers.14.mlp.down_proj",
	114	+ "m": 1024,
	115	+ "k": 3072,
	116	+ "n": 1024,
	117	+ "plain_w4a4": {
	118	+ "nmse": 0.034109581261873245,
	119	+ "nrmse": 0.1846877940251419,
	120	+ "sqnr_db": 14.671236120723465,
	121	+ "cosine": 0.9827119708061218,
	122	+ "max_abs_error": 1.5625
	123	+ },
	124	+ "macaq_w4a4": {
	125	+ "nmse": 0.021642116829752922,
	126	+ "nrmse": 0.1471125991536854,
	127	+ "sqnr_db": 16.647002628575503,
	128	+ "cosine": 0.988767147064209,
	129	+ "max_abs_error": 0.7265625
	130	+ },
	131	+ "elapsed_cpu_seconds": 0.7801028800313361,
	132	+ "nmse_reduction_factor": 1.5760741673374774
	133	+ },
	134	+ {
	135	+ "layer": "model.layers.27.self_attn.q_proj",
	136	+ "m": 1024,
	137	+ "k": 1024,
	138	+ "n": 2048,
	139	+ "plain_w4a4": {
	140	+ "nmse": 0.014878039248287678,
	141	+ "nrmse": 0.12197556824334814,
	142	+ "sqnr_db": 18.274542999566947,
	143	+ "cosine": 0.9907159805297852,
	144	+ "max_abs_error": 32.5
	145	+ },
	146	+ "macaq_w4a4": {
	147	+ "nmse": 0.007906630635261536,
	148	+ "nrmse": 0.08891923658726235,
	149	+ "sqnr_db": 21.020085491517307,
	150	+ "cosine": 0.9939277768135071,
	151	+ "max_abs_error": 17.75
	152	+ },
	153	+ "elapsed_cpu_seconds": 0.2413272900157608,
	154	+ "nmse_reduction_factor": 1.8817167431516852
	155	+ },
	156	+ {
	157	+ "layer": "model.layers.27.mlp.up_proj",
	158	+ "m": 1024,
	159	+ "k": 1024,
	160	+ "n": 3072,
	161	+ "plain_w4a4": {
	162	+ "nmse": 0.02178797498345375,
	163	+ "nrmse": 0.14760750314077448,
	164	+ "sqnr_db": 16.61783132047438,
	165	+ "cosine": 0.9861671328544617,
	166	+ "max_abs_error": 53.0
	167	+ },
	168	+ "macaq_w4a4": {
	169	+ "nmse": 0.008336487226188183,
	170	+ "nrmse": 0.09130436586597698,
	171	+ "sqnr_db": 20.790169109517755,
	172	+ "cosine": 0.9920082092285156,
	173	+ "max_abs_error": 39.0
	174	+ },
	175	+ "elapsed_cpu_seconds": 0.39534480997826904,
	176	+ "nmse_reduction_factor": 2.6135678484589
	177	+ },
	178	+ {
	179	+ "layer": "model.layers.27.mlp.down_proj",
	180	+ "m": 1024,
	181	+ "k": 3072,
	182	+ "n": 1024,
	183	+ "plain_w4a4": {
	184	+ "nmse": 0.021753111854195595,
	185	+ "nrmse": 0.14748936183398312,
	186	+ "sqnr_db": 16.62478607011158,
	187	+ "cosine": 0.9891892671585083,
	188	+ "max_abs_error": 720.0
	189	+ },
	190	+ "macaq_w4a4": {
	191	+ "nmse": 0.004722047131508589,
	192	+ "nrmse": 0.06871715310974828,
	193	+ "sqnr_db": 23.258696824728887,
	194	+ "cosine": 0.9965612888336182,
	195	+ "max_abs_error": 288.0
	196	+ },
	197	+ "elapsed_cpu_seconds": 0.8263739600079134,
	198	+ "nmse_reduction_factor": 4.606712141656655
	199	+ }
	200	+]

Adocs/benchmarks/macaq_20261027/accuracy-layer-study.metadata.json+40-0

@@ -0,0 +1,40 @@
	1	+{
	2	+ "model": "Qwen/Qwen3-0.6B",
	3	+ "revision": "c1899de289a04d12100db370d81485cdf75e47ca",
	4	+ "torch": "2.9.0+cpu",
	5	+ "execution": "CPU-only method study; not model perplexity or NPU timing",
	6	+ "dataset": "Salesforce/wikitext wikitext-2-raw-v1",
	7	+ "calibration_tokens": 1024,
	8	+ "heldout_tokens": 1024,
	9	+ "token_sha256": {
	10	+ "train": "ea0c105aae5e1f26ceb3a387594b84f26c4e53dc1ddaccdc2f15d68751f917e9",
	11	+ "test": "5972d1f241a23f491babc9456f2ec33790d009894d23fbfcf903f6b9d2efdc32"
	12	+ },
	13	+ "selected_layers": [
	14	+ "model.layers.0.self_attn.q_proj",
	15	+ "model.layers.0.mlp.up_proj",
	16	+ "model.layers.0.mlp.down_proj",
	17	+ "model.layers.14.self_attn.q_proj",
	18	+ "model.layers.14.mlp.up_proj",
	19	+ "model.layers.14.mlp.down_proj",
	20	+ "model.layers.27.self_attn.q_proj",
	21	+ "model.layers.27.mlp.up_proj",
	22	+ "model.layers.27.mlp.down_proj"
	23	+ ],
	24	+ "method_source_sha256": "f53f40ddd4006c43dde6a7fbc7f43b2f210ce2320f73562ef96bc6681401bc59",
	25	+ "reference_source_sha256": "f65f8d3229339fba1a1762edd0bbd9825e534690f15099185fc0eea1452b67f6",
	26	+ "calibration": {
	27	+ "pre_alpha": 0.75,
	28	+ "hadamard_block": 32,
	29	+ "flip_iters": 8,
	30	+ "ridge": 1.0,
	31	+ "exp_screen": [
	32	+ 1,
	33	+ 0,
	34	+ -1
	35	+ ]
	36	+ },
	37	+ "baseline": "ordinary block32 MXFP4 W4A4 RTN; activation scale algorithm matches native baseline scale_alg=0",
	38	+ "reference": "BF16-input/weight matmul accumulated in FP32 and rounded to BF16; no linear bias",
	39	+ "selection": "q_proj, up_proj, down_proj at first/middle/last transformer layer; fixed before evaluating held-out test"
	40	+}

Adocs/benchmarks/macaq_20261027/accuracy-npu-balanced.json+461-0

@@ -0,0 +1,461 @@
	1	+[
	2	+ {
	3	+ "fixture": "accuracy-fixture-model_layers_0_mlp_down_proj.pt",
	4	+ "fixture_sha256": "4f73bdf7d8a927172fea0866e27cbe850c894deb329930bd08b99aab6cff444c",
	5	+ "m": 1024,
	6	+ "k": 3072,
	7	+ "n": 1024,
	8	+ "opsnn_padded_k": 3072,
	9	+ "opsnn_padded_n": 7168,
	10	+ "plain_native": {
	11	+ "nmse": 0.030516970236956067,
	12	+ "nrmse": 0.17469107085640087,
	13	+ "sqnr_db": 15.154585858772782,
	14	+ "cosine": 0.9852403571117876,
	15	+ "max_abs_error": 0.25
	16	+ },
	17	+ "native_scale_alg2": {
	18	+ "nmse": 0.031153790077059946,
	19	+ "nrmse": 0.1765043627706124,
	20	+ "sqnr_db": 15.064891108210311,
	21	+ "cosine": 0.9843569751092435,
	22	+ "max_abs_error": 0.19921875
	23	+ },
	24	+ "native_scale_alg2_error": null,
	25	+ "tilelang_planned_macaq": {
	26	+ "nmse": 0.018508736055048453,
	27	+ "nrmse": 0.13604681567404822,
	28	+ "sqnr_db": 17.326232378143693,
	29	+ "cosine": 0.9908084532765248,
	30	+ "max_abs_error": 0.248046875
	31	+ },
	32	+ "tilelang_macaq": {
	33	+ "nmse": 0.018508736055048453,
	34	+ "nrmse": 0.13604681567404822,
	35	+ "sqnr_db": 17.326232378143693,
	36	+ "cosine": 0.9908084532765248,
	37	+ "max_abs_error": 0.248046875
	38	+ },
	39	+ "opsnn_macaq": {
	40	+ "nmse": 0.018508736055048453,
	41	+ "nrmse": 0.13604681567404822,
	42	+ "sqnr_db": 17.326232378143693,
	43	+ "cosine": 0.9908084532765248,
	44	+ "max_abs_error": 0.248046875
	45	+ },
	46	+ "cpu_reference_agreement": {
	47	+ "plain_max_abs": 0.0,
	48	+ "tilelang_max_abs": 0.0,
	49	+ "tilelang_planned_max_abs": 0.0,
	50	+ "opsnn_max_abs": 0.0
	51	+ }
	52	+ },
	53	+ {
	54	+ "fixture": "accuracy-fixture-model_layers_0_mlp_up_proj.pt",
	55	+ "fixture_sha256": "03b5ea72441a092958e49baa3ec5fe84ec79cd897de2d7d1e3a4e54a7581d92a",
	56	+ "m": 1024,
	57	+ "k": 1024,
	58	+ "n": 3072,
	59	+ "opsnn_padded_k": 2048,
	60	+ "opsnn_padded_n": 3072,
	61	+ "plain_native": {
	62	+ "nmse": 0.03272208139022038,
	63	+ "nrmse": 0.18089245807998844,
	64	+ "sqnr_db": 14.85159079464714,
	65	+ "cosine": 0.9837271166985933,
	66	+ "max_abs_error": 0.640625
	67	+ },
	68	+ "native_scale_alg2": {
	69	+ "nmse": 0.04013932103835252,
	70	+ "nrmse": 0.20034799983616638,
	71	+ "sqnr_db": 13.964299779320108,
	72	+ "cosine": 0.9798668105571983,
	73	+ "max_abs_error": 0.546875
	74	+ },
	75	+ "native_scale_alg2_error": null,
	76	+ "tilelang_planned_macaq": {
	77	+ "nmse": 0.018918712122045826,
	78	+ "nrmse": 0.13754530934221576,
	79	+ "sqnr_db": 17.231084312178428,
	80	+ "cosine": 0.9905704452715268,
	81	+ "max_abs_error": 0.3671875
	82	+ },
	83	+ "tilelang_macaq": {
	84	+ "nmse": 0.018918712122045826,
	85	+ "nrmse": 0.13754530934221576,
	86	+ "sqnr_db": 17.231084312178428,
	87	+ "cosine": 0.9905704452715268,
	88	+ "max_abs_error": 0.3671875
	89	+ },
	90	+ "opsnn_macaq": {
	91	+ "nmse": 0.018918712122045826,
	92	+ "nrmse": 0.13754530934221576,
	93	+ "sqnr_db": 17.231084312178428,
	94	+ "cosine": 0.9905704452715268,
	95	+ "max_abs_error": 0.3671875
	96	+ },
	97	+ "cpu_reference_agreement": {
	98	+ "plain_max_abs": 0.0,
	99	+ "tilelang_max_abs": 0.0,
	100	+ "tilelang_planned_max_abs": 0.0,
	101	+ "opsnn_max_abs": 0.0
	102	+ }
	103	+ },
	104	+ {
	105	+ "fixture": "accuracy-fixture-model_layers_0_self_attn_q_proj.pt",
	106	+ "fixture_sha256": "6ddc2c4ab7cccf3de0b083bb5ccb22cde417e5b4ed3ce81cd2659a579b18076c",
	107	+ "m": 1024,
	108	+ "k": 1024,
	109	+ "n": 2048,
	110	+ "opsnn_padded_k": 1536,
	111	+ "opsnn_padded_n": 2048,
	112	+ "plain_native": {
	113	+ "nmse": 0.012961810244357238,
	114	+ "nrmse": 0.11384994617634757,
	115	+ "sqnr_db": 18.873343407313314,
	116	+ "cosine": 0.9945141181596306,
	117	+ "max_abs_error": 1.015625
	118	+ },
	119	+ "native_scale_alg2": {
	120	+ "nmse": 0.011486990046041059,
	121	+ "nrmse": 0.10717737655886647,
	122	+ "sqnr_db": 19.39793755266881,
	123	+ "cosine": 0.9943654953511734,
	124	+ "max_abs_error": 1.20703125
	125	+ },
	126	+ "native_scale_alg2_error": null,
	127	+ "tilelang_planned_macaq": {
	128	+ "nmse": 0.005795284495877784,
	129	+ "nrmse": 0.07612676596229334,
	130	+ "sqnr_db": 22.36925239259129,
	131	+ "cosine": 0.9972342905702456,
	132	+ "max_abs_error": 0.46875
	133	+ },
	134	+ "tilelang_macaq": {
	135	+ "nmse": 0.005795284495877784,
	136	+ "nrmse": 0.07612676596229334,
	137	+ "sqnr_db": 22.36925239259129,
	138	+ "cosine": 0.9972342905702456,
	139	+ "max_abs_error": 0.46875
	140	+ },
	141	+ "opsnn_macaq": {
	142	+ "nmse": 0.005795284495877784,
	143	+ "nrmse": 0.07612676596229334,
	144	+ "sqnr_db": 22.36925239259129,
	145	+ "cosine": 0.9972342905702456,
	146	+ "max_abs_error": 0.46875
	147	+ },
	148	+ "cpu_reference_agreement": {
	149	+ "plain_max_abs": 0.0,
	150	+ "tilelang_max_abs": 0.0,
	151	+ "tilelang_planned_max_abs": 0.0,
	152	+ "opsnn_max_abs": 0.0
	153	+ }
	154	+ },
	155	+ {
	156	+ "fixture": "accuracy-fixture-model_layers_14_mlp_down_proj.pt",
	157	+ "fixture_sha256": "06327450f71e490385a6137915d300273b901ed344017c98b3cfaf8a7061b579",
	158	+ "m": 1024,
	159	+ "k": 3072,
	160	+ "n": 1024,
	161	+ "opsnn_padded_k": 3072,
	162	+ "opsnn_padded_n": 7168,
	163	+ "plain_native": {
	164	+ "nmse": 0.03410957884917448,
	165	+ "nrmse": 0.18468778749331122,
	166	+ "sqnr_db": 14.671236427916272,
	167	+ "cosine": 0.9830276128202399,
	168	+ "max_abs_error": 1.5625
	169	+ },
	170	+ "native_scale_alg2": {
	171	+ "nmse": 0.03799854210203062,
	172	+ "nrmse": 0.19493214743092177,
	173	+ "sqnr_db": 14.202330657302781,
	174	+ "cosine": 0.9809640218643827,
	175	+ "max_abs_error": 1.8125
	176	+ },
	177	+ "native_scale_alg2_error": null,
	178	+ "tilelang_planned_macaq": {
	179	+ "nmse": 0.021642116309140973,
	180	+ "nrmse": 0.1471125973842518,
	181	+ "sqnr_db": 16.647002733047213,
	182	+ "cosine": 0.989176496558601,
	183	+ "max_abs_error": 0.7265625
	184	+ },
	185	+ "tilelang_macaq": {
	186	+ "nmse": 0.021642116309140973,
	187	+ "nrmse": 0.1471125973842518,
	188	+ "sqnr_db": 16.647002733047213,
	189	+ "cosine": 0.989176496558601,
	190	+ "max_abs_error": 0.7265625
	191	+ },
	192	+ "opsnn_macaq": {
	193	+ "nmse": 0.021642116309140973,
	194	+ "nrmse": 0.1471125973842518,
	195	+ "sqnr_db": 16.647002733047213,
	196	+ "cosine": 0.989176496558601,
	197	+ "max_abs_error": 0.7265625
	198	+ },
	199	+ "cpu_reference_agreement": {
	200	+ "plain_max_abs": 0.0,
	201	+ "tilelang_max_abs": 0.0,
	202	+ "tilelang_planned_max_abs": 0.0,
	203	+ "opsnn_max_abs": 0.0
	204	+ }
	205	+ },
	206	+ {
	207	+ "fixture": "accuracy-fixture-model_layers_14_mlp_up_proj.pt",
	208	+ "fixture_sha256": "93015e6f4ed06d62306d6fccfc75c3fa8a2462740a20c61c72c8b4871668986d",
	209	+ "m": 1024,
	210	+ "k": 1024,
	211	+ "n": 3072,
	212	+ "opsnn_padded_k": 2048,
	213	+ "opsnn_padded_n": 3072,
	214	+ "plain_native": {
	215	+ "nmse": 0.027274659838753348,
	216	+ "nrmse": 0.16515041579951698,
	217	+ "sqnr_db": 15.642406570942896,
	218	+ "cosine": 0.9865070511505231,
	219	+ "max_abs_error": 1.109375
	220	+ },
	221	+ "native_scale_alg2": {
	222	+ "nmse": 0.030852923399939662,
	223	+ "nrmse": 0.17565000256174113,
	224	+ "sqnr_db": 15.107036790761057,
	225	+ "cosine": 0.9845036092016929,
	226	+ "max_abs_error": 1.03125
	227	+ },
	228	+ "native_scale_alg2_error": null,
	229	+ "tilelang_planned_macaq": {
	230	+ "nmse": 0.01696913017803275,
	231	+ "nrmse": 0.13026561395100686,
	232	+ "sqnr_db": 17.7034041864651,
	233	+ "cosine": 0.9915783593889278,
	234	+ "max_abs_error": 0.671875
	235	+ },
	236	+ "tilelang_macaq": {
	237	+ "nmse": 0.01696913017803275,
	238	+ "nrmse": 0.13026561395100686,
	239	+ "sqnr_db": 17.7034041864651,
	240	+ "cosine": 0.9915783593889278,
	241	+ "max_abs_error": 0.671875
	242	+ },
	243	+ "opsnn_macaq": {
	244	+ "nmse": 0.01696913017803275,
	245	+ "nrmse": 0.13026561395100686,
	246	+ "sqnr_db": 17.7034041864651,
	247	+ "cosine": 0.9915783593889278,
	248	+ "max_abs_error": 0.671875
	249	+ },
	250	+ "cpu_reference_agreement": {
	251	+ "plain_max_abs": 0.0,
	252	+ "tilelang_max_abs": 0.0,
	253	+ "tilelang_planned_max_abs": 0.0,
	254	+ "opsnn_max_abs": 0.0
	255	+ }
	256	+ },
	257	+ {
	258	+ "fixture": "accuracy-fixture-model_layers_14_self_attn_q_proj.pt",
	259	+ "fixture_sha256": "abe62ae90b968c85dfc57dbfa8bd287a1672b9a7ba9431570a121b1cff9fbb28",
	260	+ "m": 1024,
	261	+ "k": 1024,
	262	+ "n": 2048,
	263	+ "opsnn_padded_k": 1536,
	264	+ "opsnn_padded_n": 2048,
	265	+ "plain_native": {
	266	+ "nmse": 0.01466745110141977,
	267	+ "nrmse": 0.12110925274899424,
	268	+ "sqnr_db": 18.33645350967023,
	269	+ "cosine": 0.9930062509087988,
	270	+ "max_abs_error": 2.640625
	271	+ },
	272	+ "native_scale_alg2": {
	273	+ "nmse": 0.01625560235156519,
	274	+ "nrmse": 0.12749746017691957,
	275	+ "sqnr_db": 17.88996933081782,
	276	+ "cosine": 0.9918593472312802,
	277	+ "max_abs_error": 3.3125
	278	+ },
	279	+ "native_scale_alg2_error": null,
	280	+ "tilelang_planned_macaq": {
	281	+ "nmse": 0.0075766859551332265,
	282	+ "nrmse": 0.08704416094795346,
	283	+ "sqnr_db": 21.205207133813268,
	284	+ "cosine": 0.9962798805014916,
	285	+ "max_abs_error": 2.1171875
	286	+ },
	287	+ "tilelang_macaq": {
	288	+ "nmse": 0.0075766859551332265,
	289	+ "nrmse": 0.08704416094795346,
	290	+ "sqnr_db": 21.205207133813268,
	291	+ "cosine": 0.9962798805014916,
	292	+ "max_abs_error": 2.1171875
	293	+ },
	294	+ "opsnn_macaq": {
	295	+ "nmse": 0.0075766859551332265,
	296	+ "nrmse": 0.08704416094795346,
	297	+ "sqnr_db": 21.205207133813268,
	298	+ "cosine": 0.9962798805014916,
	299	+ "max_abs_error": 2.1171875
	300	+ },
	301	+ "cpu_reference_agreement": {
	302	+ "plain_max_abs": 0.0,
	303	+ "tilelang_max_abs": 0.0,
	304	+ "tilelang_planned_max_abs": 0.0,
	305	+ "opsnn_max_abs": 0.0
	306	+ }
	307	+ },
	308	+ {
	309	+ "fixture": "accuracy-fixture-model_layers_27_mlp_down_proj.pt",
	310	+ "fixture_sha256": "18f102d83016adb97aa542df1515be28ab7acea4059b91fcd8acb22427382c25",
	311	+ "m": 1024,
	312	+ "k": 3072,
	313	+ "n": 1024,
	314	+ "opsnn_padded_k": 3072,
	315	+ "opsnn_padded_n": 7168,
	316	+ "plain_native": {
	317	+ "nmse": 0.021753115704479336,
	318	+ "nrmse": 0.14748937488673322,
	319	+ "sqnr_db": 16.62478530141382,
	320	+ "cosine": 0.9903826287274079,
	321	+ "max_abs_error": 720.0
	322	+ },
	323	+ "native_scale_alg2": {
	324	+ "nmse": 0.019562300238837857,
	325	+ "nrmse": 0.13986529318897473,
	326	+ "sqnr_db": 17.085800799012137,
	327	+ "cosine": 0.990239036351425,
	328	+ "max_abs_error": 192.0
	329	+ },
	330	+ "native_scale_alg2_error": null,
	331	+ "tilelang_planned_macaq": {
	332	+ "nmse": 0.0047220473550336605,
	333	+ "nrmse": 0.06871715473616222,
	334	+ "sqnr_db": 23.258696619149188,
	335	+ "cosine": 0.9980580611788563,
	336	+ "max_abs_error": 288.0
	337	+ },
	338	+ "tilelang_macaq": {
	339	+ "nmse": 0.0047220473550336605,
	340	+ "nrmse": 0.06871715473616222,
	341	+ "sqnr_db": 23.258696619149188,
	342	+ "cosine": 0.9980580611788563,
	343	+ "max_abs_error": 288.0
	344	+ },
	345	+ "opsnn_macaq": {
	346	+ "nmse": 0.0047220473550336605,
	347	+ "nrmse": 0.06871715473616222,
	348	+ "sqnr_db": 23.258696619149188,
	349	+ "cosine": 0.9980580611788563,
	350	+ "max_abs_error": 288.0
	351	+ },
	352	+ "cpu_reference_agreement": {
	353	+ "plain_max_abs": 0.0,
	354	+ "tilelang_max_abs": 0.0,
	355	+ "tilelang_planned_max_abs": 0.0,
	356	+ "opsnn_max_abs": 0.0
	357	+ }
	358	+ },
	359	+ {
	360	+ "fixture": "accuracy-fixture-model_layers_27_mlp_up_proj.pt",
	361	+ "fixture_sha256": "e6ec478d4b4190fdef10af43e0f09cf14b869eb66bb3dcfa80f91bf0a6681297",
	362	+ "m": 1024,
	363	+ "k": 1024,
	364	+ "n": 3072,
	365	+ "opsnn_padded_k": 2048,
	366	+ "opsnn_padded_n": 3072,
	367	+ "plain_native": {
	368	+ "nmse": 0.021787974624019088,
	369	+ "nrmse": 0.1476075019232393,
	370	+ "sqnr_db": 16.61783139211963,
	371	+ "cosine": 0.9896387374240014,
	372	+ "max_abs_error": 53.0
	373	+ },
	374	+ "native_scale_alg2": {
	375	+ "nmse": 0.025910815855607353,
	376	+ "nrmse": 0.16096836911520024,
	377	+ "sqnr_db": 15.865189121392245,
	378	+ "cosine": 0.9869727854323713,
	379	+ "max_abs_error": 33.734375
	380	+ },
	381	+ "native_scale_alg2_error": null,
	382	+ "tilelang_planned_macaq": {
	383	+ "nmse": 0.008336486324689736,
	384	+ "nrmse": 0.09130436092920062,
	385	+ "sqnr_db": 20.790169579159002,
	386	+ "cosine": 0.9959198014115227,
	387	+ "max_abs_error": 39.0
	388	+ },
	389	+ "tilelang_macaq": {
	390	+ "nmse": 0.008336486324689736,
	391	+ "nrmse": 0.09130436092920062,
	392	+ "sqnr_db": 20.790169579159002,
	393	+ "cosine": 0.9959198014115227,
	394	+ "max_abs_error": 39.0
	395	+ },
	396	+ "opsnn_macaq": {
	397	+ "nmse": 0.008336486324689736,
	398	+ "nrmse": 0.09130436092920062,
	399	+ "sqnr_db": 20.790169579159002,
	400	+ "cosine": 0.9959198014115227,
	401	+ "max_abs_error": 39.0
	402	+ },
	403	+ "cpu_reference_agreement": {
	404	+ "plain_max_abs": 0.0,
	405	+ "tilelang_max_abs": 0.0,
	406	+ "tilelang_planned_max_abs": 0.0,
	407	+ "opsnn_max_abs": 0.0
	408	+ }
	409	+ },
	410	+ {
	411	+ "fixture": "accuracy-fixture-model_layers_27_self_attn_q_proj.pt",
	412	+ "fixture_sha256": "2ea2b3baf308ff4f1f81f978b28137ef5e98c111166fc24e8b2d6bb9aab7f391",
	413	+ "m": 1024,
	414	+ "k": 1024,
	415	+ "n": 2048,
	416	+ "opsnn_padded_k": 1536,
	417	+ "opsnn_padded_n": 2048,
	418	+ "plain_native": {
	419	+ "nmse": 0.014878039646034468,
	420	+ "nrmse": 0.12197556987378443,
	421	+ "sqnr_db": 18.27454288346345,
	422	+ "cosine": 0.9928836409767186,
	423	+ "max_abs_error": 32.5
	424	+ },
	425	+ "native_scale_alg2": {
	426	+ "nmse": 0.017160446816167834,
	427	+ "nrmse": 0.13099788859431222,
	428	+ "sqnr_db": 17.654714083713575,
	429	+ "cosine": 0.9913830249456285,
	430	+ "max_abs_error": 30.0
	431	+ },
	432	+ "native_scale_alg2_error": null,
	433	+ "tilelang_planned_macaq": {
	434	+ "nmse": 0.007906629422839765,
	435	+ "nrmse": 0.08891922976971722,
	436	+ "sqnr_db": 21.02008615747497,
	437	+ "cosine": 0.9961864477162093,
	438	+ "max_abs_error": 17.75
	439	+ },
	440	+ "tilelang_macaq": {
	441	+ "nmse": 0.007906629422839765,
	442	+ "nrmse": 0.08891922976971722,
	443	+ "sqnr_db": 21.02008615747497,
	444	+ "cosine": 0.9961864477162093,
	445	+ "max_abs_error": 17.75
	446	+ },
	447	+ "opsnn_macaq": {
	448	+ "nmse": 0.007906629422839765,
	449	+ "nrmse": 0.08891922976971722,
	450	+ "sqnr_db": 21.02008615747497,
	451	+ "cosine": 0.9961864477162093,
	452	+ "max_abs_error": 17.75
	453	+ },
	454	+ "cpu_reference_agreement": {
	455	+ "plain_max_abs": 0.0,
	456	+ "tilelang_max_abs": 0.0,
	457	+ "tilelang_planned_max_abs": 0.0,
	458	+ "opsnn_max_abs": 0.0
	459	+ }
	460	+ }
	461	+]

Adocs/benchmarks/macaq_20261027/accuracy-npu-balanced.metadata.json+37-0

@@ -0,0 +1,37 @@
	1	+{
	2	+ "device": "Ascend950DT_9592",
	3	+ "torch": "2.9.0+cpu",
	4	+ "torch_npu": "2.9.0.post2",
	5	+ "opsnn_wheel_sha256": "b284336b70f82a80a4c2feba59bbc7f53cf178563f988dd89e527ef44ad50e06",
	6	+ "source_cpp_sha256": {
	7	+ "experimental/matmul/common/macaq/macaq_hadamard_mxfp4.hpp": "c62ab428e2d712d38b44e2e2de71605210e73b7fc4edac087a96a8aeb48c3574",
	8	+ "experimental/matmul/common/macaq/macaq_w4a4_kernel_impl.hpp": "52a5ed13b6798df7b36cdc08335663b2ffb70b52bda12180712b0dbb1ea1c26c",
	9	+ "experimental/matmul/macaq_w4a4_matmul/macaq_w4a4_matmul.cpp": "2ddea3d5dc3eed689a3aa59a43e219247b47344180a8ea23c24f94316a05ff44",
	10	+ "experimental/matmul/macaq_w4a4_grouped_matmul/macaq_w4a4_grouped_matmul.cpp": "63637c9de33517b69800b4dd0d748c17cbe5910facab4e5db39de7f931d38c93",
	11	+ "experimental/matmul/common/macaq/tests/test_macaq_raw_diagonal.py": "72f7a4536734be62e733656ce6cc45bd4b445529d995a224bc6cf309d24c7ed7"
	12	+ },
	13	+ "tilelang_base_commit": "b7d8b14c1e14f2412e4a37ccecde85e6284d0349",
	14	+ "model": "Qwen/Qwen3-0.6B",
	15	+ "checks": 36,
	16	+ "native_scale_alg2_checks": 9,
	17	+ "tilelang_source_sha256": {
	18	+ "__init__.py": "a81e80bbf9fd868d5109d2ab725dec9afe16d0a346ccad7809e55fb6dddfd2a7",
	19	+ "matmul/__init__.py": "13c8e8bdf2323d8787bcdf944464e1975343b35fe95a95e40d8d5612cec5a1eb",
	20	+ "matmul/macaq_w4a4_matmul.py": "c0604c1535ecc4b29e9e62ec345bea133f42c0e9d8537b63b6728d00b27b6648",
	21	+ "matmul/macaq_w4a4_matmul_kernel.py": "faca71b569e0d54d6fa499e84765f0bcdcacbcbbd68651e359f990ce785b720e"
	22	+ },
	23	+ "scope": "nine real projection fixtures, actual native and MACaQ NPU outputs vs CPU algorithm references; not whole-model NPU evaluation or model throughput",
	24	+ "padding": "ops-nn zero-pads K/N to supported pairs and slices original columns; zero padding preserves block32 transform boundaries; no latency claim for padded calls",
	25	+ "fixture_sha256": {
	26	+ "accuracy-fixture-model_layers_14_mlp_up_proj.pt": "93015e6f4ed06d62306d6fccfc75c3fa8a2462740a20c61c72c8b4871668986d",
	27	+ "accuracy-fixture-model_layers_14_mlp_down_proj.pt": "06327450f71e490385a6137915d300273b901ed344017c98b3cfaf8a7061b579",
	28	+ "accuracy-fixture-model_layers_14_self_attn_q_proj.pt": "abe62ae90b968c85dfc57dbfa8bd287a1672b9a7ba9431570a121b1cff9fbb28",
	29	+ "accuracy-fixture-model_layers_27_self_attn_q_proj.pt": "2ea2b3baf308ff4f1f81f978b28137ef5e98c111166fc24e8b2d6bb9aab7f391",
	30	+ "accuracy-fixture-model_layers_27_mlp_up_proj.pt": "e6ec478d4b4190fdef10af43e0f09cf14b869eb66bb3dcfa80f91bf0a6681297",
	31	+ "accuracy-fixture-model_layers_27_mlp_down_proj.pt": "18f102d83016adb97aa542df1515be28ab7acea4059b91fcd8acb22427382c25",
	32	+ "accuracy-fixture-model_layers_0_self_attn_q_proj.pt": "6ddc2c4ab7cccf3de0b083bb5ccb22cde417e5b4ed3ce81cd2659a579b18076c",
	33	+ "accuracy-fixture-model_layers_0_mlp_down_proj.pt": "4f73bdf7d8a927172fea0866e27cbe850c894deb329930bd08b99aab6cff444c",
	34	+ "accuracy-fixture-model_layers_0_mlp_up_proj.pt": "03b5ea72441a092958e49baa3ec5fe84ec79cd897de2d7d1e3a4e54a7581d92a"
	35	+ },
	36	+ "harness_sha256": "2729878cd01f4e40b0e178e90eddaef382db2da9629e8212039efd3905eb4c36"
	37	+}

Adocs/benchmarks/macaq_20261027/accuracy-npu-bridge.json+317-0

@@ -0,0 +1,317 @@
	1	+[
	2	+ {
	3	+ "fixture": "accuracy-fixture-model_layers_0_mlp_down_proj.pt",
	4	+ "fixture_sha256": "4f73bdf7d8a927172fea0866e27cbe850c894deb329930bd08b99aab6cff444c",
	5	+ "m": 1024,
	6	+ "k": 3072,
	7	+ "n": 1024,
	8	+ "opsnn_padded_k": 3072,
	9	+ "opsnn_padded_n": 7168,
	10	+ "plain_native": {
	11	+ "nmse": 0.030516970236956067,
	12	+ "nrmse": 0.17469107085640087,
	13	+ "sqnr_db": 15.154585858772782,
	14	+ "cosine": 0.9852403571117876,
	15	+ "max_abs_error": 0.25
	16	+ },
	17	+ "tilelang_macaq": {
	18	+ "nmse": 0.018508736055048453,
	19	+ "nrmse": 0.13604681567404822,
	20	+ "sqnr_db": 17.326232378143693,
	21	+ "cosine": 0.9908084532765248,
	22	+ "max_abs_error": 0.248046875
	23	+ },
	24	+ "opsnn_macaq": {
	25	+ "nmse": 0.018508736055048453,
	26	+ "nrmse": 0.13604681567404822,
	27	+ "sqnr_db": 17.326232378143693,
	28	+ "cosine": 0.9908084532765248,
	29	+ "max_abs_error": 0.248046875
	30	+ },
	31	+ "cpu_reference_agreement": {
	32	+ "plain_max_abs": 0.0,
	33	+ "tilelang_max_abs": 0.0,
	34	+ "opsnn_max_abs": 0.0
	35	+ }
	36	+ },
	37	+ {
	38	+ "fixture": "accuracy-fixture-model_layers_0_mlp_up_proj.pt",
	39	+ "fixture_sha256": "03b5ea72441a092958e49baa3ec5fe84ec79cd897de2d7d1e3a4e54a7581d92a",
	40	+ "m": 1024,
	41	+ "k": 1024,
	42	+ "n": 3072,
	43	+ "opsnn_padded_k": 2048,
	44	+ "opsnn_padded_n": 3072,
	45	+ "plain_native": {
	46	+ "nmse": 0.03272208139022038,
	47	+ "nrmse": 0.18089245807998844,
	48	+ "sqnr_db": 14.85159079464714,
	49	+ "cosine": 0.9837271166985933,
	50	+ "max_abs_error": 0.640625
	51	+ },
	52	+ "tilelang_macaq": {
	53	+ "nmse": 0.018918712122045826,
	54	+ "nrmse": 0.13754530934221576,
	55	+ "sqnr_db": 17.231084312178428,
	56	+ "cosine": 0.9905704452715268,
	57	+ "max_abs_error": 0.3671875
	58	+ },
	59	+ "opsnn_macaq": {
	60	+ "nmse": 0.018918712122045826,
	61	+ "nrmse": 0.13754530934221576,
	62	+ "sqnr_db": 17.231084312178428,
	63	+ "cosine": 0.9905704452715268,
	64	+ "max_abs_error": 0.3671875
	65	+ },
	66	+ "cpu_reference_agreement": {
	67	+ "plain_max_abs": 0.0,
	68	+ "tilelang_max_abs": 0.0,
	69	+ "opsnn_max_abs": 0.0
	70	+ }
	71	+ },
	72	+ {
	73	+ "fixture": "accuracy-fixture-model_layers_0_self_attn_q_proj.pt",
	74	+ "fixture_sha256": "6ddc2c4ab7cccf3de0b083bb5ccb22cde417e5b4ed3ce81cd2659a579b18076c",
	75	+ "m": 1024,
	76	+ "k": 1024,
	77	+ "n": 2048,
	78	+ "opsnn_padded_k": 1536,
	79	+ "opsnn_padded_n": 2048,
	80	+ "plain_native": {
	81	+ "nmse": 0.012961810244357238,
	82	+ "nrmse": 0.11384994617634757,
	83	+ "sqnr_db": 18.873343407313314,
	84	+ "cosine": 0.9945141181596306,
	85	+ "max_abs_error": 1.015625
	86	+ },
	87	+ "tilelang_macaq": {
	88	+ "nmse": 0.005795284495877784,
	89	+ "nrmse": 0.07612676596229334,
	90	+ "sqnr_db": 22.36925239259129,
	91	+ "cosine": 0.9972342905702456,
	92	+ "max_abs_error": 0.46875
	93	+ },
	94	+ "opsnn_macaq": {
	95	+ "nmse": 0.005795284495877784,
	96	+ "nrmse": 0.07612676596229334,
	97	+ "sqnr_db": 22.36925239259129,
	98	+ "cosine": 0.9972342905702456,
	99	+ "max_abs_error": 0.46875
	100	+ },
	101	+ "cpu_reference_agreement": {
	102	+ "plain_max_abs": 0.0,
	103	+ "tilelang_max_abs": 0.0,
	104	+ "opsnn_max_abs": 0.0
	105	+ }
	106	+ },
	107	+ {
	108	+ "fixture": "accuracy-fixture-model_layers_14_mlp_down_proj.pt",
	109	+ "fixture_sha256": "06327450f71e490385a6137915d300273b901ed344017c98b3cfaf8a7061b579",
	110	+ "m": 1024,
	111	+ "k": 3072,
	112	+ "n": 1024,
	113	+ "opsnn_padded_k": 3072,
	114	+ "opsnn_padded_n": 7168,
	115	+ "plain_native": {
	116	+ "nmse": 0.03410957884917448,
	117	+ "nrmse": 0.18468778749331122,
	118	+ "sqnr_db": 14.671236427916272,
	119	+ "cosine": 0.9830276128202399,
	120	+ "max_abs_error": 1.5625
	121	+ },
	122	+ "tilelang_macaq": {
	123	+ "nmse": 0.021642116309140973,
	124	+ "nrmse": 0.1471125973842518,
	125	+ "sqnr_db": 16.647002733047213,
	126	+ "cosine": 0.989176496558601,
	127	+ "max_abs_error": 0.7265625
	128	+ },
	129	+ "opsnn_macaq": {
	130	+ "nmse": 0.021642116309140973,
	131	+ "nrmse": 0.1471125973842518,
	132	+ "sqnr_db": 16.647002733047213,
	133	+ "cosine": 0.989176496558601,
	134	+ "max_abs_error": 0.7265625
	135	+ },
	136	+ "cpu_reference_agreement": {
	137	+ "plain_max_abs": 0.0,
	138	+ "tilelang_max_abs": 0.0,
	139	+ "opsnn_max_abs": 0.0
	140	+ }
	141	+ },
	142	+ {
	143	+ "fixture": "accuracy-fixture-model_layers_14_mlp_up_proj.pt",
	144	+ "fixture_sha256": "93015e6f4ed06d62306d6fccfc75c3fa8a2462740a20c61c72c8b4871668986d",
	145	+ "m": 1024,
	146	+ "k": 1024,
	147	+ "n": 3072,
	148	+ "opsnn_padded_k": 2048,
	149	+ "opsnn_padded_n": 3072,
	150	+ "plain_native": {
	151	+ "nmse": 0.027274659838753348,
	152	+ "nrmse": 0.16515041579951698,
	153	+ "sqnr_db": 15.642406570942896,
	154	+ "cosine": 0.9865070511505231,
	155	+ "max_abs_error": 1.109375
	156	+ },
	157	+ "tilelang_macaq": {
	158	+ "nmse": 0.01696913017803275,
	159	+ "nrmse": 0.13026561395100686,
	160	+ "sqnr_db": 17.7034041864651,
	161	+ "cosine": 0.9915783593889278,
	162	+ "max_abs_error": 0.671875
	163	+ },
	164	+ "opsnn_macaq": {
	165	+ "nmse": 0.01696913017803275,
	166	+ "nrmse": 0.13026561395100686,
	167	+ "sqnr_db": 17.7034041864651,
	168	+ "cosine": 0.9915783593889278,
	169	+ "max_abs_error": 0.671875
	170	+ },
	171	+ "cpu_reference_agreement": {
	172	+ "plain_max_abs": 0.0,
	173	+ "tilelang_max_abs": 0.0,
	174	+ "opsnn_max_abs": 0.0
	175	+ }
	176	+ },
	177	+ {
	178	+ "fixture": "accuracy-fixture-model_layers_14_self_attn_q_proj.pt",
	179	+ "fixture_sha256": "abe62ae90b968c85dfc57dbfa8bd287a1672b9a7ba9431570a121b1cff9fbb28",
	180	+ "m": 1024,
	181	+ "k": 1024,
	182	+ "n": 2048,
	183	+ "opsnn_padded_k": 1536,
	184	+ "opsnn_padded_n": 2048,
	185	+ "plain_native": {
	186	+ "nmse": 0.01466745110141977,
	187	+ "nrmse": 0.12110925274899424,
	188	+ "sqnr_db": 18.33645350967023,
	189	+ "cosine": 0.9930062509087988,
	190	+ "max_abs_error": 2.640625
	191	+ },
	192	+ "tilelang_macaq": {
	193	+ "nmse": 0.0075766859551332265,
	194	+ "nrmse": 0.08704416094795346,
	195	+ "sqnr_db": 21.205207133813268,
	196	+ "cosine": 0.9962798805014916,
	197	+ "max_abs_error": 2.1171875
	198	+ },
	199	+ "opsnn_macaq": {
	200	+ "nmse": 0.0075766859551332265,
	201	+ "nrmse": 0.08704416094795346,
	202	+ "sqnr_db": 21.205207133813268,
	203	+ "cosine": 0.9962798805014916,
	204	+ "max_abs_error": 2.1171875
	205	+ },
	206	+ "cpu_reference_agreement": {
	207	+ "plain_max_abs": 0.0,
	208	+ "tilelang_max_abs": 0.0,
	209	+ "opsnn_max_abs": 0.0
	210	+ }
	211	+ },
	212	+ {
	213	+ "fixture": "accuracy-fixture-model_layers_27_mlp_down_proj.pt",
	214	+ "fixture_sha256": "18f102d83016adb97aa542df1515be28ab7acea4059b91fcd8acb22427382c25",
	215	+ "m": 1024,
	216	+ "k": 3072,
	217	+ "n": 1024,
	218	+ "opsnn_padded_k": 3072,
	219	+ "opsnn_padded_n": 7168,
	220	+ "plain_native": {
	221	+ "nmse": 0.021753115704479336,
	222	+ "nrmse": 0.14748937488673322,
	223	+ "sqnr_db": 16.62478530141382,
	224	+ "cosine": 0.9903826287274079,
	225	+ "max_abs_error": 720.0
	226	+ },
	227	+ "tilelang_macaq": {
	228	+ "nmse": 0.0047220473550336605,
	229	+ "nrmse": 0.06871715473616222,
	230	+ "sqnr_db": 23.258696619149188,
	231	+ "cosine": 0.9980580611788563,
	232	+ "max_abs_error": 288.0
	233	+ },
	234	+ "opsnn_macaq": {
	235	+ "nmse": 0.0047220473550336605,
	236	+ "nrmse": 0.06871715473616222,
	237	+ "sqnr_db": 23.258696619149188,
	238	+ "cosine": 0.9980580611788563,
	239	+ "max_abs_error": 288.0
	240	+ },
	241	+ "cpu_reference_agreement": {
	242	+ "plain_max_abs": 0.0,
	243	+ "tilelang_max_abs": 0.0,
	244	+ "opsnn_max_abs": 0.0
	245	+ }
	246	+ },
	247	+ {
	248	+ "fixture": "accuracy-fixture-model_layers_27_mlp_up_proj.pt",
	249	+ "fixture_sha256": "e6ec478d4b4190fdef10af43e0f09cf14b869eb66bb3dcfa80f91bf0a6681297",
	250	+ "m": 1024,
	251	+ "k": 1024,
	252	+ "n": 3072,
	253	+ "opsnn_padded_k": 2048,
	254	+ "opsnn_padded_n": 3072,
	255	+ "plain_native": {
	256	+ "nmse": 0.021787974624019088,
	257	+ "nrmse": 0.1476075019232393,
	258	+ "sqnr_db": 16.61783139211963,
	259	+ "cosine": 0.9896387374240014,
	260	+ "max_abs_error": 53.0
	261	+ },
	262	+ "tilelang_macaq": {
	263	+ "nmse": 0.008336486324689736,
	264	+ "nrmse": 0.09130436092920062,
	265	+ "sqnr_db": 20.790169579159002,
	266	+ "cosine": 0.9959198014115227,
	267	+ "max_abs_error": 39.0
	268	+ },
	269	+ "opsnn_macaq": {
	270	+ "nmse": 0.008336486324689736,
	271	+ "nrmse": 0.09130436092920062,
	272	+ "sqnr_db": 20.790169579159002,
	273	+ "cosine": 0.9959198014115227,
	274	+ "max_abs_error": 39.0
	275	+ },
	276	+ "cpu_reference_agreement": {
	277	+ "plain_max_abs": 0.0,
	278	+ "tilelang_max_abs": 0.0,
	279	+ "opsnn_max_abs": 0.0
	280	+ }
	281	+ },
	282	+ {
	283	+ "fixture": "accuracy-fixture-model_layers_27_self_attn_q_proj.pt",
	284	+ "fixture_sha256": "2ea2b3baf308ff4f1f81f978b28137ef5e98c111166fc24e8b2d6bb9aab7f391",
	285	+ "m": 1024,
	286	+ "k": 1024,
	287	+ "n": 2048,
	288	+ "opsnn_padded_k": 1536,
	289	+ "opsnn_padded_n": 2048,
	290	+ "plain_native": {
	291	+ "nmse": 0.014878039646034468,
	292	+ "nrmse": 0.12197556987378443,
	293	+ "sqnr_db": 18.27454288346345,
	294	+ "cosine": 0.9928836409767186,
	295	+ "max_abs_error": 32.5
	296	+ },
	297	+ "tilelang_macaq": {
	298	+ "nmse": 0.007906629422839765,
	299	+ "nrmse": 0.08891922976971722,
	300	+ "sqnr_db": 21.02008615747497,
	301	+ "cosine": 0.9961864477162093,
	302	+ "max_abs_error": 17.75
	303	+ },
	304	+ "opsnn_macaq": {
	305	+ "nmse": 0.007906629422839765,
	306	+ "nrmse": 0.08891922976971722,
	307	+ "sqnr_db": 21.02008615747497,
	308	+ "cosine": 0.9961864477162093,
	309	+ "max_abs_error": 17.75
	310	+ },
	311	+ "cpu_reference_agreement": {
	312	+ "plain_max_abs": 0.0,
	313	+ "tilelang_max_abs": 0.0,
	314	+ "opsnn_max_abs": 0.0
	315	+ }
	316	+ }
	317	+]

Adocs/benchmarks/macaq_20261027/accuracy-npu-bridge.metadata.json+30-0

@@ -0,0 +1,30 @@
	1	+{
	2	+ "device": "Ascend950DT_9592",
	3	+ "torch": "2.9.0+cpu",
	4	+ "torch_npu": "2.9.0.post2",
	5	+ "opsnn_wheel_sha256": "b284336b70f82a80a4c2feba59bbc7f53cf178563f988dd89e527ef44ad50e06",
	6	+ "source_cpp_sha256": {
	7	+ "experimental/matmul/common/macaq/macaq_hadamard_mxfp4.hpp": "c62ab428e2d712d38b44e2e2de71605210e73b7fc4edac087a96a8aeb48c3574",
	8	+ "experimental/matmul/common/macaq/macaq_w4a4_kernel_impl.hpp": "52a5ed13b6798df7b36cdc08335663b2ffb70b52bda12180712b0dbb1ea1c26c",
	9	+ "experimental/matmul/macaq_w4a4_matmul/macaq_w4a4_matmul.cpp": "2ddea3d5dc3eed689a3aa59a43e219247b47344180a8ea23c24f94316a05ff44",
	10	+ "experimental/matmul/macaq_w4a4_grouped_matmul/macaq_w4a4_grouped_matmul.cpp": "63637c9de33517b69800b4dd0d748c17cbe5910facab4e5db39de7f931d38c93",
	11	+ "experimental/matmul/common/macaq/tests/test_macaq_raw_diagonal.py": "72f7a4536734be62e733656ce6cc45bd4b445529d995a224bc6cf309d24c7ed7"
	12	+ },
	13	+ "tilelang_base_commit": "b7d8b14c1e14f2412e4a37ccecde85e6284d0349",
	14	+ "model": "Qwen/Qwen3-0.6B",
	15	+ "checks": 27,
	16	+ "scope": "nine real projection fixtures, actual native and MACaQ NPU outputs vs CPU algorithm references; not whole-model NPU evaluation or model throughput",
	17	+ "padding": "ops-nn zero-pads K/N to supported pairs and slices original columns; zero padding preserves block32 transform boundaries; no latency claim for padded calls",
	18	+ "fixture_sha256": {
	19	+ "accuracy-fixture-model_layers_14_mlp_up_proj.pt": "93015e6f4ed06d62306d6fccfc75c3fa8a2462740a20c61c72c8b4871668986d",
	20	+ "accuracy-fixture-model_layers_14_mlp_down_proj.pt": "06327450f71e490385a6137915d300273b901ed344017c98b3cfaf8a7061b579",
	21	+ "accuracy-fixture-model_layers_14_self_attn_q_proj.pt": "abe62ae90b968c85dfc57dbfa8bd287a1672b9a7ba9431570a121b1cff9fbb28",
	22	+ "accuracy-fixture-model_layers_27_self_attn_q_proj.pt": "2ea2b3baf308ff4f1f81f978b28137ef5e98c111166fc24e8b2d6bb9aab7f391",
	23	+ "accuracy-fixture-model_layers_27_mlp_up_proj.pt": "e6ec478d4b4190fdef10af43e0f09cf14b869eb66bb3dcfa80f91bf0a6681297",
	24	+ "accuracy-fixture-model_layers_27_mlp_down_proj.pt": "18f102d83016adb97aa542df1515be28ab7acea4059b91fcd8acb22427382c25",
	25	+ "accuracy-fixture-model_layers_0_self_attn_q_proj.pt": "6ddc2c4ab7cccf3de0b083bb5ccb22cde417e5b4ed3ce81cd2659a579b18076c",
	26	+ "accuracy-fixture-model_layers_0_mlp_down_proj.pt": "4f73bdf7d8a927172fea0866e27cbe850c894deb329930bd08b99aab6cff444c",
	27	+ "accuracy-fixture-model_layers_0_mlp_up_proj.pt": "03b5ea72441a092958e49baa3ec5fe84ec79cd897de2d7d1e3a4e54a7581d92a"
	28	+ },
	29	+ "harness_sha256": "506df5b86b8374788fc9e12aab2b0c884da8b383b2c5a8c0a1ed0a604963ddd4"
	30	+}

Adocs/benchmarks/macaq_20261027/accuracy-npu-dynamic-final.json+461-0

@@ -0,0 +1,461 @@
	1	+[
	2	+ {
	3	+ "fixture": "accuracy-fixture-model_layers_0_mlp_down_proj.pt",
	4	+ "fixture_sha256": "4f73bdf7d8a927172fea0866e27cbe850c894deb329930bd08b99aab6cff444c",
	5	+ "m": 1024,
	6	+ "k": 3072,
	7	+ "n": 1024,
	8	+ "opsnn_padded_k": 3072,
	9	+ "opsnn_padded_n": 7168,
	10	+ "plain_native": {
	11	+ "nmse": 0.030516970236956067,
	12	+ "nrmse": 0.17469107085640087,
	13	+ "sqnr_db": 15.154585858772782,
	14	+ "cosine": 0.9852403571117876,
	15	+ "max_abs_error": 0.25
	16	+ },
	17	+ "native_scale_alg2": {
	18	+ "nmse": 0.031153790077059946,
	19	+ "nrmse": 0.1765043627706124,
	20	+ "sqnr_db": 15.064891108210311,
	21	+ "cosine": 0.9843569751092435,
	22	+ "max_abs_error": 0.19921875
	23	+ },
	24	+ "native_scale_alg2_error": null,
	25	+ "tilelang_planned_macaq": {
	26	+ "nmse": 0.018508736055048453,
	27	+ "nrmse": 0.13604681567404822,
	28	+ "sqnr_db": 17.326232378143693,
	29	+ "cosine": 0.9908084532765248,
	30	+ "max_abs_error": 0.248046875
	31	+ },
	32	+ "tilelang_macaq": {
	33	+ "nmse": 0.018508736055048453,
	34	+ "nrmse": 0.13604681567404822,
	35	+ "sqnr_db": 17.326232378143693,
	36	+ "cosine": 0.9908084532765248,
	37	+ "max_abs_error": 0.248046875
	38	+ },
	39	+ "opsnn_macaq": {
	40	+ "nmse": 0.018508736055048453,
	41	+ "nrmse": 0.13604681567404822,
	42	+ "sqnr_db": 17.326232378143693,
	43	+ "cosine": 0.9908084532765248,
	44	+ "max_abs_error": 0.248046875
	45	+ },
	46	+ "cpu_reference_agreement": {
	47	+ "plain_max_abs": 0.0,
	48	+ "tilelang_max_abs": 0.0,
	49	+ "tilelang_planned_max_abs": 0.0,
	50	+ "opsnn_max_abs": 0.0
	51	+ }
	52	+ },
	53	+ {
	54	+ "fixture": "accuracy-fixture-model_layers_0_mlp_up_proj.pt",
	55	+ "fixture_sha256": "03b5ea72441a092958e49baa3ec5fe84ec79cd897de2d7d1e3a4e54a7581d92a",
	56	+ "m": 1024,
	57	+ "k": 1024,
	58	+ "n": 3072,
	59	+ "opsnn_padded_k": 2048,
	60	+ "opsnn_padded_n": 3072,
	61	+ "plain_native": {
	62	+ "nmse": 0.03272208139022038,
	63	+ "nrmse": 0.18089245807998844,
	64	+ "sqnr_db": 14.85159079464714,
	65	+ "cosine": 0.9837271166985933,
	66	+ "max_abs_error": 0.640625
	67	+ },
	68	+ "native_scale_alg2": {
	69	+ "nmse": 0.04013932103835252,
	70	+ "nrmse": 0.20034799983616638,
	71	+ "sqnr_db": 13.964299779320108,
	72	+ "cosine": 0.9798668105571983,
	73	+ "max_abs_error": 0.546875
	74	+ },
	75	+ "native_scale_alg2_error": null,
	76	+ "tilelang_planned_macaq": {
	77	+ "nmse": 0.018918712122045826,
	78	+ "nrmse": 0.13754530934221576,
	79	+ "sqnr_db": 17.231084312178428,
	80	+ "cosine": 0.9905704452715268,
	81	+ "max_abs_error": 0.3671875
	82	+ },
	83	+ "tilelang_macaq": {
	84	+ "nmse": 0.018918712122045826,
	85	+ "nrmse": 0.13754530934221576,
	86	+ "sqnr_db": 17.231084312178428,
	87	+ "cosine": 0.9905704452715268,
	88	+ "max_abs_error": 0.3671875
	89	+ },
	90	+ "opsnn_macaq": {
	91	+ "nmse": 0.018918712122045826,
	92	+ "nrmse": 0.13754530934221576,
	93	+ "sqnr_db": 17.231084312178428,
	94	+ "cosine": 0.9905704452715268,
	95	+ "max_abs_error": 0.3671875
	96	+ },
	97	+ "cpu_reference_agreement": {
	98	+ "plain_max_abs": 0.0,
	99	+ "tilelang_max_abs": 0.0,
	100	+ "tilelang_planned_max_abs": 0.0,
	101	+ "opsnn_max_abs": 0.0
	102	+ }
	103	+ },
	104	+ {
	105	+ "fixture": "accuracy-fixture-model_layers_0_self_attn_q_proj.pt",
	106	+ "fixture_sha256": "6ddc2c4ab7cccf3de0b083bb5ccb22cde417e5b4ed3ce81cd2659a579b18076c",
	107	+ "m": 1024,
	108	+ "k": 1024,
	109	+ "n": 2048,
	110	+ "opsnn_padded_k": 1536,
	111	+ "opsnn_padded_n": 2048,
	112	+ "plain_native": {
	113	+ "nmse": 0.012961810244357238,
	114	+ "nrmse": 0.11384994617634757,
	115	+ "sqnr_db": 18.873343407313314,
	116	+ "cosine": 0.9945141181596306,
	117	+ "max_abs_error": 1.015625
	118	+ },
	119	+ "native_scale_alg2": {
	120	+ "nmse": 0.011486990046041059,
	121	+ "nrmse": 0.10717737655886647,
	122	+ "sqnr_db": 19.39793755266881,
	123	+ "cosine": 0.9943654953511734,
	124	+ "max_abs_error": 1.20703125
	125	+ },
	126	+ "native_scale_alg2_error": null,
	127	+ "tilelang_planned_macaq": {
	128	+ "nmse": 0.005795284495877784,
	129	+ "nrmse": 0.07612676596229334,
	130	+ "sqnr_db": 22.36925239259129,
	131	+ "cosine": 0.9972342905702456,
	132	+ "max_abs_error": 0.46875
	133	+ },
	134	+ "tilelang_macaq": {
	135	+ "nmse": 0.005795284495877784,
	136	+ "nrmse": 0.07612676596229334,
	137	+ "sqnr_db": 22.36925239259129,
	138	+ "cosine": 0.9972342905702456,
	139	+ "max_abs_error": 0.46875
	140	+ },
	141	+ "opsnn_macaq": {
	142	+ "nmse": 0.005795284495877784,
	143	+ "nrmse": 0.07612676596229334,
	144	+ "sqnr_db": 22.36925239259129,
	145	+ "cosine": 0.9972342905702456,
	146	+ "max_abs_error": 0.46875
	147	+ },
	148	+ "cpu_reference_agreement": {
	149	+ "plain_max_abs": 0.0,
	150	+ "tilelang_max_abs": 0.0,
	151	+ "tilelang_planned_max_abs": 0.0,
	152	+ "opsnn_max_abs": 0.0
	153	+ }
	154	+ },
	155	+ {
	156	+ "fixture": "accuracy-fixture-model_layers_14_mlp_down_proj.pt",
	157	+ "fixture_sha256": "06327450f71e490385a6137915d300273b901ed344017c98b3cfaf8a7061b579",
	158	+ "m": 1024,
	159	+ "k": 3072,
	160	+ "n": 1024,
	161	+ "opsnn_padded_k": 3072,
	162	+ "opsnn_padded_n": 7168,
	163	+ "plain_native": {
	164	+ "nmse": 0.03410957884917448,
	165	+ "nrmse": 0.18468778749331122,
	166	+ "sqnr_db": 14.671236427916272,
	167	+ "cosine": 0.9830276128202399,
	168	+ "max_abs_error": 1.5625
	169	+ },
	170	+ "native_scale_alg2": {
	171	+ "nmse": 0.03799854210203062,
	172	+ "nrmse": 0.19493214743092177,
	173	+ "sqnr_db": 14.202330657302781,
	174	+ "cosine": 0.9809640218643827,
	175	+ "max_abs_error": 1.8125
	176	+ },
	177	+ "native_scale_alg2_error": null,
	178	+ "tilelang_planned_macaq": {
	179	+ "nmse": 0.021642116309140973,
	180	+ "nrmse": 0.1471125973842518,
	181	+ "sqnr_db": 16.647002733047213,
	182	+ "cosine": 0.989176496558601,
	183	+ "max_abs_error": 0.7265625
	184	+ },
	185	+ "tilelang_macaq": {
	186	+ "nmse": 0.021642116309140973,
	187	+ "nrmse": 0.1471125973842518,
	188	+ "sqnr_db": 16.647002733047213,
	189	+ "cosine": 0.989176496558601,
	190	+ "max_abs_error": 0.7265625
	191	+ },
	192	+ "opsnn_macaq": {
	193	+ "nmse": 0.021642116309140973,
	194	+ "nrmse": 0.1471125973842518,
	195	+ "sqnr_db": 16.647002733047213,
	196	+ "cosine": 0.989176496558601,
	197	+ "max_abs_error": 0.7265625
	198	+ },
	199	+ "cpu_reference_agreement": {
	200	+ "plain_max_abs": 0.0,
	201	+ "tilelang_max_abs": 0.0,
	202	+ "tilelang_planned_max_abs": 0.0,
	203	+ "opsnn_max_abs": 0.0
	204	+ }
	205	+ },
	206	+ {
	207	+ "fixture": "accuracy-fixture-model_layers_14_mlp_up_proj.pt",
	208	+ "fixture_sha256": "93015e6f4ed06d62306d6fccfc75c3fa8a2462740a20c61c72c8b4871668986d",
	209	+ "m": 1024,
	210	+ "k": 1024,
	211	+ "n": 3072,
	212	+ "opsnn_padded_k": 2048,
	213	+ "opsnn_padded_n": 3072,
	214	+ "plain_native": {
	215	+ "nmse": 0.027274659838753348,
	216	+ "nrmse": 0.16515041579951698,
	217	+ "sqnr_db": 15.642406570942896,
	218	+ "cosine": 0.9865070511505231,
	219	+ "max_abs_error": 1.109375
	220	+ },
	221	+ "native_scale_alg2": {
	222	+ "nmse": 0.030852923399939662,
	223	+ "nrmse": 0.17565000256174113,
	224	+ "sqnr_db": 15.107036790761057,
	225	+ "cosine": 0.9845036092016929,
	226	+ "max_abs_error": 1.03125
	227	+ },
	228	+ "native_scale_alg2_error": null,
	229	+ "tilelang_planned_macaq": {
	230	+ "nmse": 0.01696913017803275,
	231	+ "nrmse": 0.13026561395100686,
	232	+ "sqnr_db": 17.7034041864651,
	233	+ "cosine": 0.9915783593889278,
	234	+ "max_abs_error": 0.671875
	235	+ },
	236	+ "tilelang_macaq": {
	237	+ "nmse": 0.01696913017803275,
	238	+ "nrmse": 0.13026561395100686,
	239	+ "sqnr_db": 17.7034041864651,
	240	+ "cosine": 0.9915783593889278,
	241	+ "max_abs_error": 0.671875
	242	+ },
	243	+ "opsnn_macaq": {
	244	+ "nmse": 0.01696913017803275,
	245	+ "nrmse": 0.13026561395100686,
	246	+ "sqnr_db": 17.7034041864651,
	247	+ "cosine": 0.9915783593889278,
	248	+ "max_abs_error": 0.671875
	249	+ },
	250	+ "cpu_reference_agreement": {
	251	+ "plain_max_abs": 0.0,
	252	+ "tilelang_max_abs": 0.0,
	253	+ "tilelang_planned_max_abs": 0.0,
	254	+ "opsnn_max_abs": 0.0
	255	+ }
	256	+ },
	257	+ {
	258	+ "fixture": "accuracy-fixture-model_layers_14_self_attn_q_proj.pt",
	259	+ "fixture_sha256": "abe62ae90b968c85dfc57dbfa8bd287a1672b9a7ba9431570a121b1cff9fbb28",
	260	+ "m": 1024,
	261	+ "k": 1024,
	262	+ "n": 2048,
	263	+ "opsnn_padded_k": 1536,
	264	+ "opsnn_padded_n": 2048,
	265	+ "plain_native": {
	266	+ "nmse": 0.01466745110141977,
	267	+ "nrmse": 0.12110925274899424,
	268	+ "sqnr_db": 18.33645350967023,
	269	+ "cosine": 0.9930062509087988,
	270	+ "max_abs_error": 2.640625
	271	+ },
	272	+ "native_scale_alg2": {
	273	+ "nmse": 0.01625560235156519,
	274	+ "nrmse": 0.12749746017691957,
	275	+ "sqnr_db": 17.88996933081782,
	276	+ "cosine": 0.9918593472312802,
	277	+ "max_abs_error": 3.3125
	278	+ },
	279	+ "native_scale_alg2_error": null,
	280	+ "tilelang_planned_macaq": {
	281	+ "nmse": 0.0075766859551332265,
	282	+ "nrmse": 0.08704416094795346,
	283	+ "sqnr_db": 21.205207133813268,
	284	+ "cosine": 0.9962798805014916,
	285	+ "max_abs_error": 2.1171875
	286	+ },
	287	+ "tilelang_macaq": {
	288	+ "nmse": 0.0075766859551332265,
	289	+ "nrmse": 0.08704416094795346,
	290	+ "sqnr_db": 21.205207133813268,
	291	+ "cosine": 0.9962798805014916,
	292	+ "max_abs_error": 2.1171875
	293	+ },
	294	+ "opsnn_macaq": {
	295	+ "nmse": 0.0075766859551332265,
	296	+ "nrmse": 0.08704416094795346,
	297	+ "sqnr_db": 21.205207133813268,
	298	+ "cosine": 0.9962798805014916,
	299	+ "max_abs_error": 2.1171875
	300	+ },
	301	+ "cpu_reference_agreement": {
	302	+ "plain_max_abs": 0.0,
	303	+ "tilelang_max_abs": 0.0,
	304	+ "tilelang_planned_max_abs": 0.0,
	305	+ "opsnn_max_abs": 0.0
	306	+ }
	307	+ },
	308	+ {
	309	+ "fixture": "accuracy-fixture-model_layers_27_mlp_down_proj.pt",
	310	+ "fixture_sha256": "18f102d83016adb97aa542df1515be28ab7acea4059b91fcd8acb22427382c25",
	311	+ "m": 1024,
	312	+ "k": 3072,
	313	+ "n": 1024,
	314	+ "opsnn_padded_k": 3072,
	315	+ "opsnn_padded_n": 7168,
	316	+ "plain_native": {
	317	+ "nmse": 0.021753115704479336,
	318	+ "nrmse": 0.14748937488673322,
	319	+ "sqnr_db": 16.62478530141382,
	320	+ "cosine": 0.9903826287274079,
	321	+ "max_abs_error": 720.0
	322	+ },
	323	+ "native_scale_alg2": {
	324	+ "nmse": 0.019562300238837857,
	325	+ "nrmse": 0.13986529318897473,
	326	+ "sqnr_db": 17.085800799012137,
	327	+ "cosine": 0.990239036351425,
	328	+ "max_abs_error": 192.0
	329	+ },
	330	+ "native_scale_alg2_error": null,
	331	+ "tilelang_planned_macaq": {
	332	+ "nmse": 0.0047220473550336605,
	333	+ "nrmse": 0.06871715473616222,
	334	+ "sqnr_db": 23.258696619149188,
	335	+ "cosine": 0.9980580611788563,
	336	+ "max_abs_error": 288.0
	337	+ },
	338	+ "tilelang_macaq": {
	339	+ "nmse": 0.0047220473550336605,
	340	+ "nrmse": 0.06871715473616222,
	341	+ "sqnr_db": 23.258696619149188,
	342	+ "cosine": 0.9980580611788563,
	343	+ "max_abs_error": 288.0
	344	+ },
	345	+ "opsnn_macaq": {
	346	+ "nmse": 0.0047220473550336605,
	347	+ "nrmse": 0.06871715473616222,
	348	+ "sqnr_db": 23.258696619149188,
	349	+ "cosine": 0.9980580611788563,
	350	+ "max_abs_error": 288.0
	351	+ },
	352	+ "cpu_reference_agreement": {
	353	+ "plain_max_abs": 0.0,
	354	+ "tilelang_max_abs": 0.0,
	355	+ "tilelang_planned_max_abs": 0.0,
	356	+ "opsnn_max_abs": 0.0
	357	+ }
	358	+ },
	359	+ {
	360	+ "fixture": "accuracy-fixture-model_layers_27_mlp_up_proj.pt",
	361	+ "fixture_sha256": "e6ec478d4b4190fdef10af43e0f09cf14b869eb66bb3dcfa80f91bf0a6681297",
	362	+ "m": 1024,
	363	+ "k": 1024,
	364	+ "n": 3072,
	365	+ "opsnn_padded_k": 2048,
	366	+ "opsnn_padded_n": 3072,
	367	+ "plain_native": {
	368	+ "nmse": 0.021787974624019088,
	369	+ "nrmse": 0.1476075019232393,
	370	+ "sqnr_db": 16.61783139211963,
	371	+ "cosine": 0.9896387374240014,
	372	+ "max_abs_error": 53.0
	373	+ },
	374	+ "native_scale_alg2": {
	375	+ "nmse": 0.025910815855607353,
	376	+ "nrmse": 0.16096836911520024,
	377	+ "sqnr_db": 15.865189121392245,
	378	+ "cosine": 0.9869727854323713,
	379	+ "max_abs_error": 33.734375
	380	+ },
	381	+ "native_scale_alg2_error": null,
	382	+ "tilelang_planned_macaq": {
	383	+ "nmse": 0.008336486324689736,
	384	+ "nrmse": 0.09130436092920062,
	385	+ "sqnr_db": 20.790169579159002,
	386	+ "cosine": 0.9959198014115227,
	387	+ "max_abs_error": 39.0
	388	+ },
	389	+ "tilelang_macaq": {
	390	+ "nmse": 0.008336486324689736,
	391	+ "nrmse": 0.09130436092920062,
	392	+ "sqnr_db": 20.790169579159002,
	393	+ "cosine": 0.9959198014115227,
	394	+ "max_abs_error": 39.0
	395	+ },
	396	+ "opsnn_macaq": {
	397	+ "nmse": 0.008336486324689736,
	398	+ "nrmse": 0.09130436092920062,
	399	+ "sqnr_db": 20.790169579159002,
	400	+ "cosine": 0.9959198014115227,
	401	+ "max_abs_error": 39.0
	402	+ },
	403	+ "cpu_reference_agreement": {
	404	+ "plain_max_abs": 0.0,
	405	+ "tilelang_max_abs": 0.0,
	406	+ "tilelang_planned_max_abs": 0.0,
	407	+ "opsnn_max_abs": 0.0
	408	+ }
	409	+ },
	410	+ {
	411	+ "fixture": "accuracy-fixture-model_layers_27_self_attn_q_proj.pt",
	412	+ "fixture_sha256": "2ea2b3baf308ff4f1f81f978b28137ef5e98c111166fc24e8b2d6bb9aab7f391",
	413	+ "m": 1024,
	414	+ "k": 1024,
	415	+ "n": 2048,
	416	+ "opsnn_padded_k": 1536,
	417	+ "opsnn_padded_n": 2048,
	418	+ "plain_native": {
	419	+ "nmse": 0.014878039646034468,
	420	+ "nrmse": 0.12197556987378443,
	421	+ "sqnr_db": 18.27454288346345,
	422	+ "cosine": 0.9928836409767186,
	423	+ "max_abs_error": 32.5
	424	+ },
	425	+ "native_scale_alg2": {
	426	+ "nmse": 0.017160446816167834,
	427	+ "nrmse": 0.13099788859431222,
	428	+ "sqnr_db": 17.654714083713575,
	429	+ "cosine": 0.9913830249456285,
	430	+ "max_abs_error": 30.0
	431	+ },
	432	+ "native_scale_alg2_error": null,
	433	+ "tilelang_planned_macaq": {
	434	+ "nmse": 0.007906629422839765,
	435	+ "nrmse": 0.08891922976971722,
	436	+ "sqnr_db": 21.02008615747497,
	437	+ "cosine": 0.9961864477162093,
	438	+ "max_abs_error": 17.75
	439	+ },
	440	+ "tilelang_macaq": {
	441	+ "nmse": 0.007906629422839765,
	442	+ "nrmse": 0.08891922976971722,
	443	+ "sqnr_db": 21.02008615747497,
	444	+ "cosine": 0.9961864477162093,
	445	+ "max_abs_error": 17.75
	446	+ },
	447	+ "opsnn_macaq": {
	448	+ "nmse": 0.007906629422839765,
	449	+ "nrmse": 0.08891922976971722,
	450	+ "sqnr_db": 21.02008615747497,
	451	+ "cosine": 0.9961864477162093,
	452	+ "max_abs_error": 17.75
	453	+ },
	454	+ "cpu_reference_agreement": {
	455	+ "plain_max_abs": 0.0,
	456	+ "tilelang_max_abs": 0.0,
	457	+ "tilelang_planned_max_abs": 0.0,
	458	+ "opsnn_max_abs": 0.0
	459	+ }
	460	+ }
	461	+]

Adocs/benchmarks/macaq_20261027/accuracy-npu-dynamic-final.metadata.json+39-0

@@ -0,0 +1,39 @@
	1	+{
	2	+ "device": "Ascend950DT_9592",
	3	+ "torch": "2.9.0+cpu",
	4	+ "torch_npu": "2.9.0.post2",
	5	+ "opsnn_wheel_sha256": "b284336b70f82a80a4c2feba59bbc7f53cf178563f988dd89e527ef44ad50e06",
	6	+ "source_cpp_sha256": {
	7	+ "experimental/matmul/common/macaq/macaq_hadamard_mxfp4.hpp": "c62ab428e2d712d38b44e2e2de71605210e73b7fc4edac087a96a8aeb48c3574",
	8	+ "experimental/matmul/common/macaq/macaq_w4a4_kernel_impl.hpp": "52a5ed13b6798df7b36cdc08335663b2ffb70b52bda12180712b0dbb1ea1c26c",
	9	+ "experimental/matmul/macaq_w4a4_matmul/macaq_w4a4_matmul.cpp": "2ddea3d5dc3eed689a3aa59a43e219247b47344180a8ea23c24f94316a05ff44",
	10	+ "experimental/matmul/macaq_w4a4_grouped_matmul/macaq_w4a4_grouped_matmul.cpp": "63637c9de33517b69800b4dd0d748c17cbe5910facab4e5db39de7f931d38c93",
	11	+ "experimental/matmul/common/macaq/tests/test_macaq_raw_diagonal.py": "72f7a4536734be62e733656ce6cc45bd4b445529d995a224bc6cf309d24c7ed7"
	12	+ },
	13	+ "tilelang_base_commit": "b7d8b14c1e14f2412e4a37ccecde85e6284d0349",
	14	+ "model": "Qwen/Qwen3-0.6B",
	15	+ "checks": 36,
	16	+ "native_scale_alg2_checks": 9,
	17	+ "tilelang_source_sha256": {
	18	+ "__init__.py": "a81e80bbf9fd868d5109d2ab725dec9afe16d0a346ccad7809e55fb6dddfd2a7",
	19	+ "matmul/__init__.py": "13c8e8bdf2323d8787bcdf944464e1975343b35fe95a95e40d8d5612cec5a1eb",
	20	+ "matmul/macaq_w4a4_matmul.py": "039334d011a6c5714be28066c1dc12a00868ac4a9fea1c640ed974a61bf830b8",
	21	+ "matmul/macaq_w4a4_matmul_kernel.py": "fb33d479bb11db9ca3617c029d8c40dc7135af1baf0571346b6add11884bd777",
	22	+ "matmul/macaq_w4a4_global_kernel.py": "128169df74198a73da97292edc8d9aefe9caee6cfe70c3c2b2b88fe024c0d65c",
	23	+ "matmul/macaq_w4a4_dynamic_group_kernel.py": "c4a61194f87466de63e30b5f412146ce08f14eab7c5a1887cf83a2328b39212a"
	24	+ },
	25	+ "scope": "nine real projection fixtures, actual native and MACaQ NPU outputs vs CPU algorithm references; not whole-model NPU evaluation or model throughput",
	26	+ "padding": "ops-nn zero-pads K/N to supported pairs and slices original columns; zero padding preserves block32 transform boundaries; no latency claim for padded calls",
	27	+ "fixture_sha256": {
	28	+ "accuracy-fixture-model_layers_14_mlp_up_proj.pt": "93015e6f4ed06d62306d6fccfc75c3fa8a2462740a20c61c72c8b4871668986d",
	29	+ "accuracy-fixture-model_layers_14_mlp_down_proj.pt": "06327450f71e490385a6137915d300273b901ed344017c98b3cfaf8a7061b579",
	30	+ "accuracy-fixture-model_layers_14_self_attn_q_proj.pt": "abe62ae90b968c85dfc57dbfa8bd287a1672b9a7ba9431570a121b1cff9fbb28",
	31	+ "accuracy-fixture-model_layers_27_self_attn_q_proj.pt": "2ea2b3baf308ff4f1f81f978b28137ef5e98c111166fc24e8b2d6bb9aab7f391",
	32	+ "accuracy-fixture-model_layers_27_mlp_up_proj.pt": "e6ec478d4b4190fdef10af43e0f09cf14b869eb66bb3dcfa80f91bf0a6681297",
	33	+ "accuracy-fixture-model_layers_27_mlp_down_proj.pt": "18f102d83016adb97aa542df1515be28ab7acea4059b91fcd8acb22427382c25",
	34	+ "accuracy-fixture-model_layers_0_self_attn_q_proj.pt": "6ddc2c4ab7cccf3de0b083bb5ccb22cde417e5b4ed3ce81cd2659a579b18076c",
	35	+ "accuracy-fixture-model_layers_0_mlp_down_proj.pt": "4f73bdf7d8a927172fea0866e27cbe850c894deb329930bd08b99aab6cff444c",
	36	+ "accuracy-fixture-model_layers_0_mlp_up_proj.pt": "03b5ea72441a092958e49baa3ec5fe84ec79cd897de2d7d1e3a4e54a7581d92a"
	37	+ },
	38	+ "harness_sha256": "0daec03479f004eebb7e1099012d907933819f67bf76c916a9d3c180ce9d96ef"
	39	+}

Adocs/benchmarks/macaq_20261027/accuracy-npu-global.json+461-0

@@ -0,0 +1,461 @@
	1	+[
	2	+ {
	3	+ "fixture": "accuracy-fixture-model_layers_0_mlp_down_proj.pt",
	4	+ "fixture_sha256": "4f73bdf7d8a927172fea0866e27cbe850c894deb329930bd08b99aab6cff444c",
	5	+ "m": 1024,
	6	+ "k": 3072,
	7	+ "n": 1024,
	8	+ "opsnn_padded_k": 3072,
	9	+ "opsnn_padded_n": 7168,
	10	+ "plain_native": {
	11	+ "nmse": 0.030516970236956067,
	12	+ "nrmse": 0.17469107085640087,
	13	+ "sqnr_db": 15.154585858772782,
	14	+ "cosine": 0.9852403571117876,
	15	+ "max_abs_error": 0.25
	16	+ },
	17	+ "native_scale_alg2": {
	18	+ "nmse": 0.031153790077059946,
	19	+ "nrmse": 0.1765043627706124,
	20	+ "sqnr_db": 15.064891108210311,
	21	+ "cosine": 0.9843569751092435,
	22	+ "max_abs_error": 0.19921875
	23	+ },
	24	+ "native_scale_alg2_error": null,
	25	+ "tilelang_planned_macaq": {
	26	+ "nmse": 0.018508736055048453,
	27	+ "nrmse": 0.13604681567404822,
	28	+ "sqnr_db": 17.326232378143693,
	29	+ "cosine": 0.9908084532765248,
	30	+ "max_abs_error": 0.248046875
	31	+ },
	32	+ "tilelang_macaq": {
	33	+ "nmse": 0.018508736055048453,
	34	+ "nrmse": 0.13604681567404822,
	35	+ "sqnr_db": 17.326232378143693,
	36	+ "cosine": 0.9908084532765248,
	37	+ "max_abs_error": 0.248046875
	38	+ },
	39	+ "opsnn_macaq": {
	40	+ "nmse": 0.018508736055048453,
	41	+ "nrmse": 0.13604681567404822,
	42	+ "sqnr_db": 17.326232378143693,
	43	+ "cosine": 0.9908084532765248,
	44	+ "max_abs_error": 0.248046875
	45	+ },
	46	+ "cpu_reference_agreement": {
	47	+ "plain_max_abs": 0.0,
	48	+ "tilelang_max_abs": 0.0,
	49	+ "tilelang_planned_max_abs": 0.0,
	50	+ "opsnn_max_abs": 0.0
	51	+ }
	52	+ },
	53	+ {
	54	+ "fixture": "accuracy-fixture-model_layers_0_mlp_up_proj.pt",
	55	+ "fixture_sha256": "03b5ea72441a092958e49baa3ec5fe84ec79cd897de2d7d1e3a4e54a7581d92a",
	56	+ "m": 1024,
	57	+ "k": 1024,
	58	+ "n": 3072,
	59	+ "opsnn_padded_k": 2048,
	60	+ "opsnn_padded_n": 3072,
	61	+ "plain_native": {
	62	+ "nmse": 0.03272208139022038,
	63	+ "nrmse": 0.18089245807998844,
	64	+ "sqnr_db": 14.85159079464714,
	65	+ "cosine": 0.9837271166985933,
	66	+ "max_abs_error": 0.640625
	67	+ },
	68	+ "native_scale_alg2": {
	69	+ "nmse": 0.04013932103835252,
	70	+ "nrmse": 0.20034799983616638,
	71	+ "sqnr_db": 13.964299779320108,
	72	+ "cosine": 0.9798668105571983,
	73	+ "max_abs_error": 0.546875
	74	+ },
	75	+ "native_scale_alg2_error": null,
	76	+ "tilelang_planned_macaq": {
	77	+ "nmse": 0.018918712122045826,
	78	+ "nrmse": 0.13754530934221576,
	79	+ "sqnr_db": 17.231084312178428,
	80	+ "cosine": 0.9905704452715268,
	81	+ "max_abs_error": 0.3671875
	82	+ },
	83	+ "tilelang_macaq": {
	84	+ "nmse": 0.018918712122045826,
	85	+ "nrmse": 0.13754530934221576,
	86	+ "sqnr_db": 17.231084312178428,
	87	+ "cosine": 0.9905704452715268,
	88	+ "max_abs_error": 0.3671875
	89	+ },
	90	+ "opsnn_macaq": {
	91	+ "nmse": 0.018918712122045826,
	92	+ "nrmse": 0.13754530934221576,
	93	+ "sqnr_db": 17.231084312178428,
	94	+ "cosine": 0.9905704452715268,
	95	+ "max_abs_error": 0.3671875
	96	+ },
	97	+ "cpu_reference_agreement": {
	98	+ "plain_max_abs": 0.0,
	99	+ "tilelang_max_abs": 0.0,
	100	+ "tilelang_planned_max_abs": 0.0,
	101	+ "opsnn_max_abs": 0.0
	102	+ }
	103	+ },
	104	+ {
	105	+ "fixture": "accuracy-fixture-model_layers_0_self_attn_q_proj.pt",
	106	+ "fixture_sha256": "6ddc2c4ab7cccf3de0b083bb5ccb22cde417e5b4ed3ce81cd2659a579b18076c",
	107	+ "m": 1024,
	108	+ "k": 1024,
	109	+ "n": 2048,
	110	+ "opsnn_padded_k": 1536,
	111	+ "opsnn_padded_n": 2048,
	112	+ "plain_native": {
	113	+ "nmse": 0.012961810244357238,
	114	+ "nrmse": 0.11384994617634757,
	115	+ "sqnr_db": 18.873343407313314,
	116	+ "cosine": 0.9945141181596306,
	117	+ "max_abs_error": 1.015625
	118	+ },
	119	+ "native_scale_alg2": {
	120	+ "nmse": 0.011486990046041059,
	121	+ "nrmse": 0.10717737655886647,
	122	+ "sqnr_db": 19.39793755266881,
	123	+ "cosine": 0.9943654953511734,
	124	+ "max_abs_error": 1.20703125
	125	+ },
	126	+ "native_scale_alg2_error": null,
	127	+ "tilelang_planned_macaq": {
	128	+ "nmse": 0.005795284495877784,
	129	+ "nrmse": 0.07612676596229334,
	130	+ "sqnr_db": 22.36925239259129,
	131	+ "cosine": 0.9972342905702456,
	132	+ "max_abs_error": 0.46875
	133	+ },
	134	+ "tilelang_macaq": {
	135	+ "nmse": 0.005795284495877784,
	136	+ "nrmse": 0.07612676596229334,
	137	+ "sqnr_db": 22.36925239259129,
	138	+ "cosine": 0.9972342905702456,
	139	+ "max_abs_error": 0.46875
	140	+ },
	141	+ "opsnn_macaq": {
	142	+ "nmse": 0.005795284495877784,
	143	+ "nrmse": 0.07612676596229334,
	144	+ "sqnr_db": 22.36925239259129,
	145	+ "cosine": 0.9972342905702456,
	146	+ "max_abs_error": 0.46875
	147	+ },
	148	+ "cpu_reference_agreement": {
	149	+ "plain_max_abs": 0.0,
	150	+ "tilelang_max_abs": 0.0,
	151	+ "tilelang_planned_max_abs": 0.0,
	152	+ "opsnn_max_abs": 0.0
	153	+ }
	154	+ },
	155	+ {
	156	+ "fixture": "accuracy-fixture-model_layers_14_mlp_down_proj.pt",
	157	+ "fixture_sha256": "06327450f71e490385a6137915d300273b901ed344017c98b3cfaf8a7061b579",
	158	+ "m": 1024,
	159	+ "k": 3072,
	160	+ "n": 1024,
	161	+ "opsnn_padded_k": 3072,
	162	+ "opsnn_padded_n": 7168,
	163	+ "plain_native": {
	164	+ "nmse": 0.03410957884917448,
	165	+ "nrmse": 0.18468778749331122,
	166	+ "sqnr_db": 14.671236427916272,
	167	+ "cosine": 0.9830276128202399,
	168	+ "max_abs_error": 1.5625
	169	+ },
	170	+ "native_scale_alg2": {
	171	+ "nmse": 0.03799854210203062,
	172	+ "nrmse": 0.19493214743092177,
	173	+ "sqnr_db": 14.202330657302781,
	174	+ "cosine": 0.9809640218643827,
	175	+ "max_abs_error": 1.8125
	176	+ },
	177	+ "native_scale_alg2_error": null,
	178	+ "tilelang_planned_macaq": {
	179	+ "nmse": 0.021642116309140973,
	180	+ "nrmse": 0.1471125973842518,
	181	+ "sqnr_db": 16.647002733047213,
	182	+ "cosine": 0.989176496558601,
	183	+ "max_abs_error": 0.7265625
	184	+ },
	185	+ "tilelang_macaq": {
	186	+ "nmse": 0.021642116309140973,
	187	+ "nrmse": 0.1471125973842518,
	188	+ "sqnr_db": 16.647002733047213,
	189	+ "cosine": 0.989176496558601,
	190	+ "max_abs_error": 0.7265625
	191	+ },
	192	+ "opsnn_macaq": {
	193	+ "nmse": 0.021642116309140973,
	194	+ "nrmse": 0.1471125973842518,
	195	+ "sqnr_db": 16.647002733047213,
	196	+ "cosine": 0.989176496558601,
	197	+ "max_abs_error": 0.7265625
	198	+ },
	199	+ "cpu_reference_agreement": {
	200	+ "plain_max_abs": 0.0,
	201	+ "tilelang_max_abs": 0.0,
	202	+ "tilelang_planned_max_abs": 0.0,
	203	+ "opsnn_max_abs": 0.0
	204	+ }
	205	+ },
	206	+ {
	207	+ "fixture": "accuracy-fixture-model_layers_14_mlp_up_proj.pt",
	208	+ "fixture_sha256": "93015e6f4ed06d62306d6fccfc75c3fa8a2462740a20c61c72c8b4871668986d",
	209	+ "m": 1024,
	210	+ "k": 1024,
	211	+ "n": 3072,
	212	+ "opsnn_padded_k": 2048,
	213	+ "opsnn_padded_n": 3072,
	214	+ "plain_native": {
	215	+ "nmse": 0.027274659838753348,
	216	+ "nrmse": 0.16515041579951698,
	217	+ "sqnr_db": 15.642406570942896,
	218	+ "cosine": 0.9865070511505231,
	219	+ "max_abs_error": 1.109375
	220	+ },
	221	+ "native_scale_alg2": {
	222	+ "nmse": 0.030852923399939662,
	223	+ "nrmse": 0.17565000256174113,
	224	+ "sqnr_db": 15.107036790761057,
	225	+ "cosine": 0.9845036092016929,
	226	+ "max_abs_error": 1.03125
	227	+ },
	228	+ "native_scale_alg2_error": null,
	229	+ "tilelang_planned_macaq": {
	230	+ "nmse": 0.01696913017803275,
	231	+ "nrmse": 0.13026561395100686,
	232	+ "sqnr_db": 17.7034041864651,
	233	+ "cosine": 0.9915783593889278,
	234	+ "max_abs_error": 0.671875
	235	+ },
	236	+ "tilelang_macaq": {
	237	+ "nmse": 0.01696913017803275,
	238	+ "nrmse": 0.13026561395100686,
	239	+ "sqnr_db": 17.7034041864651,
	240	+ "cosine": 0.9915783593889278,
	241	+ "max_abs_error": 0.671875
	242	+ },
	243	+ "opsnn_macaq": {
	244	+ "nmse": 0.01696913017803275,
	245	+ "nrmse": 0.13026561395100686,
	246	+ "sqnr_db": 17.7034041864651,
	247	+ "cosine": 0.9915783593889278,
	248	+ "max_abs_error": 0.671875
	249	+ },
	250	+ "cpu_reference_agreement": {
	251	+ "plain_max_abs": 0.0,
	252	+ "tilelang_max_abs": 0.0,
	253	+ "tilelang_planned_max_abs": 0.0,
	254	+ "opsnn_max_abs": 0.0
	255	+ }
	256	+ },
	257	+ {
	258	+ "fixture": "accuracy-fixture-model_layers_14_self_attn_q_proj.pt",
	259	+ "fixture_sha256": "abe62ae90b968c85dfc57dbfa8bd287a1672b9a7ba9431570a121b1cff9fbb28",
	260	+ "m": 1024,
	261	+ "k": 1024,
	262	+ "n": 2048,
	263	+ "opsnn_padded_k": 1536,
	264	+ "opsnn_padded_n": 2048,
	265	+ "plain_native": {
	266	+ "nmse": 0.01466745110141977,
	267	+ "nrmse": 0.12110925274899424,
	268	+ "sqnr_db": 18.33645350967023,
	269	+ "cosine": 0.9930062509087988,
	270	+ "max_abs_error": 2.640625
	271	+ },
	272	+ "native_scale_alg2": {
	273	+ "nmse": 0.01625560235156519,
	274	+ "nrmse": 0.12749746017691957,
	275	+ "sqnr_db": 17.88996933081782,
	276	+ "cosine": 0.9918593472312802,
	277	+ "max_abs_error": 3.3125
	278	+ },
	279	+ "native_scale_alg2_error": null,
	280	+ "tilelang_planned_macaq": {
	281	+ "nmse": 0.0075766859551332265,
	282	+ "nrmse": 0.08704416094795346,
	283	+ "sqnr_db": 21.205207133813268,
	284	+ "cosine": 0.9962798805014916,
	285	+ "max_abs_error": 2.1171875
	286	+ },
	287	+ "tilelang_macaq": {
	288	+ "nmse": 0.0075766859551332265,
	289	+ "nrmse": 0.08704416094795346,
	290	+ "sqnr_db": 21.205207133813268,
	291	+ "cosine": 0.9962798805014916,
	292	+ "max_abs_error": 2.1171875
	293	+ },
	294	+ "opsnn_macaq": {
	295	+ "nmse": 0.0075766859551332265,
	296	+ "nrmse": 0.08704416094795346,
	297	+ "sqnr_db": 21.205207133813268,
	298	+ "cosine": 0.9962798805014916,
	299	+ "max_abs_error": 2.1171875
	300	+ },
	301	+ "cpu_reference_agreement": {
	302	+ "plain_max_abs": 0.0,
	303	+ "tilelang_max_abs": 0.0,
	304	+ "tilelang_planned_max_abs": 0.0,
	305	+ "opsnn_max_abs": 0.0
	306	+ }
	307	+ },
	308	+ {
	309	+ "fixture": "accuracy-fixture-model_layers_27_mlp_down_proj.pt",
	310	+ "fixture_sha256": "18f102d83016adb97aa542df1515be28ab7acea4059b91fcd8acb22427382c25",
	311	+ "m": 1024,
	312	+ "k": 3072,
	313	+ "n": 1024,
	314	+ "opsnn_padded_k": 3072,
	315	+ "opsnn_padded_n": 7168,
	316	+ "plain_native": {
	317	+ "nmse": 0.021753115704479336,
	318	+ "nrmse": 0.14748937488673322,
	319	+ "sqnr_db": 16.62478530141382,
	320	+ "cosine": 0.9903826287274079,
	321	+ "max_abs_error": 720.0
	322	+ },
	323	+ "native_scale_alg2": {
	324	+ "nmse": 0.019562300238837857,
	325	+ "nrmse": 0.13986529318897473,
	326	+ "sqnr_db": 17.085800799012137,
	327	+ "cosine": 0.990239036351425,
	328	+ "max_abs_error": 192.0
	329	+ },
	330	+ "native_scale_alg2_error": null,
	331	+ "tilelang_planned_macaq": {
	332	+ "nmse": 0.0047220473550336605,
	333	+ "nrmse": 0.06871715473616222,
	334	+ "sqnr_db": 23.258696619149188,
	335	+ "cosine": 0.9980580611788563,
	336	+ "max_abs_error": 288.0
	337	+ },
	338	+ "tilelang_macaq": {
	339	+ "nmse": 0.0047220473550336605,
	340	+ "nrmse": 0.06871715473616222,
	341	+ "sqnr_db": 23.258696619149188,
	342	+ "cosine": 0.9980580611788563,
	343	+ "max_abs_error": 288.0
	344	+ },
	345	+ "opsnn_macaq": {
	346	+ "nmse": 0.0047220473550336605,
	347	+ "nrmse": 0.06871715473616222,
	348	+ "sqnr_db": 23.258696619149188,
	349	+ "cosine": 0.9980580611788563,
	350	+ "max_abs_error": 288.0
	351	+ },
	352	+ "cpu_reference_agreement": {
	353	+ "plain_max_abs": 0.0,
	354	+ "tilelang_max_abs": 0.0,
	355	+ "tilelang_planned_max_abs": 0.0,
	356	+ "opsnn_max_abs": 0.0
	357	+ }
	358	+ },
	359	+ {
	360	+ "fixture": "accuracy-fixture-model_layers_27_mlp_up_proj.pt",
	361	+ "fixture_sha256": "e6ec478d4b4190fdef10af43e0f09cf14b869eb66bb3dcfa80f91bf0a6681297",
	362	+ "m": 1024,
	363	+ "k": 1024,
	364	+ "n": 3072,
	365	+ "opsnn_padded_k": 2048,
	366	+ "opsnn_padded_n": 3072,
	367	+ "plain_native": {
	368	+ "nmse": 0.021787974624019088,
	369	+ "nrmse": 0.1476075019232393,
	370	+ "sqnr_db": 16.61783139211963,
	371	+ "cosine": 0.9896387374240014,
	372	+ "max_abs_error": 53.0
	373	+ },
	374	+ "native_scale_alg2": {
	375	+ "nmse": 0.025910815855607353,
	376	+ "nrmse": 0.16096836911520024,
	377	+ "sqnr_db": 15.865189121392245,
	378	+ "cosine": 0.9869727854323713,
	379	+ "max_abs_error": 33.734375
	380	+ },
	381	+ "native_scale_alg2_error": null,
	382	+ "tilelang_planned_macaq": {
	383	+ "nmse": 0.008336486324689736,
	384	+ "nrmse": 0.09130436092920062,
	385	+ "sqnr_db": 20.790169579159002,
	386	+ "cosine": 0.9959198014115227,
	387	+ "max_abs_error": 39.0
	388	+ },
	389	+ "tilelang_macaq": {
	390	+ "nmse": 0.008336486324689736,
	391	+ "nrmse": 0.09130436092920062,
	392	+ "sqnr_db": 20.790169579159002,
	393	+ "cosine": 0.9959198014115227,
	394	+ "max_abs_error": 39.0
	395	+ },
	396	+ "opsnn_macaq": {
	397	+ "nmse": 0.008336486324689736,
	398	+ "nrmse": 0.09130436092920062,
	399	+ "sqnr_db": 20.790169579159002,
	400	+ "cosine": 0.9959198014115227,
	401	+ "max_abs_error": 39.0
	402	+ },
	403	+ "cpu_reference_agreement": {
	404	+ "plain_max_abs": 0.0,
	405	+ "tilelang_max_abs": 0.0,
	406	+ "tilelang_planned_max_abs": 0.0,
	407	+ "opsnn_max_abs": 0.0
	408	+ }
	409	+ },
	410	+ {
	411	+ "fixture": "accuracy-fixture-model_layers_27_self_attn_q_proj.pt",
	412	+ "fixture_sha256": "2ea2b3baf308ff4f1f81f978b28137ef5e98c111166fc24e8b2d6bb9aab7f391",
	413	+ "m": 1024,
	414	+ "k": 1024,
	415	+ "n": 2048,
	416	+ "opsnn_padded_k": 1536,
	417	+ "opsnn_padded_n": 2048,
	418	+ "plain_native": {
	419	+ "nmse": 0.014878039646034468,
	420	+ "nrmse": 0.12197556987378443,
	421	+ "sqnr_db": 18.27454288346345,
	422	+ "cosine": 0.9928836409767186,
	423	+ "max_abs_error": 32.5
	424	+ },
	425	+ "native_scale_alg2": {
	426	+ "nmse": 0.017160446816167834,
	427	+ "nrmse": 0.13099788859431222,
	428	+ "sqnr_db": 17.654714083713575,
	429	+ "cosine": 0.9913830249456285,
	430	+ "max_abs_error": 30.0
	431	+ },
	432	+ "native_scale_alg2_error": null,
	433	+ "tilelang_planned_macaq": {
	434	+ "nmse": 0.007906629422839765,
	435	+ "nrmse": 0.08891922976971722,
	436	+ "sqnr_db": 21.02008615747497,
	437	+ "cosine": 0.9961864477162093,
	438	+ "max_abs_error": 17.75
	439	+ },
	440	+ "tilelang_macaq": {
	441	+ "nmse": 0.007906629422839765,
	442	+ "nrmse": 0.08891922976971722,
	443	+ "sqnr_db": 21.02008615747497,
	444	+ "cosine": 0.9961864477162093,
	445	+ "max_abs_error": 17.75
	446	+ },
	447	+ "opsnn_macaq": {
	448	+ "nmse": 0.007906629422839765,
	449	+ "nrmse": 0.08891922976971722,
	450	+ "sqnr_db": 21.02008615747497,
	451	+ "cosine": 0.9961864477162093,
	452	+ "max_abs_error": 17.75
	453	+ },
	454	+ "cpu_reference_agreement": {
	455	+ "plain_max_abs": 0.0,
	456	+ "tilelang_max_abs": 0.0,
	457	+ "tilelang_planned_max_abs": 0.0,
	458	+ "opsnn_max_abs": 0.0
	459	+ }
	460	+ }
	461	+]

Adocs/benchmarks/macaq_20261027/accuracy-npu-global.metadata.json+38-0

@@ -0,0 +1,38 @@
	1	+{
	2	+ "device": "Ascend950DT_9592",
	3	+ "torch": "2.9.0+cpu",
	4	+ "torch_npu": "2.9.0.post2",
	5	+ "opsnn_wheel_sha256": "b284336b70f82a80a4c2feba59bbc7f53cf178563f988dd89e527ef44ad50e06",
	6	+ "source_cpp_sha256": {
	7	+ "experimental/matmul/common/macaq/macaq_hadamard_mxfp4.hpp": "c62ab428e2d712d38b44e2e2de71605210e73b7fc4edac087a96a8aeb48c3574",
	8	+ "experimental/matmul/common/macaq/macaq_w4a4_kernel_impl.hpp": "52a5ed13b6798df7b36cdc08335663b2ffb70b52bda12180712b0dbb1ea1c26c",
	9	+ "experimental/matmul/macaq_w4a4_matmul/macaq_w4a4_matmul.cpp": "2ddea3d5dc3eed689a3aa59a43e219247b47344180a8ea23c24f94316a05ff44",
	10	+ "experimental/matmul/macaq_w4a4_grouped_matmul/macaq_w4a4_grouped_matmul.cpp": "63637c9de33517b69800b4dd0d748c17cbe5910facab4e5db39de7f931d38c93",
	11	+ "experimental/matmul/common/macaq/tests/test_macaq_raw_diagonal.py": "72f7a4536734be62e733656ce6cc45bd4b445529d995a224bc6cf309d24c7ed7"
	12	+ },
	13	+ "tilelang_base_commit": "b7d8b14c1e14f2412e4a37ccecde85e6284d0349",
	14	+ "model": "Qwen/Qwen3-0.6B",
	15	+ "checks": 36,
	16	+ "native_scale_alg2_checks": 9,
	17	+ "tilelang_source_sha256": {
	18	+ "__init__.py": "a81e80bbf9fd868d5109d2ab725dec9afe16d0a346ccad7809e55fb6dddfd2a7",
	19	+ "matmul/__init__.py": "13c8e8bdf2323d8787bcdf944464e1975343b35fe95a95e40d8d5612cec5a1eb",
	20	+ "matmul/macaq_w4a4_matmul.py": "039334d011a6c5714be28066c1dc12a00868ac4a9fea1c640ed974a61bf830b8",
	21	+ "matmul/macaq_w4a4_matmul_kernel.py": "fb33d479bb11db9ca3617c029d8c40dc7135af1baf0571346b6add11884bd777",
	22	+ "matmul/macaq_w4a4_global_kernel.py": "8ad67823f4711ed04665a98bf2c00575a1604f1a7b4a4dd3df90dcfbfe707ebc"
	23	+ },
	24	+ "scope": "nine real projection fixtures, actual native and MACaQ NPU outputs vs CPU algorithm references; not whole-model NPU evaluation or model throughput",
	25	+ "padding": "ops-nn zero-pads K/N to supported pairs and slices original columns; zero padding preserves block32 transform boundaries; no latency claim for padded calls",
	26	+ "fixture_sha256": {
	27	+ "accuracy-fixture-model_layers_14_mlp_up_proj.pt": "93015e6f4ed06d62306d6fccfc75c3fa8a2462740a20c61c72c8b4871668986d",
	28	+ "accuracy-fixture-model_layers_14_mlp_down_proj.pt": "06327450f71e490385a6137915d300273b901ed344017c98b3cfaf8a7061b579",
	29	+ "accuracy-fixture-model_layers_14_self_attn_q_proj.pt": "abe62ae90b968c85dfc57dbfa8bd287a1672b9a7ba9431570a121b1cff9fbb28",
	30	+ "accuracy-fixture-model_layers_27_self_attn_q_proj.pt": "2ea2b3baf308ff4f1f81f978b28137ef5e98c111166fc24e8b2d6bb9aab7f391",
	31	+ "accuracy-fixture-model_layers_27_mlp_up_proj.pt": "e6ec478d4b4190fdef10af43e0f09cf14b869eb66bb3dcfa80f91bf0a6681297",
	32	+ "accuracy-fixture-model_layers_27_mlp_down_proj.pt": "18f102d83016adb97aa542df1515be28ab7acea4059b91fcd8acb22427382c25",
	33	+ "accuracy-fixture-model_layers_0_self_attn_q_proj.pt": "6ddc2c4ab7cccf3de0b083bb5ccb22cde417e5b4ed3ce81cd2659a579b18076c",
	34	+ "accuracy-fixture-model_layers_0_mlp_down_proj.pt": "4f73bdf7d8a927172fea0866e27cbe850c894deb329930bd08b99aab6cff444c",
	35	+ "accuracy-fixture-model_layers_0_mlp_up_proj.pt": "03b5ea72441a092958e49baa3ec5fe84ec79cd897de2d7d1e3a4e54a7581d92a"
	36	+ },
	37	+ "harness_sha256": "b6ec7a9498e46bdc07ecd93e2e2b1aa1a8e6fb4a2ff5b49405abae87967ce146"
	38	+}

Adocs/benchmarks/macaq_20261027/accuracy-npu-projections.json+132-0

@@ -0,0 +1,132 @@
	1	+[
	2	+ [
	3	+ "fixture",
	4	+ "M",
	5	+ "K",
	6	+ "N",
	7	+ "native_alg0_nmse",
	8	+ "native_alg2_nmse",
	9	+ "opsnn_nmse",
	10	+ "tilelang_nmse",
	11	+ "tilelang_prepared_nmse",
	12	+ "nmse_reduction_vs_alg0",
	13	+ "nmse_reduction_vs_alg2"
	14	+ ],
	15	+ [
	16	+ "accuracy-fixture-model_layers_0_mlp_down_proj.pt",
	17	+ "1024",
	18	+ "3072",
	19	+ "1024",
	20	+ "0.030516970236956067",
	21	+ "0.031153790077059946",
	22	+ "0.018508736055048453",
	23	+ "0.018508736055048453",
	24	+ "0.018508736055048453",
	25	+ "1.6487873697151914",
	26	+ "1.6831938163904185"
	27	+ ],
	28	+ [
	29	+ "accuracy-fixture-model_layers_0_mlp_up_proj.pt",
	30	+ "1024",
	31	+ "1024",
	32	+ "3072",
	33	+ "0.03272208139022038",
	34	+ "0.04013932103835252",
	35	+ "0.018918712122045826",
	36	+ "0.018918712122045826",
	37	+ "0.018918712122045826",
	38	+ "1.7296146365105685",
	39	+ "2.12167301766691"
	40	+ ],
	41	+ [
	42	+ "accuracy-fixture-model_layers_0_self_attn_q_proj.pt",
	43	+ "1024",
	44	+ "1024",
	45	+ "2048",
	46	+ "0.012961810244357238",
	47	+ "0.011486990046041059",
	48	+ "0.005795284495877784",
	49	+ "0.005795284495877784",
	50	+ "0.005795284495877784",
	51	+ "2.2366132764624482",
	52	+ "1.9821270300382692"
	53	+ ],
	54	+ [
	55	+ "accuracy-fixture-model_layers_14_mlp_down_proj.pt",
	56	+ "1024",
	57	+ "3072",
	58	+ "1024",
	59	+ "0.03410957884917448",
	60	+ "0.03799854210203062",
	61	+ "0.021642116309140973",
	62	+ "0.021642116309140973",
	63	+ "0.021642116309140973",
	64	+ "1.5760740937690845",
	65	+ "1.7557683157807074"
	66	+ ],
	67	+ [
	68	+ "accuracy-fixture-model_layers_14_mlp_up_proj.pt",
	69	+ "1024",
	70	+ "1024",
	71	+ "3072",
	72	+ "0.027274659838753348",
	73	+ "0.030852923399939662",
	74	+ "0.01696913017803275",
	75	+ "0.01696913017803275",
	76	+ "0.01696913017803275",
	77	+ "1.6073104250247037",
	78	+ "1.8181794279520622"
	79	+ ],
	80	+ [
	81	+ "accuracy-fixture-model_layers_14_self_attn_q_proj.pt",
	82	+ "1024",
	83	+ "1024",
	84	+ "2048",
	85	+ "0.01466745110141977",
	86	+ "0.01625560235156519",
	87	+ "0.0075766859551332265",
	88	+ "0.0075766859551332265",
	89	+ "0.0075766859551332265",
	90	+ "1.9358663125640743",
	91	+ "2.145476590665867"
	92	+ ],
	93	+ [
	94	+ "accuracy-fixture-model_layers_27_mlp_down_proj.pt",
	95	+ "1024",
	96	+ "3072",
	97	+ "1024",
	98	+ "0.021753115704479336",
	99	+ "0.019562300238837857",
	100	+ "0.0047220473550336605",
	101	+ "0.0047220473550336605",
	102	+ "0.0047220473550336605",
	103	+ "4.606712738975544",
	104	+ "4.142758165689427"
	105	+ ],
	106	+ [
	107	+ "accuracy-fixture-model_layers_27_mlp_up_proj.pt",
	108	+ "1024",
	109	+ "1024",
	110	+ "3072",
	111	+ "0.021787974624019088",
	112	+ "0.025910815855607353",
	113	+ "0.008336486324689736",
	114	+ "0.008336486324689736",
	115	+ "0.008336486324689736",
	116	+ "2.6135680879714016",
	117	+ "3.1081219168882512"
	118	+ ],
	119	+ [
	120	+ "accuracy-fixture-model_layers_27_self_attn_q_proj.pt",
	121	+ "1024",
	122	+ "1024",
	123	+ "2048",
	124	+ "0.014878039646034468",
	125	+ "0.017160446816167834",
	126	+ "0.007906629422839765",
	127	+ "0.007906629422839765",
	128	+ "0.007906629422839765",
	129	+ "1.881717082004184",
	130	+ "2.1703871395055776"
	131	+ ]
	132	+]

Adocs/benchmarks/macaq_20261027/accuracy-ppl-study.json+266-0

@@ -0,0 +1,266 @@
	1	+{
	2	+ "bf16": {
	3	+ "chunks": [
	4	+ {
	5	+ "chunk": 0,
	6	+ "scored_tokens": 1023,
	7	+ "nll": 2613.495315551758
	8	+ },
	9	+ {
	10	+ "chunk": 1,
	11	+ "scored_tokens": 1023,
	12	+ "nll": 3464.470916748047
	13	+ },
	14	+ {
	15	+ "chunk": 2,
	16	+ "scored_tokens": 1023,
	17	+ "nll": 3268.416778564453
	18	+ },
	19	+ {
	20	+ "chunk": 3,
	21	+ "scored_tokens": 1023,
	22	+ "nll": 3495.12353515625
	23	+ },
	24	+ {
	25	+ "chunk": 4,
	26	+ "scored_tokens": 1023,
	27	+ "nll": 3786.4422912597656
	28	+ },
	29	+ {
	30	+ "chunk": 5,
	31	+ "scored_tokens": 1023,
	32	+ "nll": 3405.630157470703
	33	+ },
	34	+ {
	35	+ "chunk": 6,
	36	+ "scored_tokens": 1023,
	37	+ "nll": 3722.502685546875
	38	+ },
	39	+ {
	40	+ "chunk": 7,
	41	+ "scored_tokens": 1023,
	42	+ "nll": 3582.7979125976562
	43	+ },
	44	+ {
	45	+ "chunk": 8,
	46	+ "scored_tokens": 1023,
	47	+ "nll": 2805.970687866211
	48	+ },
	49	+ {
	50	+ "chunk": 9,
	51	+ "scored_tokens": 1023,
	52	+ "nll": 3496.6533813476562
	53	+ },
	54	+ {
	55	+ "chunk": 10,
	56	+ "scored_tokens": 1023,
	57	+ "nll": 2920.6666870117188
	58	+ },
	59	+ {
	60	+ "chunk": 11,
	61	+ "scored_tokens": 1023,
	62	+ "nll": 2406.7943420410156
	63	+ },
	64	+ {
	65	+ "chunk": 12,
	66	+ "scored_tokens": 1023,
	67	+ "nll": 2533.4830017089844
	68	+ },
	69	+ {
	70	+ "chunk": 13,
	71	+ "scored_tokens": 1023,
	72	+ "nll": 2393.1245727539062
	73	+ },
	74	+ {
	75	+ "chunk": 14,
	76	+ "scored_tokens": 1023,
	77	+ "nll": 3081.54541015625
	78	+ },
	79	+ {
	80	+ "chunk": 15,
	81	+ "scored_tokens": 1023,
	82	+ "nll": 3096.2925720214844
	83	+ }
	84	+ ],
	85	+ "nll_sum": 50073.410247802734,
	86	+ "scored_tokens": 16368,
	87	+ "elapsed_cpu_seconds": 80.64238631998887,
	88	+ "perplexity": 21.31105482145711
	89	+ },
	90	+ "plain_mxfp4_w4a4": {
	91	+ "chunks": [
	92	+ {
	93	+ "chunk": 0,
	94	+ "scored_tokens": 1023,
	95	+ "nll": 2993.995475769043
	96	+ },
	97	+ {
	98	+ "chunk": 1,
	99	+ "scored_tokens": 1023,
	100	+ "nll": 3964.547393798828
	101	+ },
	102	+ {
	103	+ "chunk": 2,
	104	+ "scored_tokens": 1023,
	105	+ "nll": 3789.380096435547
	106	+ },
	107	+ {
	108	+ "chunk": 3,
	109	+ "scored_tokens": 1023,
	110	+ "nll": 3984.91650390625
	111	+ },
	112	+ {
	113	+ "chunk": 4,
	114	+ "scored_tokens": 1023,
	115	+ "nll": 4248.028869628906
	116	+ },
	117	+ {
	118	+ "chunk": 5,
	119	+ "scored_tokens": 1023,
	120	+ "nll": 3814.5045471191406
	121	+ },
	122	+ {
	123	+ "chunk": 6,
	124	+ "scored_tokens": 1023,
	125	+ "nll": 4285.850921630859
	126	+ },
	127	+ {
	128	+ "chunk": 7,
	129	+ "scored_tokens": 1023,
	130	+ "nll": 4122.245635986328
	131	+ },
	132	+ {
	133	+ "chunk": 8,
	134	+ "scored_tokens": 1023,
	135	+ "nll": 3248.092559814453
	136	+ },
	137	+ {
	138	+ "chunk": 9,
	139	+ "scored_tokens": 1023,
	140	+ "nll": 4024.273223876953
	141	+ },
	142	+ {
	143	+ "chunk": 10,
	144	+ "scored_tokens": 1023,
	145	+ "nll": 3406.539581298828
	146	+ },
	147	+ {
	148	+ "chunk": 11,
	149	+ "scored_tokens": 1023,
	150	+ "nll": 2750.1438598632812
	151	+ },
	152	+ {
	153	+ "chunk": 12,
	154	+ "scored_tokens": 1023,
	155	+ "nll": 2889.586441040039
	156	+ },
	157	+ {
	158	+ "chunk": 13,
	159	+ "scored_tokens": 1023,
	160	+ "nll": 2764.8697509765625
	161	+ },
	162	+ {
	163	+ "chunk": 14,
	164	+ "scored_tokens": 1023,
	165	+ "nll": 3561.8204956054688
	166	+ },
	167	+ {
	168	+ "chunk": 15,
	169	+ "scored_tokens": 1023,
	170	+ "nll": 3593.2583923339844
	171	+ }
	172	+ ],
	173	+ "nll_sum": 57442.05374908447,
	174	+ "scored_tokens": 16368,
	175	+ "elapsed_cpu_seconds": 100.81584860995645,
	176	+ "perplexity": 33.42860217079505
	177	+ },
	178	+ "macaq_w4a4": {
	179	+ "chunks": [
	180	+ {
	181	+ "chunk": 0,
	182	+ "scored_tokens": 1023,
	183	+ "nll": 2781.291793823242
	184	+ },
	185	+ {
	186	+ "chunk": 1,
	187	+ "scored_tokens": 1023,
	188	+ "nll": 3652.1476440429688
	189	+ },
	190	+ {
	191	+ "chunk": 2,
	192	+ "scored_tokens": 1023,
	193	+ "nll": 3466.021270751953
	194	+ },
	195	+ {
	196	+ "chunk": 3,
	197	+ "scored_tokens": 1023,
	198	+ "nll": 3660.3853759765625
	199	+ },
	200	+ {
	201	+ "chunk": 4,
	202	+ "scored_tokens": 1023,
	203	+ "nll": 3981.515594482422
	204	+ },
	205	+ {
	206	+ "chunk": 5,
	207	+ "scored_tokens": 1023,
	208	+ "nll": 3621.0074157714844
	209	+ },
	210	+ {
	211	+ "chunk": 6,
	212	+ "scored_tokens": 1023,
	213	+ "nll": 3980.6835021972656
	214	+ },
	215	+ {
	216	+ "chunk": 7,
	217	+ "scored_tokens": 1023,
	218	+ "nll": 3748.3545837402344
	219	+ },
	220	+ {
	221	+ "chunk": 8,
	222	+ "scored_tokens": 1023,
	223	+ "nll": 2963.4041137695312
	224	+ },
	225	+ {
	226	+ "chunk": 9,
	227	+ "scored_tokens": 1023,
	228	+ "nll": 3685.198028564453
	229	+ },
	230	+ {
	231	+ "chunk": 10,
	232	+ "scored_tokens": 1023,
	233	+ "nll": 3102.4110717773438
	234	+ },
	235	+ {
	236	+ "chunk": 11,
	237	+ "scored_tokens": 1023,
	238	+ "nll": 2573.624221801758
	239	+ },
	240	+ {
	241	+ "chunk": 12,
	242	+ "scored_tokens": 1023,
	243	+ "nll": 2685.5667114257812
	244	+ },
	245	+ {
	246	+ "chunk": 13,
	247	+ "scored_tokens": 1023,
	248	+ "nll": 2548.2555541992188
	249	+ },
	250	+ {
	251	+ "chunk": 14,
	252	+ "scored_tokens": 1023,
	253	+ "nll": 3272.051513671875
	254	+ },
	255	+ {
	256	+ "chunk": 15,
	257	+ "scored_tokens": 1023,
	258	+ "nll": 3286.328125
	259	+ }
	260	+ ],
	261	+ "nll_sum": 53008.246520996094,
	262	+ "scored_tokens": 16368,
	263	+ "elapsed_cpu_seconds": 102.51845485001104,
	264	+ "perplexity": 25.496195250101927
	265	+ }
	266	+}

Adocs/benchmarks/macaq_20261027/accuracy-ppl-study.metadata.json+21-0

@@ -0,0 +1,21 @@
	1	+{
	2	+ "model": "Qwen/Qwen3-0.6B",
	3	+ "revision": "c1899de289a04d12100db370d81485cdf75e47ca",
	4	+ "execution": "CPU reference quantization-model PPL; not actual NPU full-model execution",
	5	+ "dataset": "Salesforce/wikitext wikitext-2-raw-v1",
	6	+ "context": 1024,
	7	+ "calibration_tokens": 1024,
	8	+ "heldout_tokens": 16384,
	9	+ "scored_tokens": 16368,
	10	+ "token_sha256": {
	11	+ "train": "ea0c105aae5e1f26ceb3a387594b84f26c4e53dc1ddaccdc2f15d68751f917e9",
	12	+ "test": "a404e545152236c323bdaddd73c18b4b210bd67c5e5f4f811701db84f54d5217"
	13	+ },
	14	+ "quantized_decoder_linears": 196,
	15	+ "embedding_and_lm_head": "BF16 in all arms",
	16	+ "calibration": "BF16 decoder activation capture, per-linear; same fixed calibration tokens; MACaQ alpha=.75 ridge=1 flip_iters=8 exponent_offsets=1,0,-1; no held-out tuning",
	17	+ "accuracy": "ordinary W4A4 and MACaQ W4A4 use exact Ascend CPU reference activation quantization and BF16 butterfly rules; weight codes use published pure-Torch calibrator",
	18	+ "method_sha256": "f53f40ddd4006c43dde6a7fbc7f43b2f210ce2320f73562ef96bc6681401bc59",
	19	+ "reference_sha256": "f65f8d3229339fba1a1762edd0bbd9825e534690f15099185fc0eea1452b67f6",
	20	+ "torch": "2.9.0+cpu"
	21	+}

Adocs/benchmarks/macaq_20261027/accuracy.md+56-0

@@ -0,0 +1,56 @@
	1	+# Held-out MACAQ accuracy
	2	+
	3	+MACAQ lowers CPU-reference model perplexity by 23.7% versus ordinary W4A4 (33.4286 → 25.4962), and lowers actual NPU projection NMSE against both native scaling presets on all nine fixed projections. The model pilot and kernel latency sweep are separate experiments. The online MACAQ transform is fused with quantization and matmul; calibration and weight transformation are offline. Competitive latency is demonstrated for the reported small-batch workloads, not every shape or end-to-end model execution.
	4	+
	5	+These experiments compare calibrated MACAQ against ordinary MXFP4 W4A4. The independent NPU baseline quantizes both original BF16 weights and activations with native `npu_dynamic_mx_quant` (block32 E2M1, round-to-even) and calls `npu_quant_matmul`. Both supported FP4 scaling presets, `scale_alg=0` and `scale_alg=2`, are evaluated with their default type maximum. These presets were fixed before evaluation; neither was selected or tuned on the held-out results.
	6	+
	7	+## Actual NPU projection outputs
	8	+
	9	+Nine projections were fixed before evaluation: q, up and down projections at layers 0, 14 and 27 of Qwen/Qwen3-0.6B. Calibration used 1,024 WikiText-2 training tokens; evaluation used 1,024 separate test tokens. ops-nn, the TileLang standard API, the TileLang prepared callable and native alg0 matched their respective CPU algorithm references exactly for all 36 outputs. Nine additional native alg2 matmul outputs matched CPU decoding of their actual FP4 codes/scales within BF16 rtol=atol=1/128. MACAQ reduced NMSE against both native presets in all nine projections. ops-nn zero-padded unsupported K/N pairs to supported shapes and sliced the original output columns; these padded calls support an accuracy claim only. TileLang ran the original dimensions.
	10	+
	11	+![Projection accuracy](https://raw.gitcode.com/Mocchibird/ops-tilelang/blobs/d27dba41cd1bdeb69577276757cceb5dde610f42/docs/benchmarks/macaq_20261027/accuracy-npu-projections.png)
	12	+
	13	+\| Projection \| Native alg0 NMSE \| Native alg2 NMSE \| MACAQ NMSE, all interfaces \| Reduction vs alg0 \| Reduction vs alg2 \|
	14	+\|---\|---:\|---:\|---:\|---:\|---:\|
	15	+\| L0 down_proj \| 0.030517 \| 0.031154 \| 0.018509 \| 1.65x \| 1.68x \|
	16	+\| L0 up_proj \| 0.032722 \| 0.040139 \| 0.018919 \| 1.73x \| 2.12x \|
	17	+\| L0 q_proj \| 0.012962 \| 0.011487 \| 0.005795 \| 2.24x \| 1.98x \|
	18	+\| L14 down_proj \| 0.034110 \| 0.037999 \| 0.021642 \| 1.58x \| 1.76x \|
	19	+\| L14 up_proj \| 0.027275 \| 0.030853 \| 0.016969 \| 1.61x \| 1.82x \|
	20	+\| L14 q_proj \| 0.014667 \| 0.016256 \| 0.007577 \| 1.94x \| 2.15x \|
	21	+\| L27 down_proj \| 0.021753 \| 0.019562 \| 0.004722 \| 4.61x \| 4.14x \|
	22	+\| L27 up_proj \| 0.021788 \| 0.025911 \| 0.008336 \| 2.61x \| 3.11x \|
	23	+\| L27 q_proj \| 0.014878 \| 0.017160 \| 0.007907 \| 1.88x \| 2.17x \|
	24	+
	25	+## Grouped arithmetic on real projection fixtures
	26	+
	27	+The same preselected down projections at layers 0, 14 and 27 also exercise the grouped kernels. Each fixture replicates its calibrated weights and diagonal into four identical expert banks. This is arithmetic coverage on real held-out projection inputs, not an MoE-model quality evaluation or an independent accuracy dataset. M=1024 and K=3072 are unchanged; zero weight columns pad N from 1024 to 2048 for TileLang and to the supported 7168 for ops-nn. Outputs are sliced back to the original 1024 columns. These padded calls make no runtime claim.
	28	+
	29	+All 24 grouped outputs match their CPU algorithm references exactly: native CANN algorithms 0/2, MACAQ regular and prepared-16 calls with contiguous or alternating-16 routes, and MACAQ prepared-128 / ops-nn calls with valid contiguous 256-row expert runs. The adaptive TileLang kernel is explicitly confirmed as the selected prepared-16 implementation. All six MACAQ checks per projection have identical NMSE, lower than both native presets.
	30	+
	31	+\| Projection \| Native grouped alg0 NMSE \| Native grouped alg2 NMSE \| MACAQ grouped NMSE \|
	32	+\|---\|---:\|---:\|---:\|
	33	+\| Layer 0 down \| 0.030517 \| 0.031154 \| 0.018509 \|
	34	+\| Layer 14 down \| 0.034110 \| 0.037999 \| 0.021642 \|
	35	+\| Layer 27 down \| 0.021753 \| 0.019562 \| 0.004722 \|
	36	+
	37	+[Grouped raw outputs](accuracy-grouped-dynamic.json) · [Grouped source/protocol metadata](accuracy-grouped-dynamic.metadata.json).
	38	+
	39	+[Current dense projection data](accuracy-npu-dynamic-final.json) · [Current dense source/runtime metadata](accuracy-npu-dynamic-final.metadata.json).
	40	+
	41	+
	42	+## CPU model reference pilot
	43	+
	44	+All 196 decoder linear projections were quantized. Embeddings and the LM head remained BF16 in every arm. The test used the first 16,384 held-out WikiText-2 tokens, 1,024-token contexts, and 16,368 scored next-token predictions. Calibration used 1,024 training tokens. MACAQ hyperparameters were fixed: alpha=0.75, ridge=1, eight flip iterations and exponent offsets (1,0,-1). No hyperparameters were tuned on the held-out tokens. The ordinary quantizer in this model pilot uses scale algorithm 0; algorithm 2 model perplexity has not been measured.
	45	+
	46	+![CPU model perplexity](https://raw.gitcode.com/Mocchibird/ops-tilelang/blobs/04bc63d008c2676c2c50fdc9cf8a35a3aa4415c8/docs/benchmarks/macaq_20261027/accuracy-perplexity.png)
	47	+
	48	+\| Model reference \| Perplexity \|
	49	+\|---\|---:\|
	50	+\| BF16 \| 21.3111 \|
	51	+\| Ordinary MXFP4 W4A4 (alg0) \| 33.4286 \|
	52	+\| MACAQ W4A4 \| 25.4962 \|
	53	+
	54	+This is a one-model CPU reference pilot using the Ascend BF16 transform/quantization rules. It is not a full-dataset perplexity result, full-model NPU execution, or a model-throughput measurement. The nine NPU projection checks validate the arithmetic bridge; they do not turn this CPU pilot into an NPU perplexity run.
	55	+
	56	+Runtime and accuracy are separate experiments. No joint model-level speed-and-accuracy claim is made from these results.

Adocs/benchmarks/macaq_20261027/macaq-submitted-metadata.json+49-0

@@ -0,0 +1,49 @@
	1	+{
	2	+ "source_commits": {
	3	+ "ops-nn": "09ed381118aeab506fb7219daf85f3d20b6d50ef",
	4	+ "ops-tilelang": "b7d8b14c1e14f2412e4a37ccecde85e6284d0349"
	5	+ },
	6	+ "extension_wheel_sha256": "06058f56e8688cc6d4d5c89c4e146f4d9cd15e20872be399ae6d1298d23fb08a",
	7	+ "torch": "2.9.0+cpu",
	8	+ "torch_npu": "2.9.0.post2",
	9	+ "tilelang": "0.1.15",
	10	+ "baselines": [
	11	+ "plain",
	12	+ "staged"
	13	+ ],
	14	+ "samples": 5,
	15	+ "repeats": 30,
	16	+ "timing": "sum of msprof device kernel durations per invocation, excluding cache flush; separate synchronized host call timing",
	17	+ "preparation": "TileLang normalized diagonal and weight scale layout prepared once; ops-nn preparation inside extension call is included",
	18	+ "harness_sha256": "3bcff7ad8ec2a8b6c4d8458a1d55968351ad0733eae4a99b2bcbd087963b4fd9",
	19	+ "profiler_sha256": "e322b8e09d2e1e3565093bcd6e332316426a4a596375db131d04096a3feef7f9",
	20	+ "device_name": "Ascend950DT_9592",
	21	+ "device_assignment": {
	22	+ "cases_0_to_47": "card A",
	23	+ "cases_48_to_83": "card B",
	24	+ "cases_84_to_95": "card A"
	25	+ },
	26	+ "primary_comparison": "MACaQ fused public API vs independent native CANN dynamic MXFP4 quantization + dense/grouped matmul (ascendc_plain), which omits MACaQ transform",
	27	+ "primary_metric": "repeated synchronized elapsed eager public-API brackets; all outputs allocated inside call",
	28	+ "elapsed_sweep": {
	29	+ "harness_sha256": "25c7899224525617b1147305b93247c4dc5c2b039130423c3d44f0ae6fc42a9c",
	30	+ "imported_harness_sha256": "3bcff7ad8ec2a8b6c4d8458a1d55968351ad0733eae4a99b2bcbd087963b4fd9",
	31	+ "trials": 15,
	32	+ "launches_per_trial": 20,
	33	+ "warmups_per_arm": 20,
	34	+ "input_pool": "two distinct activation buffers with identical deterministic values; alternate buffers each call",
	35	+ "timing": "synchronized wall brackets of 20 eager public API calls; arm order rotates each trial; no cache flush",
	36	+ "allocation": "all three arms allocate a fresh output inside every timed call; intermediate allocations included",
	37	+ "preparation": "static packed weights, scale views and TileLang diagonal prepared once; ops-nn internal per-call diagonal preparation included"
	38	+ },
	39	+ "elapsed_device_assignment": {
	40	+ "cases_0_to_47": "card A",
	41	+ "cases_48_to_95": "card B"
	42	+ },
	43	+ "toolkit": "CANN 9.3.0 weekly.20261021.01 / V100R001C13B023",
	44	+ "profiled_sweep_output_checks": 384,
	45	+ "elapsed_sweep_output_checks": 576,
	46	+ "configurations": 96,
	47	+ "elapsed_evidence_sha256": "b2eb930a843c8f7b49ee895e0eff53ad4e05ee37c2f9e60950343e5a13a95b02",
	48	+ "evidence_sha256": "fa56ba61163624fd1c8244fe1981e2a99f429266a64237cb602c4c53f75d13c6"
	49	+}