Skip to content

fix(vpto): support signless i8*i8->i32 mad lowering (https://github.com/mouliangyu/PTOAS/issues/381)#389

Closed
TelGome wants to merge 1 commit into
mouliangyu:feature-vpto-backendfrom
TelGome:issue_381
Closed

fix(vpto): support signless i8*i8->i32 mad lowering (https://github.com/mouliangyu/PTOAS/issues/381)#389
TelGome wants to merge 1 commit into
mouliangyu:feature-vpto-backendfrom
TelGome:issue_381

Conversation

@TelGome
Copy link
Copy Markdown
Collaborator

@TelGome TelGome commented May 21, 2026

Fix issue: #381

Problem:
issue #381 的核心问题出在 VPTO LLVM lowering 的 MAD callee 分发上。tmatmul 展开后会生成 pto.mad_raw,其中 int8 case 的实际类型是 signless i8 * i8 -> i32,不是 si8 * si8 -> si32,可以在 test/tilelang_st/npu/a5/src/st/testcase/tmatmul/tmatmul.pto:106 对应的 matmul case 里看到。但 lib/PTO/Transforms/VPTOLLVMEmitter.cpp:235 原来的 MAD 类型分发只把 signed integer 当成 s4/s8/s32 处理:

  • getMadRhsFragment 只识别 signed i4/i8
  • getMadDstFragment 只识别 signed i32
  • buildMadTypedCalleeName 里也没有 s8*s8->s32 的 callee 分支

结果是 i8/i8/i32 这组类型拿不到合法的 intrinsic 名称,pto.mad_raw 在 LowerVPTOOpsPass 中无法完成 lowering,最终以 illegal op 形式报错:failed to legalize operation 'pto.mad_raw' that was explicitly marked illegal

Fix:
修复思路是把 PTO IR 中的 signless integer 语义和 MAD lowering 对齐。PTO IR 默认大量使用 signless integer,因此这里不应该只接受 si8/si32,而应该把 i8/si8 统一映射为 s8,把 i32/si32 统一映射为 s32。

具体改动在 lib/PTO/Transforms/VPTOLLVMEmitter.cpp:229:

  • 新增 isSignedOrSignlessInteger(IntegerType, width) helper
  • getMadRhsFragment 对 i4/i8 和 si4/si8 统一返回 s4/s8
  • getMadDstFragment 对 i32 和 si32 统一返回 s32
  • 在 buildMadTypedCalleeName 中补充 s8*s8->s32 到 llvm.hivm.MAD.s82s32.c310 的分发

@Zhendong404
Copy link
Copy Markdown
Collaborator

Included in PR #413

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants