add dynamic_mx_quant test kernel on tag ops-20260828 - #102
ziyang-cheng wants to merge 51 commits into
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>
- 复现仓拆为编译仓(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>
- 新增 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>
将 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>
- 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>
- 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>
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>
- PERF_newcalc_M512N256.md:逐项解释 gfsim 打屏指标 + M512N256 实测分析 - swim.json:对应 swimlane(Perfetto) 输出 Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
- 新增 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>
- RECORD.md 新增问题19:工具链 TLOAD/TSTORE 内联汇编模板把 B.IOR 的 GM 行步长 按元素数发射(漏 ×bits/8),违反 pto-spec ADR 0074 字节步长契约,导致非 1B dtype 的 GM 访存半行错位;已在工作目录 installed 头修复并实证闭环(Linx-TileOP-API issue31) - ISSUE_e4m3_overflow_inf / e8m0_cvt / rel0812_isa_dtype_support / tcvt_no_lb2 四项标题追加【已解决】(对应 emulator 侧已修复) Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…364) - tail_ocp_fp4: 按 InT if constexpr 在非窄化域 floor 取指数(half/fp32→fp32 mask、bf16 直接取 bf16 指数位), 规避 half→bf16 round-to-nearest 进位越 2^k 污染 shared scale(整块 2× 偏低) - newcalc probe: 位补(0x7F00-bits)求倒数替代 TRECIP; 接入官方精度流程 (run_precision_check PROBE_OCP_FP8_NEWCALC, fp16 in→e4m3 out, BS=32) + gen 传 --in-dtype - RECORD 问题21: fp32→e4m3 TCVT 溢出饱和漏「舍入进位越 exp_max」一路 (V∈[496,512) 如 500.0f 吐 0x78=256 非饱和 0x7e=448); 候选修法 overflow=exp>exp_max 无条件 clamp 实证 newcalc probe output 253→256/256 逐字节, 但补丁已回退未提交, 正式修复挂 SuperScalarModel issue364 - README/问题19: B.IOR 字节步长(非 1B dtype)表述更正 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
- 新增 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 / 非 mt 变体。 - 验证 (现有工具链, gfsim 08-25 build): gfrun -s softcore.multiThreadNum=4 → R2=0, 4 线程各退休 414 块 / 总 1656; gfsim --conf fourpe → Total Cycles=20037, 与 PERF §4.3 记录逐字符一致。kernel 功能正确。 - MT ELF ssrset=0, 不触发 gfsim Decoder ssrset-null 潜伏 bug (非 MT ELF ssrset=1 会崩, 与本 kernel 无关)。 - swim.json (mt4pe/noreload/单线程) 附作性能分析产物; pipe.out 因体积不提交。 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
- 去 probe 正名: probe_..._newcalc_mt.hpp -> dynamic_mx_quant_tail_ocp_fp8.hpp, 函数去 _mt 后缀,固定 SPMD 4-PE (删 -DMT 开关) - 常量整合 common.hpp (逐值等价,非改数值): RECIP_EMAX -> recip_emax_bits<e4m3>(), FP32_EXP_MASK 复用 common:38;位补两步常量保持内联 - 驱动拆分: tail_ocp_fp8.cpp (正式 4PE) + probe_ocp_fp8_newcalc.cpp (探针纯 single-PE); Makefile 加 TYPE=TAIL_OCP_FP8 - single-PE 探针 probe_..._newcalc.hpp 原样保留 - 单PE 精度实测 (M=256/N=32, gfrun res_check tid=0 前64行): scale 逐字节全对 MSE=0; data 5/2048 偏差 100% 命中 emulator issue364 (fp32->e4m3 进位越 exp_max [496,512) clamp 漏判吐 0x78=256 而非 0x7e=448),非 kernel 逻辑缺陷 - README 补实测精度证据 + 4-PE harness 崩溃现象 (driver 文件 I/O 被 4 线程各跑一次 抢 fd + BARG 断言,需按 tid 守卫);RECORD 记录问题22 (boxed 尾块 reduce stride); 新增 ISSUE_reduce_output_stride_tail.md Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
- 依 pto-spec 全仓搜索结果修正 DESIGN/RECORD:
- §2.3 y 元素 lane 表纠错为权威 {0,0.25,…,1.75}(含次正规+零),
锚定 hif4x2.asl HiF4X2FiniteDecomposition
- §1.3 补 y 规范 dtype 身份 HiF4X2(code14),TCVT-only,
与 E1M2X2(12) 数值等价但独立类型
- RECORD 问题3 定性:E6M2 非注册 tile dtype,位重构为规范意图正解
- RECORD 问题7 补 HiF4X2 规范目标 dst + matrix-functions MX 契约
- 新增 §3.3 目标形态伪码(后端补齐 subview/assemble 后):
含记法/立即数图例、复核结论(修正 E1_16 K8 误索引 bug)、
剩余 gap 清单(BF16 TROWMAX 白名单等)
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
对齐 model d8903938 rowReduce 强制 col=1 语义(Block.cpp:2069 stride=1), 把 tail_ocp_fp8 scale 链下游列向量 tile (t_hb/t_bfb/t_e8b/t_fb) 声明成 physical Cols=1,使 reduce->TCVT 形状契约通过。 - dynamic_mx_quant_tail_ocp_fp8.hpp: 列向量 tile physical Cols=1 迁移 - tcvt_reduce_shape_probe.cpp + Makefile: 最小复现探针 (TCVT_REDUCE_SHAPE_PROBE) - ISSUE_tcvt_e8m0_row_padding.md: bf16->e8m0 目的 row 被 128B 最小档撑翻倍根因链 - ISSUE_tcvt_fp4_shape_contract.md: 修正 lb2 描述 (TCVT 恒发 lb2=dst::Cols) - README/RECORD: 刷新 tail_ocp_fp8 现状(4-PE res_check output/scale=pass)+ 问题15 e8m0 regression 口径(52f56d5 修复被 930d9981 误回退, issues439)+ 问题22 并入 e8m0 row 翻倍面 + 问题23(gfrun res_check syscall-ABI 规避) Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
- 照 tail_ocp_fp8 母本重写为固定 SPMD 4-PE(get_thread_idx/run_pe/process_tile), reduce 下游列向量 tile physical Cols=1(匹配 model rowReduce col=1)。 - 保留 fp4 原有的 inf/zero/special 三守卫(TCMPS+TSEL),不随 fp8 母本省略。 - 输出 fp4 打包(每字节 2 个 e2m1),driver 改 PM=512/PN=256/RES_CHECK/4 线程。 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
单条 TSEL(x=0x1234, mask=0, k=0xABCD),期望保留 prior-dst=0x1234; 复现就地 TSEL 的 false-source 被读成 0(问题18 更深层)。Makefile 加 TYPE=TSEL_INPLACE_FALSESRC_PROBE。 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
- RECORD 加「0828-tag 补丁清单」(工具链/model/kernel 三仓补丁表 + 类型/对应问题)。 - ISSUE_tsel_inplace_falsesrc.md(已标 Linx-TileOP-API PR#41 根治)。 - ISSUE_fp4_pack_tcvt_regression.md(由 model 仓迁入;注 codex/pr-0.58.4 有原生修复)。 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
new combo(新 musl + codex gfrun)跑 res_check 时,musl 向 stdout 的 writev(1,iov,2) 无限重试挂起:codex 内存 Load 读不了 guest 高栈地址(iov=0xfffffff...→读回0→写0字节 →musl 死循环)。readBinary.h/writeBinary.h 把状态 printf+fflush 包进 #ifndef RES_CHECK 静音(文件 output.bin 走直接 write 系统调用不受影响;精度脚本按 returncode+output.bin 判定)。根因是 codex 内存/栈映射,留上游,勿硬改。见 ISSUE_gfrun_res_check_writev_hang.md。 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
- ISSUE_fp4_pack_tcvt_regression.md 按远端最新代码重写(去本地引用/去历史):坐实 31f7a8f(#314)写侧被 930d9981 revert 连坐删、官方各分支(含 codex/pr-0.58.4)均缺; 更正'codex 已原生修复'失实结论(实测 fp4 仍 fail,codex 只读侧非充分);加最小复现 探针 fp4_shape_probe WIDEN=on(6.0→码6应7)。对应 SuperScalarModel issues454。 - RECORD.md 前置 2026-09-01 状态节:端到端(fp8 pass/fp4 fail)、已落地清单、作废旧 'codex 已修复'判断、两层修复思路、工具链官方修复是 ddd07b9 非 bits8→4。 - README 状态总览:TAIL_OCP_FP4 由旧'编译崩'更新为当前'能编能跑、output=fail(模型 写侧缺)';TAIL_OCP_FP8 e8m0 提交号补注 codex 基线 60ce26fd。 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
- tail_cublas_fp8 / nontail_cublas_fp8 加 kPeNum 模板参做固定 4-PE SPMD (tail 按 M 行 boxed 切分复用 tail_ocp_fp8 骨架、nontail 按块行 kb 运行期切分), 各带 *_4pe driver + Makefile TYPE(TAIL/NONTAIL_CUBLAS_FP8_4PE) - 三个 cuBLAS kernel(tail / nontail plain / nontail bigbs)+ 共享 compute_cublas_core: cuBLAS 守卫复合条件从数据域 TAND/TOR 组合掩码改为嵌套 TSEL——修正 kernel 原写法 违反 PTO ISA(PTO-REQ-TEPL-COMPARISON-001:compare 产 packed predicate、TSEL mask 须 packed predicate、TAND 只作用 integer 且 reject packed);旧仿真器纵容,codex 会崩。 compute_cublas_core 删除过时的非合规 IDEAL VERSION 注释块 - nontail reduce 下游列向量中间 tile physical Cols=1,匹配 rowReduce 无条件 col=1 - codex 基线 gfrun 验证:tail 4 线程 output+scale pass(MaxAE=0.0098/0); nontail plain 4 线程 output pass(MaxAE=0.0117,scale 值 interleave==golden 仅问题5布局); bigbs 单线程 output 逐字节 MaxAE=0 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
- README: tail/nontail cublas 状态行标注固定 4-PE + codex gfrun 精度;新增 「cuBLAS 守卫掩码的 PTO ISA 合规写法」小节并收口整个 cuBLAS 系列(tail/nontail plain/bigbs + compute_cublas_core dead-reference 全合规) - RECORD: 删除误记的问题24(kernel 合规非编译器/model 问题);问题3 加 2026-09-01 supersede 注解(IDEAL VERSION 块已删、compute_cublas_core 现为合规 dead-reference); 问题17 措辞更正为「各 kernel 内联的 native CmpMode 版本」 - DESIGN: 指令表更正——「TAND/TOR/TNOT 守卫掩码代数」已移除(实测 0 调用), 改为嵌套 TSEL;比较行更正为原生 TCMPS<CmpMode> Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
- nontail_ocp_fp4 / _bigbs: fp4 输出 tile 从旧 [.,TileN/2] 字节列约定迁移到 element-列形 [.,TileN](与源 fp32 同列),使 fp32->fp4 打包 TCVT 走打包 specialization、不再落 ordinary 路径的 physical-Cols 校验而编译崩;存储侧仍 Post/2 字节,经 gm_y byte 域 + 字节基址折叠 y_iter 定位。与 tail_ocp_fp4 同一 迁移(0.58.3 工具链头已删 32B 列对齐 assert,无需 PW-padding)。当前 codex/ pr-0.58.4 基线下单线程与 4-PE 均可编译 + gfrun 跑通。 - nontail_ocp_fp4 增 kPeNum 模板参(默认 1=单线程零回归),按块行 kb 运行期连续 切 4 PE(与 nontail_cublas_fp8 同一范式);单 PE 路径写满全部块行、逐字节正确。 新增 driver nontail_ocp_fp4_4pe.cpp + Makefile TYPE=NONTAIL_OCP_FP4_4PE。 - 文档:RECORD 新增问题24(如实记录 4-PE gfrun res_check 落盘"仅 PE0 段有数据"的 规模相关现象,根因未定位、方向皆为待验证猜测,且在 nontail_cublas_fp8 上同样存在); README 的 NONTAIL_OCP_FP4 行同步该说明。 注:fp4 data 路径 output 在当前基线仍受模型侧 fp4 写侧缺失影响(issues454)。 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
模型 Block.cpp:2349 的 colReduce(TCOLMAX/…)按 dst->size/(validCol·elem) 反推物理行, 声明 physical row=1 时反推恰得 row=1、physCol=validCol,最干净且自洽(与 rowReduce Block.cpp:2328 无条件把物理列塌成 1 的不对称,见 RECORD 问题22 补充)。 把四个非尾轴 kernel 的 reduce 输出 + 下游行向量/累加器/scale/recip 载体从 physical row=BlockSize(或 R_sub) valid row=1 统一改为 physical row=1: - nontail_ocp_fp4(plain)+ _bigbs - nontail_cublas_fp8(plain 全列 + N_tail 尾列 [1,N_tail])+ _bigbs 等价性验证(可跑路径逐字节一致,与旧写法无差): - nontail_ocp_fp4 单 PE res_check:output/scale 与 physical row=BS 逐位相同。 - nontail_cublas_fp8 4-PE Post=256:output=pass (MaxAE=0.011719)。 - nontail_cublas_fp8_bigbs:output=pass (MaxAE=0)。 - 五个 TYPE 全部编译通过。 问题24(4-PE 落盘规模相关现象)与本改动无关、行为不变(再次佐证与 reduce 形状无关)。 RECORD 问题22 增补「rowReduce/colReduce 物理形状处理不对称」结论(含 Block.cpp 代码依据)。 docs:新增 ISSUE_gfrun_multipe_size_dependent_dump.md(问题24 的官方可复现 issue: 基线官方 SuperScalarModel 762a72c3 + 附录 A 前置补丁 diff + Bench PR PTO-ISA#102;含完整版本清单、 最小复现、已排除方向、确定性事实、未证实猜测方向;明确根因未定位);RECORD 问题24 链接该文件。 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: c58b141b00
ℹ️ 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.
Declare QEMU global before using it
Move global QEMU to the beginning of main() (or avoid mutating the module variable). Because QEMU is referenced as the default for parser.add_argument() earlier in the same function, Python rejects this entire script at parse time with SyntaxError: name 'QEMU' is used prior to global declaration, so no precision-check configuration can run.
Useful? React with 👍 / 👎.
| def compile_elf(type_name: str, compiler_dir: str): | ||
| env = os.environ.copy() | ||
| env["COMPILER_DIR"] = compiler_dir | ||
| run([ |
There was a problem hiding this comment.
Pass the compiler environment to make
Pass this env mapping into the subprocess invocation. When a caller supplies the required --compiler-dir argument without separately exporting COMPILER_DIR, the local assignment is discarded because run() invokes subprocess.run() with the inherited process environment; Makefile.common then stops with COMPILER_DIR is not set, making the documented CLI argument ineffective.
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.
Propagate precision failures through the process status
Return a nonzero status when any configuration errors or reports a failed comparison. Currently exceptions are converted into summary strings, comparison output is never interpreted, and main() falls through successfully, so automation receives exit code 0 even when QEMU failed, files were missing, or numerical comparison printed fail; this can allow broken experiment results to pass validation.
Useful? React with 👍 / 👎.
| // 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.
Skip compile-only kernel calls during result checking
Guard these throwaway instantiations with #ifndef RES_CHECK, as the non-tail cuBLAS driver already does. With res_check=on, the harness currently executes the half and fp32 calls described as compile-only—and tail_ocp_fp4.cpp and nontail_ocp_fp4.cpp do the same, including the large-block route—before writing the primary result files; in emulator/toolchain combinations where one of those explicitly runtime-blocked branches aborts, an otherwise valid primary kernel can never complete its precision check.
Useful? React with 👍 / 👎.
| int cap = budget / BlockSize; // TileSize upper on TileN | ||
| int postcap = ((Post + align - 1) / align) * align; // no need to exceed Post | ||
| if (postcap < cap) cap = postcap; | ||
| return (cap / align) * align; // floor to align (0 if cap < align) |
There was a problem hiding this comment.
Select a TileN compatible with the OCP Post dimension
Choose a legal aligned divisor of Post when dispatching to the non-tail OCP-FP4 path, or add a partial-column path there. For example, with bf16, BlockSize=32, and Post=192, this helper returns TileN=128, after which nontail_ocp_fp4_plain fails its Post % TileN == 0 assertion even though TileN=64 satisfies the alignment, tile budget, and dimension constraints; valid OCP shapes are therefore rejected solely by the tile-selection heuristic.
Useful? React with 👍 / 👎.
| } 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 | ||
| } else { // bf16 | ||
| tile_f xf32; TCVT(xf32, xin); // bf16 -> fp32(规避 TABS 拒 bf16) | ||
| tile_f abs_f; TABS(abs_f, xf32); | ||
| tile_maxf max_f; TCOLMAX(max_f, abs_f); | ||
| TCVT(max_bf, max_f); // fp32 -> bf16 |
There was a problem hiding this comment.
Mask the wide OCP maximum before narrowing to bf16
For half or fp32 input, clear the fp32 exponent/mantissa bits before converting the reduced maximum to bf16, as the tail OCP implementation already does. If a block maximum lies just below a power of two—for example the fp16 value 1.9990234375—this TCVT rounds it up to 2.0 before the subsequent bf16 mask, whereas the required exponent extraction and the golden generator truncate the original value; the resulting shared scale is one exponent too large and the entire block is quantized at half the intended magnitude. The same ordering appears in the big-BlockSize OCP path.
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.
Require exact equality for compact E8M0 scales
Use exact byte equality (or require max_ae == 0) for the compact scale result. This function documents that the compact E8M0 layout is directly byte-comparable, but the mse < 0.1 threshold lets sparse errors pass—for example, one off-by-one scale byte in a 64-byte output has MSE 0.015625 and is reported as pass—so incorrect block scales can escape precision validation.
Useful? React with 👍 / 👎.
| # --- plain builds --- | ||
| 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.
Build the supported tail OCP-FP8 kernel in compile.all
Add TYPE=TAIL_OCP_FP8 to this build list. The repository-wide compile_all.sh invokes this compile.all, the scoped Makefile defines a TAIL_OCP_FP8 target, and the accompanying README lists that production kernel among the five debugged implementations, but this script still builds only the older set of four; consequently the normal full-compilation check never compiles the newly supported kernel and cannot detect regressions in it.
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 71abf61567
ℹ️ 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.
Match the TAIL_OCP_FP4 generator dimensions
When TAIL_OCP_FP4 is selected, this entry generates an 8×64 input and golden output, but tail_ocp_fp4.cpp is compiled without PM/PN overrides and therefore processes its 512×256 defaults. Its readBinaryFile(..., sizeof(xbits)) rejects the short generated input, while the driver ignores that failure and writes 512×256-sized result files, so the comparison can only report size mismatches rather than testing this kernel. Either configure this entry as 512×256 or pass matching PM and PN definitions during compilation.
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.
Treat matching FP8 NaNs as equal
When an FP8 golden result and the kernel output both contain the E4M3 NaN encoding (0x7f or 0xff), both bytes decode to NaN, so their subtraction makes mse and max_ae NaN and this comparison reports fail even though the outputs match. This prevents the comparator from validating the special-value paths that the generator and kernels explicitly implement; compare matching NaN lanes by encoding or mask them as equal before computing the metrics.
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.
Reject incomplete reduction blocks
If the tail reduction dimension (cols) is not divisible by block_size, this floor division silently omits the final partial block; the nontail branch does the same for rows. compute_golden nevertheless emits full-sized files with every omitted element left at its initialized zero value, producing a plausible but incorrect golden result instead of reporting that the shape violates the kernels' exact-block contract.
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.
Reject unsupported algorithm and dtype combinations
The CLI accepts the algorithm and destination dtype independently even though the module documents only OCP with FP8/FP4, CUBLAS with FP8, and DYNAMIC_RANGE with FP4. Invocations such as --algo CUBLAS --dtype FP4 or --algo DYNAMIC_RANGE --dtype FP8 therefore complete successfully and write golden files by combining unrelated scaling and encoding rules, despite there being no corresponding operator semantics. Validate the pair after parsing so unsupported experiments fail rather than consuming fabricated reference data.
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.
Store tail-column scales at the full-tile offset
For a supported nontail cuBLAS shape with both full and partial column tiles, such as Post=96 and derived TileN=64, this tail scale iterator has physical width N_tail=32, so s_iter_r(0, numN) advances by numN * N_tail and writes at column 32 rather than the required numN * TileN column 64. The tail scales overwrite part of the preceding full tile and the actual trailing columns remain unwritten; fold numN * TileN into the base pointer and use column index zero, as the other folded scale stores do.
Useful? React with 👍 / 👎.
| constexpr int numKb = K / BlockSize; | ||
| // 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.
Initialize the compact scale padding
Whenever numKb is odd, including the default one-block case, scaleCols reserves an extra padding byte but the kernel stores only columns 0..numKb-1. The test drivers' statically zero-initialized arrays hide this, while a normal caller supplying reused or uninitialized output memory receives a stale padding byte instead of the required compact-layout zero. Explicitly store zero into the padded column, as dynamic_mx_quant_tail_ocp_fp4 already does; the same omission should also be corrected in the other implementations that describe their padding as merely "left zero."
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 150660a5dc
ℹ️ 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".
| make TESTCASE=dynamic_mx_quant TYPE=NONTAIL_OCP_FP4 diss | ||
|
|
||
| # --- res_check builds (for precision verification) --- | ||
| make TESTCASE=dynamic_mx_quant TYPE=TAIL_CUBLAS_FP8 res_check=on diss |
There was a problem hiding this comment.
Force res-check builds to recompile their objects
These commands run immediately after plain builds of the same sources, but changing res_check only changes compiler defines and is not represented in the object-file dependency graph. The clean recipe searches for dynamic_mx_quant.o rather than objects such as tail_cublas_fp8.o, so these invocations relink the plain objects without RES_CHECK/ENABLE_BINARY_OUTPUT; the resulting ELFs do not read inputs or write result files. Remove the relevant objects or use configuration-specific object paths before every res-check build.
Useful? React with 👍 / 👎.
| if x < 2 ** -9: | ||
| return sign << 7 |
There was a problem hiding this comment.
Preserve representable E4M3 subnormals
The E4M3 minimum positive subnormal is 2^-9, so round-to-nearest should only produce zero below the midpoint 2^-10. For example, 0.0015 is closer to 2^-9 than to zero, but this branch emits 0x00; precision data containing sufficiently small scaled values therefore gets an incorrect golden result. The following subnormal calculation also needs to quantize in 2^-9 increments rather than 2^-6 increments.
Useful? React with 👍 / 👎.
| 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.
Reject individual FP4 output mismatches
Using only aggregate MSE lets incorrect packed FP4 results pass. In the standard 8×64 case, one nibble differing by the minimum nonzero magnitude step of 0.5 contributes only 0.25/512 MSE, far below 0.1, while the computed max_ae is ignored; consequently a kernel can corrupt isolated outputs and still be reported as passing. Require exact codes or enforce an appropriate per-element/MaxAE bound in addition to MSE.
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.
Synchronize 4-PE result writes
When this driver is run with its required multiThreadNum=4, every PE executes main() and reaches these calls after computing only its own row partition. All four threads therefore open and truncate the same files concurrently, and an early PE can read the shared buffers for write() while other PEs are still updating them, yielding partial or interleaved precision results. Synchronize completion and restrict file output to one designated I/O thread, as the existing multi-thread test drivers do.
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 65bc49e848
ℹ️ 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".
| # 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.
Propagate failures from every dynamic-mx build
Make this script preserve a nonzero status from any failed make invocation. As written, its exit status is only that of the final FP4 probe, so if any production or res_check target on lines 12–21 fails while that probe succeeds, the newly wired benchmark/one-level-arch/compile_all.sh treats the operator as successfully compiled and the full-compilation check misses the regression.
Useful? React with 👍 / 👎.
| 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.
Honor kPeNum when routing to the big-block kernel
When kPeNum=4 and the shape selects the big-block route (for example cuBLAS FP8 with BlockSize=128), this call drops kPeNum; the big-block implementation has neither get_thread_idx() nor a partition and loops over every kb. In a four-thread SPMD launch, all four PEs therefore recompute the entire tensor and concurrently store to the same y and scale addresses instead of dividing the work. Partition the big-block loop, restrict it to one PE, or reject this combination; the OCP-FP4 dispatcher has the same omission.
Useful? React with 👍 / 👎.
65bc49e to
cd71c3c
Compare
把 dynamic_mx_quant_tail_ocp_fp8 迁移到运行期动态 shape:M/N 经 tiling 指针
运行期传入,所有切分参数运行期计算;BlockSize/kPeNum 保留模板参(属性)。
对齐 normalization/rms_norm 的动态入口范式。
设计要点:
- physical tile 形状(TileM×列宽)仍编译期锁定(仅由 BlockSize 决定,决定寄存器分配);
Valid 有效尺寸全部下放运行期:tile 声明 ValidRow=-1(DYNAMIC)+ 运行期 ctor 传 vr,
列 valid 保持静态(TEPL B.DIM 立即数约束)。
- global_tensor 用 RowMajor<-1,-1>,构造传运行期 (M,N),stride_t 以 dynamicCol=N
作行 stride → [vr,BlockSize] 多行 strided 列块 load 正确;不能用 global_iterator
(依赖编译期 RowStride)。
- 尾块范式跃迁:full-tile 与尾块共用同一 Valid=-1 类型,仅 ctor 传不同 vr,消除
boxed 编译期特例;PE 分派退化为运行期公式(无需 switch(tid) 编译期展开)。
- 动态 valid 尾块走 physical 列=1 规整路径,天然规避静态版 boxed 尾块的问题22 契约缺陷。
验证(gfrun,codex 基线工具链):多组 shape 全链路
(整除/尾块/纯尾块/行余数/空PE/单PE零回归/奇numKb/大N/BS64),动态 kernel 计算全部正确。
其中 2 个异常经黄金标准对照证明与动态化无关:
- M=513 的少数行 e4m3 边界差异:静态版逐字节完全相同、错同样行(家族既有=问题21,非动态)。
- M=130/BS64 的 scale 写空:PE=1 时 scale 逐字节正确,4-PE SPMD 写终止竞态(问题24 同族 harness)。
新增 TYPE=TAIL_OCP_FP8_DYN;driver 按 rms_norm 范式运行期传 tiling={M,N}。
文档同步(只记状态 + 更新问题,不记 shape 矩阵):
- README 状态总览新增 TAIL_OCP_FP8_DYN 行(动态入口范式 + 天然规避问题22 + 两个既有边界非动态引入)。
- RECORD 问题24 补复现范围:尾轴 fp8 动态版(BS64/M130)亦触发变体签名(output 满、scale_output 写空、
单 PE 正确)→ 现象跨 axis/scaleAlg/dstType。
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: cd71c3cbfa
ℹ️ 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".
| cmp_dir = COMPARE_ROOT / elf_name | ||
| cmp_dir.mkdir(parents=True, exist_ok=True) |
There was a problem hiding this comment.
Remove stale outputs before each precision run
Clear output.bin and scale_output.bin when preparing this directory, including when --skip-gen is used. On a repeated run, if QEMU exits successfully without rewriting either file—for example because a driver ignores a failed writeBinaryFile call—compare_results() reads the previous run's files and can report a pass for a run that produced no results.
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.
Reject isolated FP8 output mismatches
Require exact FP8 codes or enforce a per-element/MaxAE bound in addition to aggregate MSE. For the standard 8×32 case, changing one expected 1.0 value to 1.125 contributes only 0.125²/256 ≈ 0.000061 MSE, so this condition reports pass despite a wrong FP8 output byte; localized kernel corruptions therefore escape the precision check.
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.
Infer FP4 for suffixed driver names
Recognize FP4 anywhere in the driver suffix (or make dtype optional before inference). When the comparator is invoked directly without --dtype for the introduced dynamic_mx_quant_nontail_ocp_fp4_4pe.elf, dtype defaults to FP8 and the basename does not end with fp4, so packed FP4 bytes are decoded as E4M3 values and the reported metrics/status do not validate the kernel's actual output.
Useful? React with 👍 / 👎.
依据 pto-spec d0ce06ad(ADR-0101「Matrix Scale Cell Layouts, HiF4 Scale Words, and CScale」,target 0.58.4)核实 PTO-ISA matmul 的 scale 排布契约: matmul 消费 MX scale 两类载体(Shared 普通 Tile 纯 planar [group_M,G]/[G,N]、 或 CUBE_M32 CellReg 网格),两者均无 parity zip。真实消费方 matmul_shared_lowp.hpp 逐行坐实(gmAScale/gmBScale = 纯 RowMajor global_tensor + 普通 global_iterator, 全链零交织),与 kernel 现有 compact-planar 三方一致。 结论:AscendC 的 DataCopy<DIST_INTLV_B8> / Reg::Interleave 交织是 Ascend 硬件 打包约定,非 PTO-ISA scale 契约。原「问题5:非尾轴 scale 须交织、卡缺 TINTERLEAVE」 定性反转为「规范定义无需交织,问题不存在」。 改动: - gen_dynamic_mx_quant_data.py: compact_scale_bytes 非尾轴分支 parity 交织 → 纯 planar [scaleRows, cols](尾轴本就 planar,未动)。kernel 未改一行。 - RECORD.md 问题5 重写(结论反转 + 实证核实 + 历史 intrinsic 缺口降级存档)。 - README.md / DESIGN.md / 4 个 nontail kernel + 2 个 driver 的过时「MISSING INTERLEAVE / 卡缺 TINTERLEAVE」注释全部更正为「planar = PTO-ISA Shared 契约」。 验证:nontail_cublas_fp8_4pe(Axis512/Post256/BS32,numKb=16,交织本会大幅重排 16 块行)4-PE gfrun → output=pass (MaxAE=0.0117)、scale=pass (MaxAE=0,逐字节精确)。 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: 1b8c745268
ℹ️ 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.
Launch the four-PE precision case with four threads
The registered TAIL_OCP_FP4 check is launched with this single-PE QEMU argument set, but its kernel hardcodes kPeNum = 4 and assigns each PE a separate row segment; with only PE0 running, three quarters of the output and scale remain unwritten. The precision runner therefore cannot exercise this configuration correctly unless it supplies the documented four-thread setting or uses a single-PE kernel variant.
Useful? React with 👍 / 👎.
| uint8_t *y_u8 = reinterpret_cast<uint8_t *>(y); | ||
|
|
||
| using gm_x = global_tensor<InT, RowMajor<M, N>>; | ||
| using gm_y = global_tensor<uint8_t, RowMajor<M, N / 2>>; |
There was a problem hiding this comment.
Restrict the FP4 kernel to packed output types
If a caller overrides the public OutT template with a supported non-FP4 type such as __fp8_e4m3, recip_emax_bits<OutT>() accepts it and the kernel compiles, but this global layout and the stores below still use N / 2 and BlockSize / 2 byte offsets. FP8 output then overlaps blocks and rows instead of occupying one byte per element; add an assertion limiting OutT to the two packed FP4 types or derive the storage layout from its width.
Useful? React with 👍 / 👎.
| types_to_run = args.types or list(CONFIGS.keys()) | ||
| if args.algo: | ||
| types_to_run = [t for t in types_to_run if CONFIGS[t]["algo"] == args.algo] |
There was a problem hiding this comment.
Reject filters that select no configurations
When the requested filters have no intersection—most directly with the accepted --algo DYNAMIC_RANGE, for which CONFIGS contains no entry—this comprehension produces an empty list, the test loop is skipped, and the script prints an empty summary before exiting successfully. A mistyped or currently unsupported experiment can therefore appear to pass without compiling or running anything; reject an empty types_to_run selection with a nonzero status.
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.
Propagate NaNs while reducing cuBLAS golden blocks
When a block contains a NaN after a finite value, Python's order-dependent max can retain the finite value—for example, max([1.0, nan]) is 1.0. The golden generator then emits a normal scale and reciprocal instead of the cuBLAS path's required 0xff scale and NaN reciprocal, so special-value tests depend on the NaN's position and can disagree with the bit-domain reduction used by the large-block kernel; detect any NaN explicitly before reducing the finite magnitudes.
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.
Point the comparator default at the generated compare tree
When the comparator is invoked directly without --cmp-root, this default resolves from .../dynamic_mx_quant/src to benchmark/one-level-arch/test/kernel/compare, while both the runner and generated artifacts use benchmark/one-level-arch/compare. The documented standalone -d/-l usage therefore cannot find its output and golden files unless callers manually override the supposedly optional path; derive the default from the one-level benchmark root instead.
Useful? React with 👍 / 👎.
官方定性 issue #489「小 Post 只有 PE0 段有值」为多 PE 收尾/输出协议 bug (非数值分叉,per-PE 计算正确)。5 个 dmxq 4-PE driver 原本让 4 个 PE 都 writeBinaryFile(整 buffer, O_TRUNC) 互相覆盖,且无屏障 → PE0 抢先写+ exit_group 杀掉还在算的 worker,小 shape 只留 PE0 段、大 shape 碰巧全过。 修复(纯 driver 层,runtime 不动): - 新增 test/common/spmd_res_check.h:SpmdResCheck<kPeNum>,把文件 I/O 收敛到 leader(tid0)。leader_load 输入屏障(worker 等 leader 读完,复用 matmul_shared idiom);leader_should_write 仅 PE0 落盘。 - 5 driver(tail_ocp_fp8/tail_ocp_fp4/tail_cublas_fp8_4pe/nontail_cublas_fp8_4pe/ nontail_ocp_fp4_4pe)套用 helper;Makefile 加 PM/PN/PAXIS/PPOST 宏转发。 不用输出屏障:本模型无硬件 barrier 原语,gfrun 下「leader 忙等 worker」自旋会 hang;改由构造保证——kernel 按 tid 连续切分,leader 恒拿 ceiling 份额,配合 execWidth=1 块级锁步 + leader 多做 I/O → PE0 必最后完成,落盘时 worker 分片已 在共享 buffer。 验证(逐字节):nontail_cublas 4-PE Axis=512/Post=256(4 PE 全活)output=pass MaxAE=0.0117(fp8舍入)、scale=pass MaxAE=0,4 段全填齐;tail_ocp_fp8 4-PE 512/256 同样 pass;5 driver 全编译过。(既有 gfrun 缺陷:Post=64 4-PE res_check hang——原始全写代码亦 hang;Axis=128 decode 断言——均与本修复无关。) 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: 837d2c5594
ℹ️ 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".
| 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.
Serialize dynamic four-PE result dumping
When this driver runs with its default PPE=4, all four PEs execute main() and reach these calls after computing only their own row partition. Each PE therefore opens and truncates the same two files while other PEs may still be updating the shared buffers, producing partial or overwritten precision results. Use the SpmdResCheck<4> leader-only I/O protocol already used by the static four-PE drivers.
Useful? React with 👍 / 👎.
No description provided.