Feat/dmxq rel0812 - #49
Conversation
- 3 scale algorithms: OCP, cuBLAS, DynamicRange - 2 axis modes: tail (2D loop) and non-tail (3D loop) - 6 ELF configurations (3 alg × 2 axis) - Multi-file architecture: common.hpp + tail.hpp + nontail.hpp - Register in compile_all.sh
…rnel - 4 个已调试 kernel（计算逻辑对齐 AscendC，已知问题记录在 RECORD.md）： TAIL_CUBLAS_FP8、NONTAIL_CUBLAS_FP8、TAIL_OCP_FP4、NONTAIL_OCP_FP4 - 4 个未调试 kernel 模板加 UNDEBUGGED 标注： tail/nontail OCP-FP8、tail/nontail DynRange-FP4 - emax 改为按输出 dtype 派生（constexpr trait），cuBLAS 补齐 exp/man/NaN/zero 守卫 - 删除未调试 kernel 的 test 入口及 bs/tm4 遗留探针 driver - README 引入已调试/未调试状态定义；RECORD 新增问题6（fp32→fp4 直转待 ISA 确认） Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
- 新增 nontail_{ocp_fp4,cublas_fp8}_bigbs 两个模板（方案 A：切分归约轴 +
running-TMAX 累积），规避非尾轴 TileN 上对齐下界与 TileSize 上界的双重约束，
大 BlockSize 下任意 TileN 均可满足对齐；逐 op 对齐 AscendC ComputeScale{Ocp,CuBlas}
- common.hpp 与 4 个已 review 的 plain kernel 同步（scale boxed compact 平铺）
- 三文档定位对齐：DESIGN 保留从 0 设计（删除实现期约束 §7.5）、README 记录当前现状
（含 bigbs 变体）、RECORD 记录约束与工具链缺口
- Makefile 注册两个 bigbs TYPE 分支
Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
- tail_ocp_fp4 的 fp4 输出 tile 切分有两种等价但互斥的方案：当前默认的 列装箱补齐物理宽（PW），与旧的 2-block scratch-HBM concat 配对 + 零块方案 - 后者保留在 dynamic_mx_quant_tail_ocp_fp4.hpp.bak，选哪个待 toolchain<->emulator skew 解除、两者可运行期比对后再决策 - 在 kernel 头注释与 README 状态表中说明两方案并存与决策条件 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
TileM/TileN/R_sub 从调用方模板参数改为编译期从算子输入 + InT 预算推导 (common.hpp 新增 max_tilem/pick_tilen/max_rsub 等 constexpr helper)。 非尾轴统一入口 if constexpr 在无合法 TileN 时自动路由到 _bigbs 方案A, 删除独立 bigbs driver/Makefile TYPE,改在 nontail driver 加 BS=128 编译期 调用覆盖该路由。InT 仅作预算感知,数据路径仍 bf16(static_assert 守住)。 尾轴 ocp-fp4 的 Contig 传 PW(补齐物理宽)而非 BlockSize。默认 BS=32 推导值 与改造前一致,零行为回归;4 个保留 TYPE 编译+链接+反汇编通过。 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
fp16(__half)/bf16(__bf16)/fp32(float) 三种输入经 if constexpr 分派，镜像
AscendC Compute() 的 ComputeMaxExp{Ocp,Cublas}{Bf16,Half,Fp32}：类型差异集中
在输入正则化一处（OCP→uint16 bf16 指数域，cuBLAS→fp32 amax），scale-finalize
与 ComputeData 共享。6 个 kernel（4 debugged + 2 bigbs）放开 InT static_assert
并透传 InT；bf16 走原指针 reinterpret 路径，零回归。golden 生成器加 --in-dtype。
新增 align32_probe / dtype_probe 探针与 ALIGN32_PROBE/DTYPE_PROBE Makefile TYPE；
ISSUE_32B_align.md 记录 pto_tile.hpp:649 的 RowMajor/NoneBox 连续轴 32B 对齐
static_assert 对打包 fp4（单 MX block=16B）的约束、最小复现与报错。
runtime 待工具链↔仿真器 skew 解除后验证，当前仅编译 + 反汇编 + 逐 op 复核。
Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
镜像 AscendC ocp_new：以 max_exp(bf16=2^E_max) × 2^-emax 得共享 scale 2^(E_max-emax)，再 Cast<bf16->e8m0> 直接产出 E8M0 scale 字节，替换旧的 clamp/移位/inf-Select 路径。3 个 OCP kernel（tail、nontail plain、bigbs） scale 直接存 e8m0，去掉窄化 TCVT。recip 保持 uint16 域的 inf/zero/special 三选，eq_zero/eq_inf 取自原始 max_exp。边界1（inf/nan 依赖 cast 得 0xFF）、 边界2（微小非零下溢 recip 差异）记于 RECORD 问题7，待 skew 解除跑 golden。 编译验证：3 config 全部 res_check=on diss 干净；反汇编含 TMULS + TCVT→e8m0， 无 U16 移位（旧 scale 路径已消失）。 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
对 dynamic_mx_quant 全量构建 + gfrun 后，剔除已知问题清单，定位 4 处报错根因 均在 emulator（TABS/BF16、TROWMAX/U16 白名单）或 LinxV5 后端（%Z unknown operand、-O0 spill/reload 类不对称），非 kernel 逻辑。 - RECORD.md 新增问题9-12：单问题成节，交叉引用体现 fp16 规避→逐道撞墙→ 收敛到已知 skew 的验证链条 - ISSUE_rel0812_defects.md：对外提交版，独立缺陷、去链条、带 commit 锚定 - 新增 fp16 cuBLAS driver + Makefile TYPE 作问题9 复现入口 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
问题8 是独立可复现的 emulator↔LinxV5 交互缺陷（非内联 tile-helper 的 tile 参数 经 TSTORE/TLOAD S64 通用块栈传参，与 emulator 按源 tile dtype 精确匹配的口径冲突）， 默认 bf16 路径被缺陷1 掩盖、经 fp16 越过后首个命中。CmpMode 缺陷已单独提交，不收录。 Co-Authored-By: Claude Opus 4 <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: bd33df1e4b
ℹ️ 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".
| env = os.environ.copy() | ||
| env["COMPILER_DIR"] = compiler_dir | ||
| run([ |
There was a problem hiding this comment.
When --compiler-dir is supplied but COMPILER_DIR is not already exported, this still invokes make with the original process environment: the local env is constructed and then dropped before calling run(). In that scenario Makefile.common fails on the missing COMPILER_DIR, so the required CLI option does not actually let the precision pipeline compile; pass this environment into subprocess.run (or through run()).
Useful? React with 👍 / 👎.
| parser.add_argument("--qemu", default=QEMU, help="QEMU path") | ||
| args = parser.parse_args() | ||
|
|
||
| global QEMU |
There was a problem hiding this comment.
As written, this file does not parse: parser.add_argument("--qemu", default=QEMU, ...) reads QEMU earlier in main(), then this later global QEMU declaration makes Python raise SyntaxError: name 'QEMU' is used prior to global declaration before any precision check can run. Move the declaration to the top of main() or avoid assigning the module global so the script starts.
Useful? React with 👍 / 👎.
| constexpr int TileN = pick_tilen<BlockSize, Post, OutT, InT, /*IsCublas=*/false>(); | ||
| if constexpr (TileN >= nontail_align_lower<OutT>()) { | ||
| nontail_ocp_fp4_plain<Axis, Post, BlockSize, TileN, OutT, InT>(x, y, scale); |
There was a problem hiding this comment.
For legal Post sizes that are not an exact multiple of the derived aligned tile, such as dynamic_mx_quant_nontail_ocp_fp4<32, 96> with the default BlockSize, pick_tilen returns 128 and this dispatch calls nontail_ocp_fp4_plain, whose Post % TileN static_assert then fires. The public entry only rejects an invalid reduction axis, so this drops many Post multiples of 32 unless a boxed tail-column path is added or TileN is capped differently.
Useful? React with 👍 / 👎.
| constexpr int BigTileN = nontail_align_lower<OutT>(); | ||
| constexpr int Rsub = max_rsub<BlockSize, BigTileN, InT, /*IsCublas=*/true>(); | ||
| dynamic_mx_quant_nontail_cublas_fp8_bigbs<Axis, Post, BlockSize, BigTileN, Rsub, OutT, | ||
| InT, MaxLowBoundBits>(x, y, scale); |
There was a problem hiding this comment.
When the public non-tail cuBLAS entry auto-routes large BlockSize cases here, Post values that are not multiples of 32 (for example Axis=128, Post=48, BlockSize=128) hit the big-BS implementation's Post % TileN == 0 static_assert even though the plain cuBLAS path has an N_tail boxed path and does not require this. This silently narrows the public API for large BlockSize shapes; either add the same tail-column path to bigbs or route such shapes differently.
Useful? React with 👍 / 👎.
| static uint16_t scale[8 * 512] __attribute__((aligned(4096))) = {}; | ||
|
|
||
| int main() { | ||
| dynamic_mx_quant_tail_cublas_fp8<8, 512, 8, 512>(x, reinterpret_cast<__fp8_e4m3*>(y), scale); |
There was a problem hiding this comment.
The exposed BS512_CUBLAS tooling target still calls dynamic_mx_quant_tail_cublas_fp8 as if the fourth template argument were a tile/block-size integer, but the new signature makes that slot OutT; the call also passes a uint16_t * scale buffer to a uint8_t * parameter. Building TYPE=BS512_CUBLAS therefore fails instead of probing the 512-wide tail case, so the call and scale buffer need to be updated to the new public signature.
Useful? React with 👍 / 👎.
- 复现仓拆为编译仓(SuperNPUBench)/执行仓(SuperScalarModel)，明确复现跨两仓 - 补全工具链三 commit(gfrun 319294f / llvm eb64de8 / TileOP-API 72f8255)，标注与 README release_ver0812 验证版本一致 - 运行时缺陷(1/2/5)复现命令补 gfrun 执行步骤：diss 只编译+反汇编、不跑 gfrun - 缺陷4 复现命令变量名 EXTRA_CXXFLAGS→CFLAGS(Makefile 无 EXTRA_CXXFLAGS，原命令被忽略仍以 -O2 编译) - 新增「复现方式」小节，区分编译期缺陷(3/4)与运行时缺陷(1/2/5) Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: a9e27e7d88
ℹ️ 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".
| for elf in elf_paths: | ||
| r = check_elf(elf, args.cmp_root, args.dtype, args.scale_layout) | ||
| results.append(r) | ||
| print(f"{r['elf']}: output={r['output']}, scale={r['scale']}") |
There was a problem hiding this comment.
If any output or scale comparison returns fail, size_mismatch, or missing files, this loop still only prints the result and the script falls through with exit code 0. run_precision_check.py invokes this comparator as a subprocess, so precision regressions are not machine-detectable; aggregate the statuses and sys.exit(1) when any checked result is not a 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 includes an extra padding column, but this kernel only stores the real kb columns and never writes column numKb (unlike the tail OCP FP4 path). Callers that pass an uninitialized or reused scale buffer will leave stale bytes where the AscendC/golden compact layout expects 0x00 padding.
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.
For compact scale output this still accepts mse < 0.1, so small discrete byte mismatches are reported as pass (for example, one off-by-one scale byte in any registered compact buffer longer than 10 bytes). The comment says this layout is a direct byte compare, so use exact equality or at least require max_ae == 0 before reporting pass.
Useful? React with 👍 / 👎.
| dynamic_mx_quant_tail_cublas_fp8<8, 32, 32, __fp8_e4m3, __half>( | ||
| xh16, reinterpret_cast<__fp8_e4m3*>(y_h16), scale_h16); |
There was a problem hiding this comment.
These calls are described as compile-only, but they are ordinary calls inside main() and therefore run under the QEMU precision pipeline before the RES_CHECK outputs are written. On the registered BF16 config this also executes the fp16/fp32 branches that the surrounding docs call runtime-skew/diagnostic-only, so a side-branch emulator failure can abort an otherwise valid precision run; gate these calls away from RES_CHECK or replace them with non-executed explicit instantiations.
Useful? React with 👍 / 👎.
| inline void compute_cublas_core( | ||
| pto::Tile<pto::Location::Vec, float, R, C, pto::BLayout::RowMajor, ValidR, ValidC> &max_abs, | ||
| pto::Tile<pto::Location::Vec, uint16_t, R, C, pto::BLayout::RowMajor, ValidR, ValidC> &scale_byte, | ||
| pto::Tile<pto::Location::Vec, uint16_t, R, C, pto::BLayout::RowMajor, ValidR, ValidC> &recip_out) { |
There was a problem hiding this comment.
compute_cublas_core takes Tile references but is only marked inline; when the Linx backend leaves this helper as a call, tile arguments are spilled/reloaded through S64 local TSTORE/TLOAD and the documented emulator path rejects those synthesized stores before cuBLAS fp16/fp32 runtime checks can complete. Mark this helper and its cuBLAS scale wrappers always_inline (or otherwise eliminate the call boundary) so tile arguments are not passed through stack spills.
Useful? React with 👍 / 👎.
- 新增 dynamic_hi_f4_quant 算子骨架（tail）+ 探针（e6m2/segreduce/framework_emit） - mx_quant DESIGN/README：OCP scale 文档改为 bf16 乘 2^-emax + Cast<e8m0> 直转 - bs512_cublas 测试改用紧凑 uint8 E8M0 scale（[M, scaleCols]）+ 模板参数精简 - ISSUE_rel0812 拆分为 isa_dtype_support / toolchain_defects 两份 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
新增 probe_dynamic_mx_quant_tail_ocp_fp8.hpp（half→e4m3，BlockSize=32， OCP bf16→e8m0 直转，尾轴，half 值域归约规避 U16-TROWMAX）+ driver + Makefile TYPE=PROBE_OCP_FP8。kernel 为未加规避的复现版：gfrun 直接挂 M47 ValidateLocalTlsu 断言（TCVT_T 唯一漏发 lb2 → boxed valid-col-1 的 col 塌成 1 vs TSTORE 声明 Cols=32）。 RECORD 问题13 + ISSUE_tcvt_no_lb2.md 记录根因、各仓复现分支、两条已实测 解法（A kernel 侧盖章 TMULS/TMAX；B toolchain 侧 TCVT_T 补 B.DIM lb2）， 均收敛到已知 startup skew 墙，非 kernel 缺陷。缺陷根仓为 Linx-TileOP-API 组件源，本次不改 toolchain 以保留可复现。 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
指向含 probe kernel 的复现提交 1bd945a（该 commit 即可原样复现 M47 断言）。 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 67c3c5cfba
ℹ️ 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".
| // 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 cuBLAS invocation, this writes a planar [scaleRows, Post] buffer although the public mxScale layout is parity-interleaved. This already corrupts the registered NONTAIL_CUBLAS_FP8 case with one real and one padding block row: the kernel emits all real scales followed by zeros, while compact_scale_bytes() expects each scale immediately followed by its padding byte. Consumers therefore associate scale bytes with the wrong columns; perform the even/odd zip before the store rather than exposing the planar workspace as the result.
Useful? React with 👍 / 👎.
| except Exception as e: | ||
| print(f" ERROR: {e}", file=sys.stderr) | ||
| results.append((type_name, f"ERROR: {e}")) |
There was a problem hiding this comment.
When compilation, ELF lookup, or QEMU execution fails for any selected configuration, this handler records the error but lets main() finish normally, so the precision pipeline exits with status 0 despite not completing the experiment. Track these exceptions and return a nonzero status after printing the summary so automation can detect infrastructure and execution failures.
Useful? React with 👍 / 👎.
将 dynamic_mx_quant OCP 尾轴探针的指数位提取链从 scratch-HBM 位重解释往返 （TSTORE bf16 -> TLOAD u16 -> TANDS -> TSTORE u16 -> TLOAD bf16）改为 v0.58 零指令寄存器 bitcast reinterpret_tile<uint16_t>(max_bf) + 就地 TANDS + 引用别名 读回 bf16 域。旧 scratch-HBM 写法以注释保留便于回退。 此版本是问题14（reinterpret 零指令 bitcast 运行期不可见，emulator ValidateScalarLogicalTepl dtype 相等断言误杀）与问题15（TCVT bf16->e8m0 未建模） 的复现探针。scratch-HBM 版（历史 commit 1bd945a）对应问题13（TCVT 不发 lb2）。 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
- 新增 ISSUE_e8m0_cvt.md（问题15）：emulator 未实现浮点→e8m0(SF8) 正向转换， CubeEngine SF8 落 assert 分支、FloatPointUtils 无 SF8 目标条目 - 新增 ISSUE_reinterpret_dtype_tag.md（问题14）：ValidateScalarLogicalTepl dtype 相等断言误杀零指令 reinterpret_tile 位重解释 - 复现基线 Linx-TileOP-API 指向 personal 仓 cb47f6d（含 lb2 补丁） - README/RECORD：问题3（原生 CmpMode TCMP/TCMPS）、问题4（reinterpret_tile 寄存器 bitcast）主线已闭合，业务 kernel 待迁移 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
- probe_dynamic_mx_quant_tail_ocp_fp8_newcalc.hpp：接口/入参/输出与原探针一致， 仅把倒数求法从 TCVT+TRECIP 换成 bf16 指数字段位补（TXORS 0xFFFF + TSUBS 0x80FF = 0x7F00-bits），对 2 的幂精确，与原探针逐字节一致。 - 新增 probe_ocp_fp8_newcalc.cpp 驱动 + Makefile TYPE=PROBE_OCP_FP8_NEWCALC。 - 原探针 probe_ocp_fp8.cpp 改用 0x4400(4.0) 固定值填充，静态初始化避免运行期噪声。 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
- ISSUE_e4m3_overflow_inf.md：fp32→e4m3 TCVT 用 IEEE-with-inf 语义， float8_params exp_max=15 保留给 inf/nan，≥2^8 溢出成 +inf(0x78)；OCP E4M3 无 inf、max=448。含源码定位（softfloat float8_params + parts_uncanon 溢出分支）、 arm_althp 修复模板、非均匀块决定性证据。复现基线指向 SuperNPUBench c3401fd。 - ISSUE_e8m0_cvt.md / ISSUE_reinterpret_dtype_tag.md：复现基线表补 fork 仓地址。 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 71203cd386
ℹ️ 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".
| // Reg::Interleave. Blocked on TINTERLEAVE/TDEINTERLEAVE not being | ||
| // exposed in the -D__linx header (RECORD 问题5); insert the even/odd | ||
| // zip here once available. | ||
| TSTORE(gs, scale_e8m0); |
There was a problem hiding this comment.
When Axis / BlockSize is odd, scaleRows includes a padding row, but this loop writes only the real kb < numKb rows and never stores the required 0x00 E8M0 bytes to the padding row. The registered driver hides this because its static scale buffer is zero-initialized, while a caller supplying a reused or uninitialized buffer receives stale padding bytes; explicitly clear the padded row as the tail OCP path does.
Useful? React with 👍 / 👎.
|
|
||
| using gm_x = global_tensor<InT, RowMajor<M, N>>; | ||
| using gm_xu = global_tensor<uint16_t, RowMajor<M, N>>; | ||
| using gm_y = global_tensor<uint8_t, RowMajor<M, N / 2>>; |
There was a problem hiding this comment.
Because OutT is public and emax_bits() also accepts FP8 types, an instantiation such as dynamic_mx_quant_tail_ocp_fp4<8, 64, 32, __fp8_e4m3> is accepted, but this hard-coded N / 2 layout and the corresponding PW / 2 output tile write only half of an FP8 result. Add an OutT assertion for the two packed FP4 types, or derive the output dimensions from the element packing.
Useful? React with 👍 / 👎.
| return sign << 7 | ||
| exp = math.floor(math.log2(x)) | ||
| biased_exp = exp + 7 | ||
| if biased_exp <= 0: | ||
| # round-half-to-even (numpy.rint) to match ttk _mx_round_mantissa rint |
There was a problem hiding this comment.
For normalized values in the E4M3 subnormal range, this divides by 2^-6, although the decoder below defines each subnormal mantissa step as (1/8) * 2^-6 = 2^-9. Consequently even exactly representable values such as 2^-9, 2^-8, and 2^-7 are encoded as zero, and the preceding cutoff also flushes values that should round to the smallest subnormal. This makes golden.bin incorrect whenever quantization produces small nonzero FP8 outputs; compute the mantissa in units of 2^-9 and use the half-step 2^-10 as the zero cutoff.
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.
When this comparator is invoked directly without --cmp-root, the default resolves from src through only three parents to benchmark/one-level-arch/test/kernel/compare, while both Makefile.common's CHK_DIR and the generator place files under benchmark/one-level-arch/compare. Thus the advertised standalone -d mode reports missing files even after a successful run; derive the default from the one-level-arch root as the pipeline does.
Useful? React with 👍 / 👎.
- dynamic_mx_quant_tail_ocp_fp4: InT if constexpr 分派（bf16/half/fp32）， scale 归约与 data 路径按输入 dtype 走对应计算域，镜像 AscendC ComputeMaxExp* - 新增 ISSUE_tcvt_fp4_shape_contract：emulator ValidateOperandContract 对 「宽类型→打包 fp4」TCVT 逐 conjunct 比 physical row/col 结构性必崩（实证崩在 row 8≠4）；附加宽 tile 骗过断言但输出数据错的反证，论证修复须在 emulator 侧 - RECORD 问题16 记录同一缺陷；DESIGN/README 指针引用 - 删除 ENV_reproducible_setup.md（不再需要） Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 2796655844
ℹ️ 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".
| auto sh_i16 = reinterpret_tile<int16_t>(shared_bf); | ||
| TXORS(sh_i16, sh_i16, static_cast<uint16_t>(0xFFFF)); | ||
| TSUBS(sh_i16, sh_i16, static_cast<uint16_t>(0x80FF)); |
There was a problem hiding this comment.
When a block's extracted maximum exponent is zero or non-finite, this raw bit-complement is not the required reciprocal finalization: a BF16-subnormal block gets 0x7f00 instead of reciprocal zero, so its tiny inputs are amplified into nonzero FP4 values, while Inf/NaN also receives the wrong reciprocal. The AscendC path additionally selects zero, NaN, and the special exponent after subtraction; apply those selections here and in the mirrored tail-row path.
Useful? React with 👍 / 👎.
| void dynamic_mx_quant_nontail_ocp_fp4(InT *x, OutT *y, uint8_t *scale) { | ||
| static_assert(std::is_same_v<InT, __bf16> || std::is_same_v<InT, __half> || | ||
| std::is_same_v<InT, float>, | ||
| "InT must be one of {__bf16, __half, float}"); |
There was a problem hiding this comment.
When callers instantiate this public FP4 API with an FP8 OutT, such as dynamic_mx_quant_nontail_ocp_fp4<32, 64, 32, __fp8_e4m3>, emax_bits() accepts it but the implementation still hard-codes Post / 2 storage and a TileN / 2 output tile. The result therefore allocates and writes only half of the unpacked FP8 output; assert that OutT is one of the packed FP4 types before dispatching.
Useful? React with 👍 / 👎.
| } else if constexpr (std::is_same_v<InT, float>) { | ||
| tile_x abs_f; | ||
| TABS(abs_f, xin); | ||
| tile_maxf max_f; | ||
| TROWMAX(max_f, abs_f); | ||
| TCVT(max_bf, max_f); // fp32 -> bf16 |
There was a problem hiding this comment.
For the supported InT=float path, converting the reduced maximum to BF16 before masking can round a value immediately below a power of two up to that power—for example, the largest FP32 below 2.0 may become BF16 2.0. AscendC's FP32 OCP path masks the FP32 exponent and shifts it into the BF16 exponent field without rounding, so this implementation can choose a scale twice as large and quantize the block differently; extract the reduced FP32 exponent bits directly here and in the tail-row copy.
Useful? React with 👍 / 👎.
| # confirmation, RECORD 问题6). fp4 emit itself is verified (RECORD 问题2). | ||
|
|
||
| # --- plain builds --- | ||
| make TESTCASE=dynamic_mx_quant TYPE=TAIL_CUBLAS_FP8 diss |
There was a problem hiding this comment.
If any production or res_check build fails but the final FP4_PROBE build succeeds, Bash returns the final command's zero status and compile_all.sh reports dynamic_mx_quant compilation completed. This hides failures in the eight configurations the script is intended to validate; retain all builds if desired, but aggregate their statuses and exit nonzero when any make fails.
Useful? React with 👍 / 👎.
| clean: | ||
| @find $(OBJ_ROOT) -type f -name "*.o" -exec rm -rf {} \; |
There was a problem hiding this comment.
Every build of this target runs clean as a prerequisite through Makefile.common, and this override deletes every .o beneath the shared one-level output root rather than only this operator's objects. A dynamic-mx-quant build therefore invalidates all previously compiled operators and can race with a concurrent build by removing an object just before it links; scope the deletion to $(OBJ) and this target's common artifacts.
Useful? React with 👍 / 👎.
- plain/bigbs 就地展开 compute_cublas_core，换 reinterpret_tile + 原生 TCMPS<CmpMode>，替代 scratch-HBM 往返 + min/max-EQ 模拟 - 新增独立 bigbs harness（TYPE=NONTAIL_CUBLAS_FP8_BIGBS）+ BS 参数化 golden； BS=128（Axis=128/Post=32→R_sub=32/TileN=32）gfrun R2=0，data 逐字节匹配、 scale 值逐字节匹配（仅问题5 parity 交织布局差） - RECORD 问题5 补「布局无关」验证小节：data 逐字节（recip 取自寄存器、不经落盘 布局）+ 单调判别实验（planar 字节精确、无列错位）坐实「仅布局差」；问题17/README 加指针 - bigbs static_assert 收紧到 R_sub*TileN<=2048（32b 中间量固有，实测更正） - 新增 ROW32_PROBE 探针（TYPE=ROW32_PROBE）作调试工具 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
同一 TileLogicalShapeMatch 契约（TCVT src/dst physical Rows/Cols 全等）在当前 工具链落两层：编译期 template_asm.hpp:115 static_assert（崩 Cols 32≠64，打包 fp4 tile_o=PW/2 vs 源 PW）+ 放宽后 emulator Block.cpp:1039 运行期（崩 Row）。 新增 fp4_shape_probe.cpp 单条 TCVT 探针（TYPE=FP4_SHAPE_PROBE，WIDEN 切换 OCOL=32/64）闭环实证：变体A 不一致→编译崩、变体B 加宽→数据错。修复须 工具链头 + emulator 双侧删 physical row==row/col==col 只留 valid。 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
full loop 与 M_tail loop 的 scale 倒数 finalize 从简单位取补 （TXORS 0xFFFF + TSUBS 0x80FF）替换为就地展开 common::finalize_recip_u16 （规避问题8）：reinterpret_tile<uint16_t> 视图上 TCMPS 判 eq_inf/eq_zero/ eq_special，TSUB(0x7f00-shared) 得基础倒数后依次 TSEL 覆盖 inf→NaN、 zero→0、special→0x0040，对齐 AscendC 的特殊值语义。 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
内联展开 compute_cublas_scale_tail + compute_cublas_core（规避问题8 tile 入参 S64 栈往返），reinterpret_tile 零指令视图替 scratch-HBM（问题4）， 原生 TCMPS<CmpMode> 替 min/max 模拟（问题3）。fp16 driver gfrun R2=0 逐字节对齐 golden（含宽数据 |max| 0.24→25.1、scale 117→123 复验）。 至此 3 个 cuBLAS kernel 全部迁移，compute_cublas_core 已无 kernel 调用。 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
值域归约（TABS+TCOLMAX 求块 |max| → TCVT bf16 → reinterpret_tile+TANDS 取指数位 → TMULS 乘 2^-emax → TCVT e8m0）替换旧指数位域归约，1:1 对齐 tail_ocp_fp4 母本。 finalize_recip_u16 全内联（inf/zero/special 三 TSEL），消除函数调用（问题8）与 scratch-HBM reinterpret_u16_to_bf16（问题4）；bigbs 用 fp32 统一累加器跨 R_sub 子块 running-TMAX，删 static helper ocp_scale_from_maxexp_not_tail_boxed_bigbs。 README 同步更新迁移状态。 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
- 新增 tsel_inplace_probe.cpp（单条就地 TSEL）+ Makefile TYPE=TSEL_INPLACE_PROBE - ISSUE 以官方 gfrun@63dbb5a2 + TileOP-API cdeb624a(0.58.1 可编) 为复现基线， 实测崩 ValidateCompareSelectTepl AccumulateBlockInfo.cpp:383 - RECORD/README 记本地 ab822e7a/1f398190 修复及官方未采纳 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: f98f439aef
ℹ️ 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".
| # --- e6m2 (L1 base) cast emit probe --- | ||
| # Stage 1: honest one-level path — EXPECTED TO FAIL COMPILE (type_traits<__fp8_e6m2> | ||
| # unregistered; __type_code has no e6m2 entry). See RECORD 问题3. | ||
| make TESTCASE=dynamic_hi_f4_quant TYPE=E6M2_PROBE diss |
There was a problem hiding this comment.
When the framework build fails, or when E6M2_PROBE unexpectedly starts compiling, this script still returns the status of the final E6M2_PROBE_S2 command and can report success. The Stage 1 failure is intentional, but it should be asserted as an expected failure while unexpected failures from the framework and Stage 2 are propagated, otherwise this compile.all cannot reliably validate any of its three probes.
Useful? React with 👍 / 👎.
| dynamic_mx_quant_nontail_ocp_fp4<128, 64, 128>( | ||
| x_bs128, reinterpret_cast<__fp4_e2m1x2*>(y_bs128), scale_bs128); |
There was a problem hiding this comment.
When this driver is built with RES_CHECK, the large-BlockSize call here and the five following fp16/fp32 calls are still executed by QEMU even though they are described as compile-only. Any emulator limitation in those side routes can terminate the process before the primary kernel's result files are written; guard this block with #ifndef RES_CHECK, as the non-tail cuBLAS driver already does.
Useful? React with 👍 / 👎.
|
|
||
| def scale_recip_cublas(group_vals, emax: int) -> tuple: | ||
| # cuBLAS is FP8-only; emax unused (dstMax folded into inv_dst_max). | ||
| max_abs = max(abs(v) for v in group_vals) |
There was a problem hiding this comment.
When a cuBLAS reduction block contains a NaN after a finite value, Python's max commonly retains the finite value because comparisons with NaN are false (for example, max(abs(v) for v in [1.0, nan]) is 1.0). The golden then emits a finite scale instead of the 0xff non-finite scale this function's later guard intends, with the result depending on the NaN's position; detect any NaN explicitly before reducing the finite magnitudes.
Useful? React with 👍 / 👎.
| return "size_mismatch", float("nan"), float("nan") | ||
| out_f32 = np.array([fp8_e4m3_to_f32(int(b)) for b in out_bytes], dtype=np.float32) | ||
| gold_f32 = np.array([fp8_e4m3_to_f32(int(b)) for b in gold_bytes], dtype=np.float32) | ||
| diff = out_f32 - gold_f32 |
There was a problem hiding this comment.
When both files contain the expected FP8 NaN code at the same position, decoding produces two nan values and their subtraction remains nan, so the aggregate MSE is nan and an otherwise byte-identical result is reported as fail. This prevents the comparator from validating the Inf/NaN behavior supported by the generator and kernels; compare NaN masks separately and compute errors only over the remaining finite lanes.
Useful? React with 👍 / 👎.
| // L2：8 个 1bit → 一个字节值 [0,255] | ||
| Tile<Location::Vec, uint32_t, TileM, PW, BLayout::RowMajor, ValidR, 8> e8u, pos8, pow8, bits8; | ||
| TCVT(e8u, e1_8); // bf16 0/1 → uint32 0/1 | ||
| TCI(pos8, static_cast<uint32_t>(0)); // 0,1,2,...,7（逐列递增） |
There was a problem hiding this comment.
When TileM > 1, TCI generates indices from the tile's linear physical position rather than restarting at each row, so after the first row this produces shifts beginning at PW instead of 0; the mirrored 16-bit path has the same problem. Consequently the supposedly complete scale packer loses or misplaces L2/L3 bits for every later row in a tile once the placeholder stages are replaced; reduce the generated indices modulo 8 and 16 before shifting, as other row-wise TCI users do.
Useful? React with 👍 / 👎.
kBudgetElems 原按 sizeof(__half) 预算, 令 kTilemMax=128, 但 data pass 的 fp32 tile_f 物理字节 = TileM*BlockSize*4 = 16KB, 超出 TSize 仅编码 <=8KB 2 的幂 的上限 (pto_tile.hpp TilesizeCode -> __tilesize_unknown), 大 M 编不过。改按 sizeof(float) 预算 -> kTilemMax=64 (tile_f=8KB 合法)。M=512/N=256 现可编, gfsim 35236 cyc / 0.269 cyc-per-elem。探针 cpp 参数化为 -DPM/-DPN 便于扫尺寸。 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: c3858ec549
ℹ️ 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".
| // Reg::Interleave. Blocked on TINTERLEAVE/TDEINTERLEAVE not being | ||
| // exposed in the -D__linx header (RECORD 问题5); insert the even/odd | ||
| // zip here once available. | ||
| TSTORE(gs, scale_e8m0); |
There was a problem hiding this comment.
For every non-tail OCP invocation, this writes planar [scaleRows, Post] bytes instead of the public parity-interleaved mxScale layout. Even the registered NONTAIL_OCP_FP4 case has one real and one padding block row, so it emits all real scales followed by zeros while the generator and consumers expect each real scale adjacent to its padding byte; zip the even/odd block rows before storing, including in the big-BlockSize implementation.
Useful? React with 👍 / 👎.
| TABS(abs_h, xin); | ||
| tile_maxh max_h; | ||
| TROWMAX(max_h, abs_h); | ||
| TCVT(max_bf, max_h); // half -> bf16 |
There was a problem hiding this comment.
For InT=__half, this conversion uses the wrapper's default LINX_RNONE, whereas the copied AscendC ComputeMaxExpOcpHalf explicitly uses CAST_TRUNC. If a block maximum is just below a power of two (for example the largest FP16 value below 2.0), rounding half-to-BF16 can produce 2.0 while truncation retains the lower exponent, selecting a scale twice as large; request truncation or extract the FP16 exponent without a rounding conversion here and in the mirrored tail/non-tail paths.
Useful? React with 👍 / 👎.
- PERF_newcalc_M512N256.md：逐项解释 gfsim 打屏指标 + M512N256 实测分析 - swim.json：对应 swimlane(Perfetto) 输出 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 48a59b7bf8
ℹ️ 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".
| # unreliable due to toolchain<->emulator skew; that is a separate documented caveat. | ||
| CONFIGS = { | ||
| "TAIL_CUBLAS_FP8": {"M": 8, "K": 32, "algo": "CUBLAS", "kernel": "tail", "dtype": "FP8", "driver": "tail_cublas_fp8", "blocked": False, "scale_layout": "compact"}, | ||
| "TAIL_OCP_FP4": {"M": 8, "K": 64, "algo": "OCP", "kernel": "tail", "dtype": "FP4", "driver": "tail_ocp_fp4", "blocked": False, "scale_layout": "compact"}, |
There was a problem hiding this comment.
With the pinned PTO 0.58.1 toolchain, selecting this configuration always reaches the documented fp32 -> fp4 TileLogicalShapeMatch static assertion, and README.md explicitly states that TAIL_OCP_FP4 cannot produce even an object file. Because it is marked unblocked and included in the default configuration list, every default precision run attempts a configuration that cannot reach QEMU or comparison; mark it blocked or remove it until the shape contract is fixed.
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.
When Axis / BlockSize is odd, scaleRows reserves an additional zero padding row, but this loop stores only rows kb < numKb; the mirrored big-BlockSize implementation behaves the same way. The registered static buffers happen to start zeroed, while callers providing an uninitialized or reused scale buffer expose stale bytes in the required padding row, so explicitly clear that row rather than relying on caller initialization.
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.
The precision documentation describes byte/nibble comparison against the golden output, but this threshold can report pass even when every FP8 code is wrong: for example, comparing 1.0 with the adjacent E4M3 value 1.125 in every lane gives MSE 0.015625. The FP4 comparator repeats the same pattern and similarly hides sparse code errors; compare the raw FP8 bytes and unpacked FP4 nibbles exactly (while reporting decoded metrics separately) so quantization regressions are not accepted.
Useful? React with 👍 / 👎.
| print("[3/4] Running QEMU...") | ||
| elf_path = find_elf(cfg) | ||
| run_qemu(elf_path) |
There was a problem hiding this comment.
The compare directory is reused across runs, but neither generation nor this execution step removes output.bin and scale_output.bin. In particular, --skip-compile can reuse a non-RES_CHECK ELF that exits successfully without writing files, after which the comparator reads outputs left by an earlier run and can report a false pass; unlink both result files immediately before QEMU so each comparison proves that the current execution produced them.
Useful? React with 👍 / 👎.
- 新增 ISSUE_linx_tileop_fp4_tile_size_bits.md: 根因=Linx-TileOP-API type_traits<__fp4_e2m1x2>::bits=8(打包容器) vs 模型 ElementBits(FP4)=4, 致 tile size 512B(翻倍)、运行期 Block.cpp:1155 row 8≠16;修复须工具链头 fp4 tile size 位宽 8→4(256B),模型侧 31f7a8f 已就位,kernel 不改 - tail_ocp_fp4 tile_o 改元素列(PW/BlockSize)对齐 31f7a8f,过编译期 TileLogicalShapeMatch - 新增 tabs_trowmax_probe: 实证 TABS(bf16)/TROWMAX(uint16) 已解、TCMPS(uint32) 未解 - Makefile 增 TABS_TROWMAX_PROBE 条目 - RECORD: 问题9/10 标已解决、17 未解决、14/16/18 补代码仓 issue 编号 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 2ade027ec7
ℹ️ 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".
| parser.add_argument("--qemu", default=QEMU, help="QEMU path") | ||
| args = parser.parse_args() | ||
|
|
||
| global QEMU |
There was a problem hiding this comment.
Every invocation of this precision pipeline fails during Python parsing because main() reads QEMU in default=QEMU before declaring that name global; Python raises SyntaxError: name 'QEMU' is used prior to global declaration, so none of the generation, compilation, QEMU, or comparison stages can run. Move the declaration to the start of main() or avoid mutating the module global.
Useful? React with 👍 / 👎.
| compile_operator "$REPO_ROOT/test/kernel/fa" "fa" | ||
| compile_operator "$REPO_ROOT/test/kernel/sort" "sort" | ||
| compile_operator "$REPO_ROOT/test/kernel/deepseek" "deepseek" | ||
| compile_operator "$REPO_ROOT/test/kernel/quant/dynamic_mx_quant" "dynamic_mx_quant" |
There was a problem hiding this comment.
The commit adds both dynamic quant operators but registers only dynamic_mx_quant in the script advertised as compiling all operators. A repo-wide search of the reviewed tree finds dynamic_hi_f4_quant/compile.all but no caller outside its own directory, so running this umbrella script can report full compilation success without ever building the newly added hi-f4 framework or probes; add a corresponding compile_operator entry.
Useful? React with 👍 / 👎.
| // fp4 output tile: ELEMENT-column shape (physical PW, valid BlockSize), | ||
| // matching tile_f so TCVT dst/src pass TileLogicalShapeMatch. gfrun packs | ||
| // two 4-bit elements per byte via BytesOf(fp4) (SuperScalarModel 31f7a8f). | ||
| using tile_o = Tile<Location::Vec, OutT, TileM, PW, BLayout::RowMajor, TileM, BlockSize>; |
There was a problem hiding this comment.
With the default packed __fp4_e2m1x2 output, declaring tile_o with PW columns makes the conversion emit one FP4 nibble per byte instead of two per byte, while gm_y and the caller allocate only N/2 bytes. The repository's own ISSUE_tcvt_fp4_shape_contract.md records that this widened shape produces 512 bytes instead of 256 for the registered M=8,N=64 case and corrupts the packed output; the tile must retain physical PW/2 and valid BlockSize/2 dimensions, with the shape-contract problem fixed outside the kernel.
Useful? React with 👍 / 👎.
| if constexpr (IsCublas && !kRegBitcast) { | ||
| return 4; // fp32 scratch-HBM roundtrip binds a 32b tile | ||
| } else { | ||
| return static_cast<int>(sizeof(InT)); // bf16/fp16 = 2, fp32 = 4 |
There was a problem hiding this comment.
For BF16/FP16 OCP calls this returns a 4096-element budget based on the 2-byte input, but the active tail, plain non-tail, and big-BlockSize FP4 implementations materialize full physical float tiles before reduction and quantization. Shapes such as tail M=64,BlockSize=32 or non-tail BlockSize=128,Post=64 therefore derive 4096-element FP32 tiles (16 KiB), exceeding the documented 8192-byte tile limit and failing to build; size these migrated OCP paths using their 4-byte widest intermediate.
Useful? React with 👍 / 👎.
1 participant
No description provided.