feat(dmxq): 迁移 dynamic_mx_quant 算子套件到 ops-20260904 + 新 multi_thread 布局 - #111
ziyang-cheng wants to merge 9 commits into
Conversation
将 dynamic_mx_quant / dynamic_hi_f4_quant 全套（kernel 头 + 测试 driver/probe +
gen/compare 脚本 + RECORD/DESIGN/ISSUE 文档）移植到 tag ops-20260904，并按上游新
kernels 布局归入 multi_thread：
- kernels/multi_thread/quant/{dynamic_mx_quant,dynamic_hi_f4_quant}/
- test/kernel/multi_thread/quant/{dynamic_mx_quant,dynamic_hi_f4_quant}/
- include 前缀统一为 "multi_thread/quant/..."；Makefile include 深度对应 +1。
4-PE res_check 收尾改用上游官方 test/common/multi_thread_res_check.h（输入屏障
res_check_publish_inputs + 输出屏障 res_check_wait_for_all + PE0 落盘），移除本地
spmd_res_check.h；readBinary/writeBinary 保持上游原版（gfrun writev EFAULT 修复已
在配套 model 侧覆盖旧 printf 挂起）。
compile_all.sh 注册 multi_thread/quant/dynamic_mx_quant。10 个正式 kernel（cuBLAS/
OCP-FP4/OCP-FP8 全族 plain+4PE+bigbs+dyn）+ res_check 在配套工具链（llvm 1ae4ee39
+ TileOP-API 804eb03）下全部干净编译。
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: 0ad545eb1b
ℹ️ 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.
Python rejects this module at parse time because main() reads QEMU when constructing the --qemu argument before declaring that name global. python3 -m py_compile reports SyntaxError: name 'QEMU' is used prior to global declaration, so none of the precision-check modes can run until the declaration is moved before the first reference or the global assignment is removed.
Useful? React with 👍 / 👎.
| #ifdef RES_CHECK | ||
| writeBinaryFile(CHK_DIR "/output.bin", (uint8_t*)y, sizeof(y)); | ||
| writeBinaryFile(CHK_DIR "/scale_output.bin", (uint8_t*)scale, sizeof(scale)); |
There was a problem hiding this comment.
With the default PPE=4, all four SPMD threads execute these writes to the same filenames without waiting for the other PEs to finish. writeBinaryFile opens each file with O_TRUNC, so an early PE can snapshot incomplete shared output and concurrent PEs repeatedly truncate it; the other four-PE drivers in this change avoid this with res_check_wait_for_all followed by a tid == 0 guard. Apply the same synchronization and single-writer protocol here before treating the dumped files as precision results.
Useful? React with 👍 / 👎.
| BENCH_ROOT = TEST_DIR.parents[3] | ||
| COMPARE_ROOT = BENCH_ROOT / "compare" |
There was a problem hiding this comment.
TEST_DIR is the dynamic_mx_quant test directory, so TEST_DIR.parents[3] resolves to benchmark/one-level-arch/test, not benchmark/one-level-arch. Consequently generation uses test/compare and find_elf() searches under test/output, while Makefile.common emits both artifacts under the one-level-arch root; every pipeline run will therefore fail to find the ELF after compilation.
Useful? React with 👍 / 👎.
| # 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.
Selecting TAIL_OCP_FP4 generates an 8×64 input and golden result, but compile_elf() supplies no PM/PN overrides and tail_ocp_fp4.cpp now defaults to 512×256. The executable therefore reads and writes buffers for 512×256 while the comparison files describe 8×64, guaranteeing a short input and output/scale size mismatches rather than a valid precision measurement.
Useful? React with 👍 / 👎.
| def compile_elf(type_name: str, compiler_dir: str): | ||
| env = os.environ.copy() | ||
| env["COMPILER_DIR"] = compiler_dir |
There was a problem hiding this comment.
The required --compiler-dir value is written only into this local env mapping, but run() has no environment parameter and invokes subprocess.run() with the inherited process environment. When COMPILER_DIR was not already exported, the documented command fails in Makefile.common; when it was exported, the CLI argument is silently ignored and may use the wrong compiler.
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 scales are documented here as requiring a direct exact byte comparison, but the pass condition uses aggregate mse < 0.1. For a typical 4096-byte scale tensor, even a byte that differs by one produces an MSE of only about 0.00024 and is reported as passing; multiple incorrect scale entries can therefore be silently accepted. Base this status on exact equality or zero maximum error instead.
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.
Any generation, compilation, execution, or lookup exception is caught and recorded here, after which the script prints the summary and exits successfully; numerical comparisons are likewise invoked with check=False and their fail text is never converted into a failing exit status. As a result, automation sees status 0 even when every requested precision case errors or reports mismatches, allowing broken experiments to appear successful.
Useful? React with 👍 / 👎.
| // Compile-only: instantiate the bigbs auto-route at BS=128 (not res-checked). | ||
| 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.
These calls are labeled compile-only but are not guarded by #ifndef RES_CHECK, so a precision run executes the BS=128 route plus four additional dtype/route instantiations before writing the result of the actual 32×64 case. Those unrelated kernels operate on large throwaway buffers and can trigger a runtime or emulator failure that prevents the requested result from being dumped, while also substantially distorting execution time; guard the compile-coverage block as the neighboring nontail cuBLAS driver does.
Useful? React with 👍 / 👎.
| // scale even-pads the quant-axis block count: scaleRows = ceil_even(numKb). | ||
| // The trailing padding block-row is left zero. Layout is PLAIN planar | ||
| // [scaleRows, Post] = PTO-ISA Shared B-scale [G,N] (ADR-0101); no interleave. | ||
| constexpr int scaleRows = ((numKb + 1) / 2) * 2; |
There was a problem hiding this comment.
When numKb is odd, scaleRows reserves one extra row and the documented layout requires that padding row to contain zero, but every store is inside loops bounded by the real block range [0, numKb). The padding row is therefore only zero in these tests because their static buffers are initially zero; callers that provide reused or uninitialized output memory receive nondeterministic padding bytes, unlike the tail OCP-FP4 implementation which explicitly stores zero into its odd padding slot.
Useful? React with 👍 / 👎.
| make TESTCASE=dynamic_mx_quant TYPE=TAIL_CUBLAS_FP8 diss | ||
| make TESTCASE=dynamic_mx_quant TYPE=NONTAIL_CUBLAS_FP8 diss | ||
| make TESTCASE=dynamic_mx_quant TYPE=TAIL_OCP_FP4 diss | ||
| make TESTCASE=dynamic_mx_quant TYPE=NONTAIL_OCP_FP4 diss |
There was a problem hiding this comment.
The operator is now registered with the repository-wide compilation script, but this build list only exercises the four legacy/plain type names and a probe. The Makefile also exposes the newly added TAIL_OCP_FP8, TAIL_OCP_FP8_DYN, TAIL_CUBLAS_FP8_4PE, NONTAIL_CUBLAS_FP8_4PE, and NONTAIL_OCP_FP4_4PE drivers, none of which are invoked here, so compile_all.sh can report the migrated operator successfully compiled while the principal new multi-thread sources have never been built.
Useful? React with 👍 / 👎.
RECORD.md 补问题25:hosted musl __libc_start_main 发 ppoll,gfrun 白名单 无 handler → 启动即 abort;根因/与问题23 区别/do_ppoll 修复/tail_ocp_fp8 512x256 逐字节 pass 验证。ISSUE_gfrun_ppoll_libc_startup.md 组件清单以 Bench PR#111 链接给出 + 复现步骤。 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
tail_ocp_fp8_dyn.cpp 是迁移时唯一漏切官方 helper 的 4-PE driver:原先 4 个 PE 都 readBinaryFile + 两次 writeBinaryFile(O_TRUNC)同写共享 buffer, 无输入/输出屏障 → 落盘竞态(问题24 记录的 _dyn scale 写空变体即此), 512x256 碰巧过=PE0 I/O 最重自然最后完成的时序运气。 改用与其余 5 个 driver 一致的 multi_thread_res_check.h:PE0-only 读输入 + res_check_publish_inputs 输入屏障 + res_check_wait_for_all 输出屏障 + PE0-only 落盘。验证 4-PE 逐字节:512x256 与 256x64(均无尾块)output/scale 全 pass (MSE=0,output MaxAE=0.0117=fp8 LSB)。 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
RECORD 问题26: PR#510 恢复 FP4 写侧(前置澄清)+ 4 处新 model 缺口 (compare-select bit-width/TSTORE 描述符 packed/E2M1 编码 RNE 漏 code0/ NORM TSTORE packed 行 stride)+ 工具链 tile-size, tail_ocp_fp4 逐字节 pass + 349/349 无回归。ISSUE_fp4_pack_tcvt_regression 标记 PR#510 已修; ISSUE_linx_tileop_fp4_tile_size_bits 标记 804eb03 仍存在+本地已修。 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
- 问题25 ppoll handler 缺失【SuperScalarModel issue554】 - 问题26 emulator NORM Local TSTORE 不支持打包 4-bit(描述符 size + 源行 stride 均按字节口径,应按元素位宽,对照 pto-spec DerivedTileRows/ TileMemoryElementAddress)【issue557】+ ISSUE_gfrun_norm_tstore_packed4bit.md - 问题27 TCVT fp→E2M1/E1M2 编码最近邻漏 code0(小值抬到最小正档偏离 RNE) + pto-spec 参考模型无 E2M1 编码器【issue558】+ ISSUE_gfrun_tcvt_e2m1_rne_code0.md - 问题14 补 compare/select sibling(IsCompatibleDataTile)说明(同 issue254 根) - 问题16 更新:工具链侧已由官方 ddd07b9 解决(Linx-TileOP-API issue30) - ISSUE_fp4_pack_tcvt_regression 标已修复(PR#510); ISSUE_linx_tileop_fp4_tile_size_bits 标已修复(ddd07b9) 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: 18bf0167f6
ℹ️ 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".
| } else if constexpr (std::is_same_v<InT, float>) { | ||
| tile_x abs_f; TABS(abs_f, xin); | ||
| tile_maxf max_f; TCOLMAX(max_f, abs_f); | ||
| TCVT(max_bf, max_f); // fp32 -> bf16 |
There was a problem hiding this comment.
For InT=float, a block maximum just below a power of two can round up when this TCVT narrows it to bf16; masking the exponent afterward then selects the next exponent, doubling the emitted scale and halving the quantized values. The tail implementation avoids this by applying FP32_EXP_MASK to the FP32 maximum before narrowing, and the documented FP32 OCP path likewise requires direct FP32 exponent extraction, so this branch should do the same rather than derive the exponent from a rounded bf16 value.
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.
This new script runs each make sequentially without set -e or accumulated status, so its exit code is only that of the final FP4_PROBE build. If any plain or res_check kernel build fails but that probe succeeds, direct callers receive status 0 and benchmark/one-level-arch/compile_all.sh enters its if bash compile.all success branch and prints that DynamicMxQuant compilation completed, masking the broken kernel.
Useful? React with 👍 / 👎.
| global_iterator<gm_s, tile_sstore_r> s_iter_r(scale + kb * Post); | ||
| auto gs = s_iter_r(0, numN); |
There was a problem hiding this comment.
When Post % TileN != 0, this iterator's tile has Cols=N_tail, so indexing it with numN advances only numN * N_tail columns rather than the required numN * TileN. For example, Post=96 and TileN=64 writes the 32 tail scales over columns 32–63 and leaves columns 64–95 unwritten, even though the data tail is loaded and stored at column 64. Fold numN * TileN into the scale base (or otherwise address that absolute column) before storing the tail scales.
Useful? React with 👍 / 👎.
| biased_exp = exp + 7 | ||
| 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.
E4M3 subnormals represent mant * 2^-9, as the comparator's decoder also reflects, but this encoder divides subnormal inputs by 2^-6. For example, the exactly representable value 3 * 2^-9 is encoded as zero instead of mantissa code 3, and the preceding cutoff also zeros values that should round to the first subnormal. Any precision case containing elements much smaller than their block maximum can therefore receive an incorrect golden.bin and report a valid kernel as failing.
Useful? React with 👍 / 👎.
- nontail_ocp_fp4_4pe（plain 4-PE，fp16 in）4 线程 gfrun 端到端跑通： output=pass (MSE=0.019) / scale=pass (MSE=0.0068)。scale 7/1024 off-by-one 系 golden >>16 截断 vs spec-RNE 差异（model 正确），非工具问题，记入 README。 - 移植 nontail_ocp_fp4_bigbs.cpp 单 PE driver + Makefile TYPE=NONTAIL_OCP_FP4_BIGBS （方案A 切归约轴，bf16 in）：单线程 gfrun output/scale 全逐字节 pass (MSE=0)。 - README 新增「golden 与 spec 的实现差异」小节，更新 nontail_ocp_fp4 / bigbs 状态。 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
tail_cublas_fp8_4pe / nontail_cublas_fp8_4pe / nontail_cublas_fp8_bigbs 三者随 gfrun ppoll 启动缺口修复(issue554)+官方屏障就位, 端到端 gfrun res_check 全 R2=0 逐字节 pass: - tail_cublas_fp8_4pe (4-PE, fp16, 512x256): output MSE=0/MaxAE=0.0098, scale byte-exact - nontail_cublas_fp8_4pe (4-PE, fp16, 512x256): output MSE=0/MaxAE=0.0117, scale byte-exact - nontail_cublas_fp8_bigbs (单PE, bf16, 128x32 BS128): output/scale 全逐字节 MSE=0 driver 本就带屏障, 本轮为新 tag 上的验证, 无 kernel/driver 代码改动。 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: a16151d8bb
ℹ️ 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".
| COMPARE_ROOT = BENCH_ROOT / "compare" | ||
|
|
||
| QEMU = os.environ.get("QEMU", "/remote/lms60/c00622284/qemu/LinxBlockModel/build/qemu-linx") | ||
| QEMU_ARGS = ["-blk_optimize", "force_tb_chained", "-s", "4096M"] |
There was a problem hiding this comment.
When --type TAIL_OCP_FP4 reaches execution, this argument list leaves the model at its default single PE, while tail_ocp_fp4.cpp hardcodes a four-PE kernel and its RES_CHECK path waits for done[1..3]; the run therefore hangs rather than producing results. The repository's multithread runner demonstrates the required CLI form at test/kernel/multi_thread/res_check_all.py:202 (-s softcore.multiThreadNum=4), so select that setting for this configuration.
Useful? React with 👍 / 👎.
| parser.add_argument("-l", dest="elf_list", default=None, | ||
| help="ELF list file") | ||
| parser.add_argument("--cmp-root", dest="cmp_root", | ||
| default=os.path.abspath(os.path.dirname(__file__) + "/../../../compare"), |
There was a problem hiding this comment.
Invoking this comparator directly without --cmp-root resolves the default to benchmark/one-level-arch/test/kernel/multi_thread/compare, but Makefile.common:113 defines CHK_DIR under benchmark/one-level-arch/compare. Consequently the documented -d/-l modes report every output as missing files unless callers discover and override the incorrect default.
Useful? React with 👍 / 👎.
| # dtype defaults to FP8 but is inferred from the ELF/driver name when unset. | ||
| compare_out = compare_fp4 if dtype.upper() == "FP4" or elf_name.endswith("fp4") else compare_fp8 |
There was a problem hiding this comment.
When the standalone comparator is given an FP4 variant such as dynamic_mx_quant_nontail_ocp_fp4_bigbs.elf or dynamic_mx_quant_nontail_ocp_fp4_4pe.elf without an explicit --dtype, the default remains FP8 and neither name satisfies endswith("fp4"). Their packed FP4 bytes are therefore decoded as E4M3 values, making the reported MSE and pass/fail status invalid; infer FP4 from the complete driver-name token rather than only a terminal suffix.
Useful? React with 👍 / 👎.
| def reduction_groups(rows: int, cols: int, kernel: str, block_size: int): | ||
| groups = [] | ||
| if kernel == "tail": | ||
| numKb = cols // block_size |
There was a problem hiding this comment.
When the generator is called with tail cols or nontail rows that are not divisible by block_size, this integer division silently omits the final partial block. Because compute_golden() preinitializes every quantized value and scale to zero, the uncovered input elements are emitted as plausible but incorrect zero golden data; the kernels themselves require whole reduction blocks, so the generator should reject these dimensions instead of producing an invalid reference.
Useful? React with 👍 / 👎.
新工具链头把 tile 上限从 8KB 抬到 256KB（StorageBytes 须为 [128B,256KB] 内 2 的幂，pto_tile.hpp TilesizeCode），bigbs（方案A 切归约轴）赖以存在的 tile-size 墙消失。改动： - nontail_ocp_fp4 / nontail_cublas_fp8 public 入口去掉 bigbs else 分支， pick_tilen 对大 BS 无解时回退 TileN=align 走单块 load；小 BS 首选值不变。 - 放宽两 plain 的 tile-size 断言 4096 -> 65536（32b 中间量 *4 = 256KB 上限）。 - 删除只被 bigbs 路由使用的 common.hpp::max_rsub。 - 两个 _bigbs.hpp 移入 bak/；fp4 bigbs driver 改调 public 入口。 BS=128 plain 单块实测 gfrun 逐字节 == 旧 bigbs；全 8 用例零回归 （两 bigbs 转 plain MSE=0 byte-exact，其余 6 例结果与收敛前完全一致）。 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: 7a5816b121
ℹ️ 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".
| // Compile-only: exercise the half / fp32 input branches (not res-checked). | ||
| 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 RES_CHECK is enabled, these two calls labeled compile-only still execute after the checked bf16 kernel. A precision run must therefore successfully execute the unrelated half and fp32 implementations—even though their outputs are discarded—so an emulator/runtime failure in either branch prevents the valid checked result from being written. Guard these calls with #ifndef RES_CHECK, as done for the corresponding nontail cuBLAS driver.
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 numKb is odd, scaleCols allocates an extra column whose required E8M0 value is zero, but the loop only stores columns [0, numKb). The bundled drivers hide this because their static buffers begin zeroed; callers supplying reused or uninitialized scale memory receive arbitrary padding bytes, unlike the tail FP4 kernel which explicitly writes the odd padding column.
Useful? React with 👍 / 👎.
| // === data pass (NEWCALC: 位补求倒数, 复用 xh) === | ||
| auto sh_u16 = reinterpret_tile<int16_t>(shared_bf); | ||
| TXORS(sh_u16, sh_u16, RECIP_XOR_NOT); // 0xFFFF - bits (重打 I16 标签) | ||
| TSUBS(sh_u16, sh_u16, RECIP_COMPL_SUB); // -> 0x7F00 - bits = 2^(8-E_max) | ||
| t_fb recip_f; TCVT(recip_f, shared_bf); // bf16 -> fp32 |
There was a problem hiding this comment.
For a block containing infinity or a propagated NaN, shared_bf is non-finite, but this unguarded bit-complement turns its bits into an unrelated reciprocal (for example, infinity produces a negative-infinity-like value) before quantization. The AscendC OCP implementation and the adjacent FP4 kernel instead select BF16_NAN_PATTERN for this case, so the formal static FP8 kernel—and the dynamic copy of this sequence—returns saturated or otherwise incorrect data for special-value inputs unless the inf/zero/special guards are restored.
Useful? React with 👍 / 👎.
| def scale_recip_ocp(group_vals, emax: int) -> tuple: | ||
| exp = [(f32_to_bf16_bits(v) & BF16_EXP_MASK) for v in group_vals] | ||
| max_exp = max(exp) |
There was a problem hiding this comment.
When data is generated with --in-dtype fp16 --algo OCP, this always truncates the promoted FP16 value to BF16, while the actual FP16 OCP path uses TCVT's round-to-nearest-even conversion. A block maximum just below a power of two can therefore produce a golden scale one exponent below the kernel's scale and correspondingly different quantized data; the repository's documented FP16 case already observes these off-by-one scale bytes. Make the OCP reference aware of the input dtype and apply BF16 RNE for FP16 instead of routing every input through the truncating helper.
Useful? React with 👍 / 👎.
| if constexpr (seg_tail > 0) { | ||
| process_tile.template operator()<TileM, seg_tail>(row_begin + seg_full * TileM); | ||
| } |
There was a problem hiding this comment.
When a four-PE shape gives a per-PE SubM that is not divisible by TileM, this instantiates the boxed seg_tail path even though the current reduce-to-TCVT model cannot execute it. For example, M=1000, N=32, and BlockSize=32 gives each PE 250 rows and a 58-row tail, which hits the documented logical-shape assertion at runtime instead of producing output. Until that path is implemented with compatible dynamic-valid tiles, constrain such shapes at compile time rather than accepting every positive M and dispatching into a known-crashing path.
Useful? React with 👍 / 👎.
收敛提交（7a5816b）删了 _bigbs kernel、移入 bak/，但 README/DESIGN 仍把 bigbs 描述成活的路由目标，与代码自相矛盾。本次补齐： - README 顶部加退休横幅；状态总览去掉"另加 2 个大 BlockSize 专用模板"； 路由节改为"恒走 plain 单块"；「大 BlockSize 变体」节标注为历史记录。 - DESIGN §4.1 路由描述改为"大 BS 回退 plain，方案A/max_rsub 已退休"。 历史条目（状态表/RECORD 引用）保留并标注为退休前记录，不全量重写。 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: ffbf0600d1
ℹ️ 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".
| // reduce-axis block count, even-aligned (padding block-row left zero). | ||
| constexpr int scaleRows = ((numKb + 1) / 2) * 2; |
There was a problem hiding this comment.
When Axis / BlockSize is odd (including the default 32/32 case), scaleRows reserves an additional padding row, but the loops only write rows below numKb. The test driver hides this by using a zero-initialized static buffer; callers that pass reused or uninitialized output memory receive arbitrary bytes in the compact scale tensor's required padding row. Explicitly write zero to that row before returning.
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 E4M3 NaN encoding, such as 0x7f, both decoded values are NaN, so subtraction and np.mean produce NaN and mse < 0.1 is false. The comparator therefore reports a failure for byte-identical special-value results; handle matching NaNs explicitly or compare the encoded bytes before calculating finite-value error metrics.
Useful? React with 👍 / 👎.
| args = parser.parse_args() | ||
| gen_all(args.out_dir, args.M, args.K, args.algo, args.kernel, args.dtype, | ||
| args.seed, args.scale_layout, args.in_dtype, args.block_size) |
There was a problem hiding this comment.
The CLI independently accepts every algorithm and output dtype, even though this module documents only OCP with FP8/FP4, CUBLAS with FP8, and DYNAMIC_RANGE with FP4. For example, --algo CUBLAS --dtype FP4 succeeds and emits a plausible golden file while scale_recip_cublas still uses the FP8 1/448 constant, so the reference cannot correspond to any supported kernel. Validate the pair after parsing rather than silently creating invalid experiment data.
Useful? React with 👍 / 👎.
1 participant
将 dynamic_mx_quant / dynamic_hi_f4_quant 全套（kernel 头 + 测试 driver/probe + gen/compare 脚本 + RECORD/DESIGN/ISSUE 文档）移植到 tag ops-20260904，并按上游新 kernels 布局归入 multi_thread：
4-PE res_check 收尾改用上游官方 test/common/multi_thread_res_check.h（输入屏障 res_check_publish_inputs + 输出屏障 res_check_wait_for_all + PE0 落盘），移除本地 spmd_res_check.h；readBinary/writeBinary 保持上游原版（gfrun writev EFAULT 修复已 在配套 model 侧覆盖旧 printf 挂起）。
compile_all.sh 注册 multi_thread/quant/dynamic_mx_quant。10 个正式 kernel（cuBLAS/ OCP-FP4/OCP-FP8 全族 plain+4PE+bigbs+dyn）+ res_check 在配套工具链（llvm 1ae4ee39