support vpto backend#699

Open

mouliangyu wants to merge 266 commits into

hw-native-sys:mainfrom

mouliangyu:feature-vpto-backend-merge

This pull request is big! We're only showing the most recent 250 commits

Commits on May 25, 2026

Support more syntax in tilelang dsl

Zhendong404
authored and
mouliangyu
committed
Update openspec

Zhendong404
authored and
mouliangyu
committed
tilelang dsl first working version

Zhendong404
authored and
mouliangyu
committed
Update openspec

Zhendong404
authored and
mouliangyu
committed
Support more dsl syntax

Zhendong404
authored and
mouliangyu
committed
Amend

Zhendong404
authored and
mouliangyu
committed
Update docs

Zhendong404
authored and
mouliangyu
committed
docs: add TileOp expand design and demo

qukelin
authored and
mouliangyu
committed
feat(transforms): add Tile→Vector template lowering pipeline

qukelin
authored and
mouliangyu
committed
brough back the lost pass and pto op after rebase
mouliangyu
committed
feat: tile lang dsl

Zhendong404
authored and
mouliangyu
committed
Update DSL user guide

Zhendong404
authored and
mouliangyu
committed
Support more tile attributes

Zhendong404
authored and
mouliangyu
committed
Support more event IDs

Zhendong404
authored and
mouliangyu
committed
Align copy_ubuf_to_ubuf

Zhendong404
authored and
mouliangyu
committed
Support more load/store ops

Zhendong404
authored and
mouliangyu
committed
Support colmajor indexing syntax

Zhendong404
authored and
mouliangyu
committed
Enhance Tile indexing syntax documentation to clarify restrictions on slice forms

Zhendong404
authored and
mouliangyu
committed
Update openspec

Zhendong404
authored and
mouliangyu
committed
Support more predicate ops

Zhendong404
authored and
mouliangyu
committed
Support keyword arguments for vdup and vci operations in TileLang DSL

Zhendong404
authored and
mouliangyu
committed
Support more dma operations

Zhendong404
authored and
mouliangyu
committed
Align DSL interface with VPTO v0.3

Zhendong404
authored and
mouliangyu
committed
Add more memory op support

Zhendong404
authored and
mouliangyu
committed
Fix failed lit case

Zhendong404
authored and
mouliangyu
committed
Dump source location for DSL frontend errors

Zhendong404
authored and
mouliangyu
committed
Fix inline_proc codegen

Zhendong404
authored and
mouliangyu
committed
Support optional RoundMode/SatMode/Part arguments in vcvt op

Zhendong404
authored and
mouliangyu
committed
Fix tile attribute extraction

Zhendong404
authored and
mouliangyu
committed
support expand tload/tstore to tile lib in ptoas

qukelin
authored and
mouliangyu
committed
Fix TileLang ST ptoas invocation

qukelin
authored and
mouliangyu
committed
Fold TileLang dead branches after intrinsic folding

qukelin
authored and
mouliangyu
committed
Fix tilelang ST tadd simulator flow

qukelin
authored and
mouliangyu
committed
fixup ptr normalize
mouliangyu
committed
Refactor tilelang ST framework and document full execution flow

qukelin
authored and
mouliangyu
committed
Add local tilelang ST test common header

qukelin
authored and
mouliangyu
committed
resolve confilct

qukelin
authored and
mouliangyu
committed
update ptoas option

qukelin
authored and
mouliangyu
committed
Support scalar mod operation in DSL

Zhendong404
authored and
mouliangyu
committed
Support custom tile padding value

Zhendong404
authored and
mouliangyu
committed
Fix missing valid_shape/padvalue info when expand tileop

Zhendong404
authored and
mouliangyu
committed
Use pad_value.eval() instead of pad_value.value

Zhendong404
authored and
mouliangyu
committed
Fix semantics of vdup op

Zhendong404
authored and
mouliangyu
committed
Fix vecscope inference for inline_proc

