Add one-level DynamicMxQuant kernel with precision check - #83
Add one-level DynamicMxQuant kernel with precision check#83ziyang-cheng wants to merge 3 commits into
Conversation
组件合入前的功能/精度看护基线，基于 upstream/main。第一批仅收 tail_cublas_fp8： - kernel: dynamic_mx_quant_tail_cublas_fp8.hpp + dynamic_mx_quant_common.hpp - test: tail_cublas_fp8.cpp(8×32 基础场景)+ Makefile(仅 TAIL_CUBLAS_FP8 块) - harness: gen / compare / run_precision_check(CONFIGS 裁剪为仅 TAIL_CUBLAS_FP8) 验证(gfrun 代 QEMU，工具链含 B.IOR 字节步长修复)：gen(seed=8)→编译 res_check=on →gfrun(R2=0)→compare：output=pass(MSE=0)、scale=pass(MSE=0)，逐字节精确。 compare/ 生成数据不入库。 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
… 及 compile_all 集成 承接第一批（510f44e）继续迁入已对齐 AscendC 的 driver，并把 dynamic_mx_quant 挂进顶层全量构建： - tail_ocp_fp4（kernel + driver，8×64）：编译通过出 ELF；gfrun 精度当前被 emulator fp4 TSTORE 合法性缺陷阻塞（IsLegalLocalTileDescriptor 用 BytesOf(FP4)=1B 而非 ElementBits/8，AccumulateBlockInfo.cpp:50/98）——不改 emulator/工具链。 - nontail_cublas_fp8（kernel plain + bigbs + driver，32×32）：编译通过出 ELF； gfrun 精度 FAIL（output MSE≈1586/MaxAE=441、scale MSE≈7421，e4m3 饱和征兆），待定位。 - compile_all 集成：one-level compile_all.sh 挂 compile_operator dynamic_mx_quant； 新建裁剪版 compile.all，仅枚举已迁移 3 driver（TAIL_CUBLAS_FP8/TAIL_OCP_FP4/ NONTAIL_CUBLAS_FP8）的 plain + res_check，刻意不含尚未迁移的 NONTAIL_OCP_FP4/FP4_PROBE。 - Makefile / run_precision_check.py 追加对应 TYPE 块与 CONFIG，逐字段对齐源仓正式格式。 产物（compare/*.bin、output/）不入库。 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 88f68b0410
ℹ️ About Codex in GitHub
Codex has been enabled to automatically review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
When you sign up for Codex through ChatGPT, Codex can also answer questions or update the PR, like "@codex address that feedback".
| global QEMU | ||
| QEMU = args.qemu |
There was a problem hiding this comment.
Running this precision-check script fails immediately with SyntaxError: name 'QEMU' is used prior to global declaration, because main() reads QEMU as the argument default on line 126 before declaring it global here. Move the declaration to the start of main() or avoid mutating the module global so none of the advertised precision checks are blocked at startup.
Useful? React with 👍 / 👎.
| env = os.environ.copy() | ||
| env["COMPILER_DIR"] = compiler_dir | ||
| run([ |
There was a problem hiding this comment.
On a clean shell without COMPILER_DIR already exported, --compiler-dir has no effect: this function constructs a modified env but never passes it to run()/subprocess.run(), so Makefile.common aborts with COMPILER_DIR is not set. Forward this environment or pass COMPILER_DIR=... in the make arguments so the documented invocation works.
Useful? React with 👍 / 👎.
| // The zip needs TINTERLEAVE/TDEINTERLEAVE, which LinxISA 0.57 defines but | ||
| // the -D__linx header does not expose (RECORD 问题5). Once exposed, insert | ||
| // a TINTERLEAVE of even/odd block-rows right here before the store. | ||
| TSTORE(gs, scale_u8); // store scale early; scale_byte now dead |
There was a problem hiding this comment.
For every non-tail invocation, this store writes planar [scaleRows, Post] bytes, while the operator's required mxScale layout is parity-interleaved [scaleRows/2, Post, 2] as the surrounding code and golden generator document. Consumers therefore read scales for the wrong columns/parities—even the registered one-block case places values contiguously rather than at alternating byte offsets—so this kernel cannot produce a compatible scale tensor until the store layout is interleaved.
Useful? React with 👍 / 👎.
| ], check=False) | ||
| return result.stdout |
There was a problem hiding this comment.
When output or scale comparison reports fail, size_mismatch, or missing files, the comparator only prints that status and exits successfully; this call also explicitly ignores its return code, so run_precision_check.py records stdout and ultimately exits 0. Any CI or experiment driver invoking this pipeline will therefore treat incorrect kernel results—including the known non-tail scale-layout mismatch—as passing; return and propagate a nonzero status whenever either comparison is not pass.
Useful? React with 👍 / 👎.
| // AscendC scale layout: uint8 E8M0, one byte per block, compact [M, scaleCols] | ||
| // with the block count even-aligned (scaleColNum_ = CeilDiv(numKb,2)*2). The | ||
| // trailing padding column is left zero. Mirrors dynamic_mx_quant_tail_axis_fp8.h:168. | ||
| constexpr int scaleCols = ((numKb + 1) / 2) * 2; |
There was a problem hiding this comment.
When K / BlockSize is odd, scaleCols allocates an extra padding column that must contain E8M0 byte 0x00, but this kernel only stores the numKb real columns and never initializes the extra one. The test driver's static zero-initialized buffer hides this; callers using an uninitialized or reused output buffer receive arbitrary padding that no longer matches the required mxScale tensor, so explicitly write the padding column as the FP4 implementation already does.
Useful? React with 👍 / 👎.
| diff = out_u8.astype(np.int32) - gold_u8.astype(np.int32) | ||
| mse = float(np.mean(diff.astype(np.float64)**2)) | ||
| max_ae = float(np.max(np.abs(diff))) | ||
| status = "pass" if mse < 0.1 else "fail" |
There was a problem hiding this comment.
Compact E8M0 scale bytes are expected to match exactly, but this MSE threshold accepts isolated encoded-exponent errors: for the 16-byte tail test, one scale byte differing by one yields MSE 0.0625 and is reported as pass, and the 64-byte non-tail test can hide several such errors. Use exact array equality (or require zero MaxAE) so sparse scale corruption cannot pass the precision baseline.
Useful? React with 👍 / 👎.
| dynamic_mx_quant_tail_cublas_fp8<8, 32, 32, __fp8_e4m3, __half>( | ||
| xh16, reinterpret_cast<__fp8_e4m3*>(y_h16), scale_h16); | ||
| dynamic_mx_quant_tail_cublas_fp8<8, 32, 32, __fp8_e4m3, float>( | ||
| xf32, reinterpret_cast<__fp8_e4m3*>(y_f32), scale_f32); |
There was a problem hiding this comment.
When this driver is built with RES_CHECK, these supposedly compile-only half/fp32 calls still execute after the bf16 kernel under test; tail_ocp_fp4.cpp has the same pattern, whereas the non-tail driver correctly guards such calls with #ifndef RES_CHECK. The comments state these paths are runtime-blocked by toolchain/emulator skew, so they can make QEMU fail before outputs are written and prevent otherwise valid bf16 precision checks; exclude them from result-check builds.
Useful? React with 👍 / 👎.
| parser.add_argument("--cmp-root", dest="cmp_root", | ||
| default=os.path.abspath(os.path.dirname(__file__) + "/../../../compare"), | ||
| help="compare root directory") |
There was a problem hiding this comment.
Invoking the comparator directly without --cmp-root resolves this default to benchmark/one-level-arch/test/kernel/compare, because the script lives under test/kernel/quant/dynamic_mx_quant/src; generated data actually lives under benchmark/one-level-arch/compare. The standalone command therefore reports every output as missing unless callers know to override the ostensibly optional argument, so derive the default from the one-level architecture root.
Useful? React with 👍 / 👎.
- probe_dynamic_mx_quant_tail_ocp_fp8_newcalc.hpp: fp16 in -> e4m3 out, OCP tail, BlockSize=32, 倒数走位补(非 TRECIP)的单线程 probe。 - probe_dynamic_mx_quant_tail_ocp_fp8_newcalc_mt.hpp: 按 M 切分的 SPMD 4-PE 变体, kernel 内靠 get_thread_idx() 自我切分, main 只调一次。 - probe_ocp_fp8_newcalc.cpp: 单驱动, -DMT 宏 dispatch 选 _mt(4-PE)/ 单线程版。 - Makefile: 新增 PROBE_OCP_FP8_NEWCALC TYPE 块。 - run_precision_check.py: 注册 PROBE_OCP_FP8_NEWCALC 配置(M8 K32 OCP FP8 fp16-in)。 验证(现有工具链, gfsim 08-25 build): 单线程 + MT 两版均编译通过(EXIT=0)。 MT ELF 271120B, ssrset=0(不触发 gfsim Decoder ssrset-null 潜伏 bug)。 gfrun -s softcore.multiThreadNum=4 -> R2=0, 4 线程各退休 414 块 / 总 1656。 gfsim --conf fourpe -> Total Cycles=20037, 与 PERF §4.3 记录一致。 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: ce012306ca
ℹ️ About Codex in GitHub
Codex has been enabled to automatically review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
When you sign up for Codex through ChatGPT, Codex can also answer questions or update the PR, like "@codex address that feedback".
| "--dtype", cfg["dtype"], | ||
| "--scale-layout", cfg.get("scale_layout", "broadcast"), | ||
| "-o", str(cmp_dir), |
There was a problem hiding this comment.
When PROBE_OCP_FP8_NEWCALC is selected, its config requests in_dtype="fp16", but this command omits --in-dtype, so the generator uses its bf16 default. The probe driver then reads those bf16 bit patterns as __half while the golden was computed from bf16 values, making the advertised probe precision check compare results for a different input tensor; pass cfg.get("in_dtype", "bf16") through here.
Useful? React with 👍 / 👎.
| if biased_exp <= 0: | ||
| # round-half-to-even (numpy.rint) to match ttk _mx_round_mantissa rint | ||
| mant = round(x / (2 ** -6)) |
There was a problem hiding this comment.
For scaled magnitudes in the E4M3 subnormal range, the mantissa must be rounded in units of 2**-9, not 2**-6. For example, this function encodes 2**-8 as 0x00 instead of 0x02 and encodes 7 * 2**-9 as 0x01 instead of 0x07; inputs with sufficient within-block dynamic range therefore produce an incorrect golden tensor and can report a correct kernel as failing.
Useful? React with 👍 / 👎.
| make TESTCASE=dynamic_mx_quant TYPE=TAIL_CUBLAS_FP8 diss | ||
| make TESTCASE=dynamic_mx_quant TYPE=TAIL_OCP_FP4 diss | ||
| make TESTCASE=dynamic_mx_quant TYPE=NONTAIL_CUBLAS_FP8 diss |
There was a problem hiding this comment.
Because the script neither enables set -e nor combines the make statuses, its exit status is only that of the final non-tail res-check build. If any of the first five variants fails while the last succeeds, compile_all.sh reports dynamic_mx_quant compilation completed even though required ELFs were not built; make each failure terminate the script or accumulate and return a nonzero status.
Useful? React with 👍 / 👎.
| TROWMAX(max_h, abs_h); | ||
| TCVT(max_bf, max_h); // half -> bf16 |
There was a problem hiding this comment.
For the advertised InT=__half specialization, converting the reduced maximum to bf16 before clearing its mantissa can round a value just below a power of two upward; for example, half 1.9990234375 becomes bf16 2.0, raising the shared-scale exponent by one and quantizing the entire block with half the intended reciprocal. The fp32 branch and the duplicated tail-row branches have the same issue; extract/floor the exponent in the original or fp32 bit domain as the newcalc probe already does.
Useful? React with 👍 / 👎.
| parser.add_argument("--algo", type=str, default="OCP", | ||
| choices=["OCP", "CUBLAS", "DYNAMIC_RANGE"]) | ||
| parser.add_argument("--kernel", type=str, default="tail", | ||
| choices=["tail", "nontail"]) | ||
| parser.add_argument("--dtype", type=str, default="FP8", | ||
| choices=["FP8", "FP4"]) |
There was a problem hiding this comment.
The module documents that cuBLAS supports only FP8 and dynamic-range supports only FP4, but these independent choices accept --algo CUBLAS --dtype FP4 and --algo DYNAMIC_RANGE --dtype FP8 and generate apparently valid files. In particular, the cuBLAS implementation always uses the FP8 destination maximum even when FP4 was requested, so such a command silently creates a nonsensical golden baseline rather than reporting an invalid operator configuration.
Useful? React with 👍 / 👎.
|
|
||
| clean: | ||
| @find $(OBJ_ROOT) -type f -name "*.o" -exec rm -rf {} \; |
There was a problem hiding this comment.
This clean recipe runs as a prerequisite of every target through Makefile.common, but it deletes every .o anywhere under the shared one-level output tree. Consequently, when the newly added operator runs last in compile_all.sh, each of its six builds removes the object files produced for all previously compiled operators, defeating incremental builds and potentially interfering with another concurrent build; scope the deletion to this case's object directory instead.
Useful? React with 👍 / 👎.
| diff = out_f32 - gold_f32 | ||
| mse = float(np.mean(diff**2)) | ||
| max_ae = float(np.max(np.abs(diff))) | ||
| status = "pass" if mse < 0.1 else "fail" |
There was a problem hiding this comment.
When both files contain the same valid E4M3 NaN encoding (0x7f or 0xff), decoding produces paired NaNs, so subtraction and mean yield NaN and the status is always fail even though the bytes agree. This prevents correct special-value kernel results from passing; exclude lanes where both operands are NaN from the numeric metric or compare their encodings explicitly while still rejecting a one-sided NaN.
Useful? React with 👍 / 👎.
1 participant
No description provided.