Zhendong404
authored and
mouliangyu
committed
Fix tshrs/tshls scalar type legality check

Zhendong404
authored and
mouliangyu
committed
fix(dsl): extend scalar arithmetic/bitwise lowering and tests (#71)

Zhendong404
authored and
mouliangyu
committed
fix(dsl): align vbitsort and vmrgsort4 with vpto (#67)

Zhendong404
authored and
mouliangyu
committed
[PTOAS] wire MLIR IR printing through ptoas pipelines

Zhendong404
authored and
mouliangyu
committed
Use select_kernel for expand helper selection

qukelin
authored and
mouliangyu
committed
Recover static view strides in ExpandTileOp

qukelin
authored and
mouliangyu
committed
Refine tilelang ST compare flow

qukelin
authored and
mouliangyu
committed
fix(tilelang-dsl): enforce vcvt attrs by type pair and align vcvt lowering (#63)

Zhendong404
authored and
mouliangyu
committed
chore: drop trowargmax template from issue #63 PR

Zhendong404
authored and
mouliangyu
committed
fix(dsl): support module-level literal constants in kernels (#62)

Zhendong404
authored and
mouliangyu
committed
feat(dsl): allow inline_proc capture of module literal globals

Zhendong404
authored and
mouliangyu
committed
Update dsl user guide

Zhendong404
authored and
mouliangyu
committed
tilelang-dsl: support pass as frontend no-op

Zhendong404
authored and
mouliangyu
committed
align(tilelang-dsl): support PartitionTensorView slice binding and docs

Zhendong404
authored and
mouliangyu
committed
udpate tload and tstore.

sundyCoder
authored and
mouliangyu
committed
update tload, tstore and related ST.

sundyCoder
authored and
mouliangyu
committed
Add a skill to resolve DSL issue

Zhendong404
authored and
mouliangyu
committed
fix(tilelang-dsl): emit stable float bit-pattern constants

Zhendong404
authored and
mouliangyu
committed
fix(tilelang-dsl): enforce explicit pointer surface for psts

Zhendong404
authored and
mouliangyu
committed
fix(dsl): handle valid_shape subscript and guard unsupported tuple subscripts (#90)

Zhendong404
authored and
mouliangyu
committed
support pto.set_mov_pad_val
mouliangyu
committed
Add VPTO membar support
mouliangyu
committed
warning the channel vsts
mouliangyu
committed
update docs for channel vsts
mouliangyu
committed
fix(dsl): support explicit set_mov_pad_val for DMA padding

Zhendong404
authored and
mouliangyu
committed
fix: support unsigned dtypes in ExpandTileOp

qukelin
authored and
mouliangyu
committed
fix(tilelang-dsl): keep vbitsort and vmrgsort4 out of inferred vecscope

qukelin
authored and
mouliangyu
committed
feat(tilelang-dsl): add select_kernel diagnostics report

Zhendong404
authored and
mouliangyu
committed
specify dist of vsts/vlds
mouliangyu
committed
fix(dsl): materialize vtrc round mode in tilelang

Zhendong404
authored and
mouliangyu
committed
Add tcvt TileLib support and ST coverage

qukelin
authored and
mouliangyu
committed
Add f16-to-f32 tcvt support

qukelin
authored and
mouliangyu
committed
Add f32-to-f16 tcvt support

qukelin
authored and
mouliangyu
committed
let the dist of store op more flexiable
mouliangyu
committed
fix(dsl): fix vecscope inference

Zhendong404
authored and
mouliangyu
committed
support ctrl reg configure ops
mouliangyu
committed
udpate tload and tstore, and add testcase.

chenjinlin
authored and
mouliangyu
committed
feat(dsl): complete predicate ops and enforce PredicateDist enum

Zhendong404
authored and
mouliangyu
committed
Add vbitcast implementation

Zhendong404
authored and
mouliangyu
committed
Fix TileLang enum support for issue 131

Zhendong404
authored and
mouliangyu
committed
Adding testing for TCI, assuring correctness

OmarZohir
authored and
mouliangyu
committed
test(tilelang_st): add self-hosted runner batch entrypoint

Zhendong404
authored and
mouliangyu
committed
fix(tilelang-dsl): emit signless arith constants for unsigned scalars (#143)

Zhendong404
authored and
mouliangyu
committed
Revert "Adding testing for TCI, assuring correctness"

Zhendong404
authored and
mouliangyu
committed
fix(tilelang-dsl): support i64 vregs in DSL v1 (#151)

Zhendong404
authored and
mouliangyu
committed
bugfix: i8/i16 vcadd will widen the return type
mouliangyu
committed
bugfix: i8/i16 vcadd res type
mouliangyu
committed
bugfix: update i8/i16 vcadd dsl docs
mouliangyu
committed
fix(tilelang-dsl): align f16->i32 vcvt contract (#152)

Zhendong404
authored and
mouliangyu
committed
Support bf16->f16 convert

Zhendong404
authored and
mouliangyu
committed
Support materialize PadValue with eval(dtype) interface

Zhendong404
authored and
mouliangyu
committed
fix(dsl): allow PadValue.eval with static dtype bindings

Zhendong404
authored and
mouliangyu
committed
fix(tilelang-dsl): auto-cast mov pad scalars for issue 170 (#170)

Zhendong404
authored and
mouliangyu
committed
fixup ci
mouliangyu
committed
fix(dsl): support integer string scalar literals (#174)

Zhendong404
authored and
mouliangyu
committed
feat: vpto ci
mouliangyu
committed
feat: ignore local workspace
mouliangyu
committed
feat: support tilelang dsl ci
mouliangyu
committed
remove the fork repo pr protects
mouliangyu
committed
fix(vpto): normalize signed integer vector decls

Zhendong404
authored and
mouliangyu
committed
[WIP] feat: new dma load/store op (#141)

mouliangyu
and
mouliangyu
committed
Support packed vcvt part modes

Zhendong404
authored and
mouliangyu
committed
Support p0/p1/p2/p3 vcvt part mode

Zhendong404
authored and
mouliangyu
committed
Support the whole mem barrier types

Zhendong404
authored and
mouliangyu
committed
chore(ci): fix license headers for pr199

Zhendong404
authored and
mouliangyu
committed
fix(expand-tileop): add cmp_mode attribute handling for TCmpOp and TCmpSOp

Zhendong404
authored and
mouliangyu
committed
feat: support dsl ut ci
mouliangyu
committed
feat(tilelang-dsl): add vscatter surface lowering

Zhendong404
authored and
mouliangyu
committed
fix(dsl): require hex strings for integer bit patterns (#174)

Zhendong404
authored and
mouliangyu
committed
docs: add license headers to issue 174 updates

Zhendong404
authored and
mouliangyu
committed
Add bitcast after arith.constant

Zhendong404
authored and
mouliangyu
committed
fix: handle type compatibility in tryCloneOpLibInlineBridgeOp

Zhendong404
authored and
mouliangyu
committed
fix(dsl): align vexpdif surface with VPTO

Zhendong404
authored and
mouliangyu
committed
fix: avoid false A5 tprelu vec overflow in memory planning

FangRui0
authored and
mouliangyu
committed
fix(dsl): unify type conversion logics in frontend

Zhendong404
authored and
mouliangyu
committed
fix(ptoas): make vpto backend own tile op expansion (#162)

Zhendong404
authored and
mouliangyu
committed
feat: support pto.pbitcast op
mouliangyu
committed
fix(dsl): fix the order mode lowering of pto.vci

Zhendong404
authored and
mouliangyu
committed
feat: support copy_ubuf_to_ubuf / dma_copy lowering
mouliangyu
committed
fix(vpto): align vrelu i32 support (#220)

Zhendong404
authored and
mouliangyu
committed
Fix DSL frontend vecscope auto inference

Zhendong404
authored and
mouliangyu
committed
feat(tilelang-dsl): support constructor calls on static dtype bindings

Zhendong404
authored and
mouliangyu
committed
fix(dsl): fix DSL frontend punpack lowering

Zhendong404
authored and
mouliangyu
committed
bugfix: vbr/cmps with vreg<i8> accepts i16 as scalar operand
mouliangyu
committed
feat: vci support si8/si16/si32/f16/f32
mouliangyu
committed
feat: remove dead code
mouliangyu
committed
feat: dsl supports for trems
mouliangyu
committed
align(tilelang-dsl): require enum dist for vlds/vsts

Zhendong404
authored and
mouliangyu
committed
chore(tileops): add missing license header for tcvt template

Zhendong404
authored and
mouliangyu
committed
Add a mixed Tile/VPTO online softmax kernel

qukelin
authored and
mouliangyu
committed
Refine mixed Tile/VPTO softmax kernels and CI fixes

qukelin
authored and
mouliangyu
committed
fix(pto): avoid false A5 trowarg vec overflow (#558)

Zhendong404
authored and
mouliangyu
committed
Fix tcolarg verifier element width constraints

Zhendong404
authored and
mouliangyu
committed
feat(tileop): Add unary tileop templates (#168)

liggest
authored and
mouliangyu
committed
feat(vpto): add vmatmal op and dependent copy op

FangRui0
authored and
mouliangyu
committed
feature dma remove sid (#243)

mouliangyu
and
mouliangyu
committed
fix(dsl): forward trandom rounds in tile expansion (#245) (#250)

Zhendong404
authored and
mouliangyu
committed
Add OP for TPartAdd & TPartMul (#213)

authored and
mouliangyu
committed
Add OP for TPartMin & TPartMax (#230)

authored and
mouliangyu
committed
fix(vpto): legalize integer cast ops after VPTO lowering (#240)

Zhendong404
authored and
mouliangyu
committed
fix(vpto): restore byte-scalar vdup HIVM ABI

Zhendong404
authored and
mouliangyu
committed
feat: cancel ci on update (#260)

mouliangyu
and
mouliangyu
committed
fix(vpto): normalize scalar memref access before llvm emit (#247)

Zhendong404
authored and
mouliangyu
committed
feat: disable vpto ci
mouliangyu
committed
docs(vpto): add Cluster Programming Model section to Part I

Happybot
authored and
mouliangyu
committed
docs(vpto): add ASCII cluster topology diagram to Cluster Programming Model section

Happybot
authored and
mouliangyu
committed
docs(vpto): fix SC semaphore layout in cluster diagram

Happybot
authored and
mouliangyu
committed
[feat] add rowsum rowmin rowmax rowargmax rowargmin (#255)

authored and
mouliangyu
committed
add tcolmax tcolmin tileops lib implementation (#221)

ChaoyangJi
authored and
mouliangyu
committed
Add scalar ops (#191)

LLMZhangYC
authored and
mouliangyu
committed
feat: improve the func call in dsl. and support vdiv/vmod func
mouliangyu
committed
add texpand/tfillpad/tfillpad_inplace/tfillpad_expand op (#167)

authored and
mouliangyu
committed
bugfix: vprelu should accept mask operand (#266)

mouliangyu
and
mouliangyu
committed
ST tests for TReLU, TLReLU , TPrelu, Tsel and Tselsand operators have been passed. (#212)

authored and
mouliangyu
committed
add trowexpand op (#149)

pbbb205
authored and
mouliangyu
committed
feat: enable vpto sim
mouliangyu
committed
fix(dsl): allow integer vdiv vector types

Zhendong404
authored and
mouliangyu
committed
feat: support vmrgsort4 emit
mouliangyu
committed
feat: support 64bits dma op
mouliangyu
committed
add tcolexpand series op (#169)

authored and
mouliangyu
committed
[Add] tadd tsub tmul tdiv tmax tmin tshl tshr tor tand txor tcmp trem tfmod

zwd060924
authored and
mouliangyu
committed
[Add] tcmp trem tfmod

zwd060924
authored and
mouliangyu
committed
[Delete] tcmp tfmod trem

zwd060924
authored and
mouliangyu
committed
[Delete] tcmp tfmod trem

zwd060924
authored and
mouliangyu
committed
[Fix] pass CI

zwd060924
authored and
mouliangyu
committed
[Fix] delete tcmp tfmod trem in Cmakelists

zwd060924
authored and
mouliangyu
committed
[Fix] CI error in tmin

zwd060924
authored and
mouliangyu
committed
Fix VPTO vcvt and vaxpy mask lowering

Zhendong404
authored and
mouliangyu
committed
Fix installed TileLang resources and disable wheel CI

Zhendong404
authored and
mouliangyu
committed
Fix TileLang soft-math helper lookup in installed layout

Zhendong404
authored and
mouliangyu
committed
bugfix: fixup the mask lack (#291)

mouliangyu
and
mouliangyu
committed
[Fix] TPartMin & TPartMax add mem_bar

caojian5
authored and
mouliangyu
committed
Fix saturation immediate encoding in VPTO LLVM emitter

qukelin
authored and
mouliangyu
committed
Fix lit cases

Zhendong404
authored and
mouliangyu
committed
fix(dsl): allow mixed index and integer scalar binary ops (#246)

Zhendong404
authored and
mouliangyu
committed
feat(vpto): support vecscope auto infer in the vpto backend

Zhendong404
authored and
mouliangyu
committed
Support cross-file inline_proc import

Zhendong404
authored and
mouliangyu
committed
fix(vpto): support merged align state across scf.if (#300) (#308)

Zhendong404
authored and
mouliangyu
committed
docs: refresh Tile Instruction specs and release bundles

qukelin
authored and
mouliangyu
committed
docs: move tensor_view_addr into tile spec

qukelin
authored and
mouliangyu
committed
docs: refresh v0.4 release spec notes

qukelin
authored and
mouliangyu
committed
fix(vpto): force V300 ctrl mode for sat-sensitive vcvt

Zhendong404
authored and
mouliangyu
committed
feat: add cube dma laod/store ops

FangRui0
authored and
mouliangyu
committed
bugfix: fixup mix fatobj
mouliangyu
committed
fix(test/vpto): use pto.acc_store_gm for ACC to GM in cube vpto ST

FangRui0
authored and
mouliangyu
committed
feat: add left/right_load_mx

FangRui0
authored and
mouliangyu
committed
docs: add NZ fractal layout section for cube buffers in Cluster Programming Model

peanutchan
authored and
mouliangyu
committed
docs: add Cube Internal Buffer Layout NZ Fractal Format section

peanutchan
authored and
mouliangyu
committed
[Feat] enhance vcvt behavior for rowsum (#295)

lwwangcgz
authored and
mouliangyu
committed
[Add] tcmp tfmod trem (#321)

authored and
mouliangyu
committed
fix tfillpad_inplace_template (#299)

authored and
mouliangyu
committed
Add trems/tcmps/tfmods (#283)

LLMZhangYC
authored and
mouliangyu
committed
Delete useless code and abbreviate tile_buf type in test cases (#332)

authored and
mouliangyu
committed
add tcolargmax/min tileops lib implementation (#294)

ChaoyangJi
authored and
mouliangyu
committed
Add OP For TSort32 (#258)

authored and
mouliangyu
committed
Fix default render dtype selection for AnyInt

Zhendong404
authored and
mouliangyu
committed
Support scalar/index auto conversion

Zhendong404
authored and
mouliangyu
committed
Add precision_mode attr plumbing for tile math ops

qukelin
authored and
mouliangyu
committed
Fix PrecisionMode Python bindings

qukelin
authored and
mouliangyu
committed
Remove PrecisionMode binding smoke test

qukelin
authored and
mouliangyu
committed
add op: tprelu trandom (#328)

yuqiha
authored and
mouliangyu
committed
feat(tileop): Add more conversion paths for TCVT tileop template (#248)

authored and
mouliangyu
committed
add tilebuf static check

pbbb205
authored and
mouliangyu
committed
cleanup(ptoas): drop duplicate pass registration

Zhendong404
authored and
mouliangyu
committed
fix(dsl): support builtin vector operand matching (#340)

Zhendong404
authored and
mouliangyu
committed
test: add fixpipe cube-cv mixed st case

FangRui0
authored and
mouliangyu
committed
test: add fixpipe cc-to-gm and cc-to-ub st case

FangRui0
authored and
mouliangyu
committed
test: add fixpipe quant ub cv st case

FangRui0
authored and
mouliangyu
committed
feat: add fp_load wrapper op

FangRui0
authored and
mouliangyu
committed
feat: fold fpc into acc_store wrappers

FangRui0
authored and
mouliangyu
committed
feat: sync testcase
mouliangyu
committed
feat: enable smoke test of tilelang dsl ci
mouliangyu
committed
feat(tilelang-dsl): add cube surface and vpto lowering support

Zhendong404
authored and
mouliangyu
committed
Enhance verification for multiple PTO operations in A5 architecture (#303)

authored and
mouliangyu
committed
Implement ptoas VPTO fatobj emission
mouliangyu
committed
Support VPTO section split input
mouliangyu
committed
test: switch validation suites to fatobj
mouliangyu
committed
chore: fix PR386 headers for VPTO helpers
mouliangyu
committed
Remove legacy VPTO target CLI knobs
mouliangyu
committed
Add get_vms4_sr four-result support
mouliangyu
committed
Add HP (HIGH_PRECISION) support for TLog (#344)

authored and
mouliangyu
committed
Organize PTOAS test frameworks
mouliangyu
committed
Fix VPTO lit checks for fatobj workflow
mouliangyu
committed
Move misplaced VPTO lit cases
mouliangyu
committed
Fix stale PTO lit cases
mouliangyu
committed
Add texp high precision implementation (#343)

ChaoyangJi
authored and
mouliangyu
committed
feat: fixpipe design
mouliangyu
committed
Add VPTO MAD semantic ops
mouliangyu
committed
Finalize cube semantic docs and ops
mouliangyu
committed
Rename VPTO MTE semantic ops (#364)

mouliangyu
and
mouliangyu
committed
feat: fixup rebase mistakes
mouliangyu
committed
Fix VPTO cube matmul validation cases

FangRui0
authored and
mouliangyu
committed
feat: support simt
mouliangyu
committed
add tsqrt high precision implementation (#367)

ChaoyangJi
authored and
mouliangyu
committed
Add OP for TMrgSort (#262)

authored and
mouliangyu
committed
Add HP (HIGH_PRECISION) support for TDiv (#348)

authored and
mouliangyu
committed
Support index-type operands in tile op template expansion

TelGome
authored and
mouliangyu
committed
support dynamic set/wait flag

LLMZhangYC
authored and
mouliangyu
committed
bugfix: fixup llvm build in ci
mouliangyu
committed
Add lit for dynamic flagId

LLMZhangYC
authored and
mouliangyu
committed
vpto-rebase: fixup ptoas outputFile
mouliangyu
committed
vpto-rebase: unify set/wait flag asm format
mouliangyu
committed
vpto-rebase: fixup memref addrspace
mouliangyu
committed
vpto-rebase: fixup license
mouliangyu
committed

Commits on May 26, 2026

feat: remove dead pass and docs
mouliangyu
committed
feat: rename pto.aicore to pto.kernel
mouliangyu
committed