trending/tilelang
1
0
Fork 0
mirror of https://github.com/tile-ai/tilelang.git synced 2026-10-02 06:34:36 +08:00
Code Issues Packages Projects Releases Wiki Activity
Files
main
tilelang/cmake/BackendOptions.cmake
T
+17
Chaofan LinLeiWang1999matt-tianxuruifansilentCoder-devDenver JinYongqi ZhuoMeng YuhaoDenver JinmengyuhaoLei WangDenver Jincjcaojian5AndersSepcntКаримxuebozhang525-altpenguin_wwyAlfredSoham Pandaandyluo7Xiangwen Wang小云xy200303Ziming Wangedragain
 e5a02f9aab [Public Release 9/30] Introduce Ascend 950 backend (#3308) 
* [Ascend] Refresh comparison snapshot at upstream #3174

Comparison-only snapshot of Ascend at 925bfbb9aff0fb57a4ed9fecc3591b8dcac21c8f.
Use upstream #3174 (62bba8d20d) as the logical synchronization point.

Preserve the complete tree of the previous comparison head d0d2f40be354cc0a4d6ed037b34278c8746fa031,
including submodule references. No source changes or additional integration
resolutions are introduced. This snapshot is for upstream comparison only.

* [Ascend] Keep the CUDA default facade and reach Ascend through its own dialect

Upstream #3186 removes the shared `Kernel`, so `tilelang.language` can no longer
adapt to the host at runtime (the removed `check_ascend_availability()` probe was
the only thing making one import work for both). The fork has to pick a dialect
for the default facade, and it now keeps upstream's choice: `tilelang.language`
is the CUDA dialect, and Ascend is reached explicitly, like every other backend.

That is the option with no upstream divergence: the 136 fork-owned Ascend files
(examples/ascend, testing/ascend, tilelang/ascend) take the explicit import,
while upstream's files stay untouched. Previously the fork had converted six
upstream files to `tilelang.cuda.language` to work around the Ascend facade;
those are reverted to upstream's `import tilelang.language as T`, and the
materialize-launch test added by #3186 passes unmodified (19 tests).

Ascend coverage for the launch refactor moves to a fork-owned file,
`testing/ascend/transform/test_ascend_materialize_kernel_launch.py`, which pins
the decoupled lowering (grid -> thread_extent, placeholders dropped), the
`cthread` launch dimension, and the backward-compatible default of
`lower_grid_binding`.

Effect on testing/python/ (H100, CUDA 13.1): 155 failed / 3347 passed before, 17
failed / 3557 passed after. The 138 removed failures were all "module
'tilelang.language' has no attribute ..." and "Kernel() got an unexpected
keyword argument 'threads'" from the dialect swap. 6 of the remaining 17 fail
identically on upstream main (missing nvidia-cutlass-dsl); the other 11 are the
pre-existing shared-code deltas from the reducer and thread-storage-sync work,
unrelated to the launch refactor.

* [Ascend] Point the remaining explicit dialect imports at the Ascend package

Follow-up to the facade change: files that pulled individual dialect symbols
with `from tilelang.language import <name>` still went through the CUDA facade.

- `from tilelang.language import simd as S` (6 files) -> the Ascend `simd`
  module, which is an Ascend extension.
- `from tilelang.language import kernel as K` (test_simtvf_thread_binding) ->
  the Ascend kernel module. It has to be this one: the SimtVF-aware
  get_thread_binding/get_thread_extent accessors now live there, so importing
  the shared module returned the launch placeholder `tx` instead of the
  SimtVF var `simtvf_tx`.
- test_simtvf_frame_type asserted SimtVFFrame is importable from
  `tilelang.language`; it is an Ascend frame, so it now imports it from the
  Ascend dialect.

Verified the three SimtVF shapes these cover still lower through the new
launch materialization (warp-reduce style, vf var scope, and the
simtvf_tx/ty/tz naming contract); none of them trips the new
"body references thread index" diagnostic, so it is not a false positive on
legitimate T.SimtVF code.

* [Ascend] Restore CANN 9.2 codegen compatibility lost by the snapshot base

The snapshot's Ascend tree is asc@925bfbb9, which predates the two commits
that made generated kernels compile on CANN 9.2 / bisheng:

  - 1037a5fc "use the C API exclusively" (#399), which also moved the block
    index off the CUDA builtin ("reserve block_idx", "modify block_idx due to
    cann")
  - bbcf1e47 "keep assert failure handling out of line" (#438)

CANN 9.2 rejects the `blockIdx.x` builtin outside a SIMT context, and says so
in c_api/utils/sys_var.h: "use block_idx for pure Vector, pure Cube, and
Mix(1, 1). For Mix(1, 2), please use block_idx on the Cube core and
block_idx * asc_get_sub_block_num() + asc_get_sub_block_id() on Vector cores."
Every mixed kernel is a Mix(1, 2) launch, so the AIC body could not compile:
"can only use thread or block infos in simt function or simt callee".

Three places emitted or consumed `blockIdx.x` where CANN now requires the C API
`block_idx`:

- src/ascend/codegen/codegen_ascend.cc: map the blockIdx.x thread var to
  `block_idx`, reserve that name so the name supply cannot hand it to another
  var, keep skipping the block index when capturing SIMT_VF helper arguments,
  and include `block_idx` in the int32 cast-back list.
- src/tl_templates/ascend/debug.h: the debug/assert helpers are emitted into
  both a `__aicore__` and a `__simt_callee__` namespace. The former is not a
  SIMT function, so `TL_DEFINE_ASCEND_DEBUG` now takes the block index as a
  parameter -- `block_idx` for the aicore namespace, `blockIdx.x` for the simt
  one -- instead of hardcoding `blockIdx.x` in the macro body.
- tilelang/profiler/bench.py: get_msprof_cache_flush_kernel is Ascend-only
  (target="ascend", T.SimtVF, l2_cache_ctrl="WTS_FV") and was reaching those
  through the default language facade. Since the facade is the CUDA dialect
  again, it imports tilelang.ascend.language instead. This was the one
  remaining shared file the facade change had missed.

testing/ascend/ on 8x Ascend950DT (CANN 9.2.0): 812 passed / 255 failed before,
857 passed / 210 failed after, against 804/250 for the pre-refactor baseline.
The 45 bisheng device-compilation failures are gone; the only failures this
branch still adds over the baseline are the 5 in
test_tilelang_ascend_l1_l0_alignment_guards.py, where CUDA's GemmMMA.infer_layout
now raises before the Ascend L1->L0 alignment guard gets to run.

examples/ascend/example_gemm.py: all 7 sections compile and PASS, with max_diff
matching the reference asc tree (fp8 2.14e-04 / 2.44e-04 run to run).

* feat(ascend): Use the C API exclusively to eliminate compilation overhead caused by mixed API usage. (#399)

* Use capi to eliminate the loss of compilation time caused by mixed api.

* reserve block_idx

* modify block_idx due to cann

* refactor(ascend): finish C API lowering and simplify SIMD wrappers

* fix(ascend): return benchmark latency for blockscaled GEMM

* refactor(ascend): trim C API migration scope

Restore the device assert tests, Bisheng architecture tests, and CUDA debug helper to the merge base. Remove the unused Python GM/UB DMA wrappers and update the remaining pad-value documentation.

* test(ascend): expect C API locks in counter-channel spill

* refactor(ascend): share SIMD call emission and precision selection

Use one printer for SIMD expression operands and trailing control enums, and share SFU wrapper selection between zeroing and merging calls. Preserve argument validation, print order, and the existing float32 precision fallbacks.

* fix(ci): inherit runner Torch stack in Ascend workflows

* fix(ci): bootstrap pre-commit without pipx

* Revert "fix(ci): bootstrap pre-commit without pipx"

This reverts commit 04c742d292eec79863b8400e915c81e3483b20f7.

* fix(ci): install pipx before lint

* fix(ci): install pipx with configured Python

* fix(ci): avoid uninstalling shared pipx dependencies

---------

Co-authored-by: xuruifan <xuruifan@deepseek.com>
Co-authored-by: silentCoder-dev <silentcoder@foxmail.com>
Co-authored-by: Denver Jin <denverjin@deepseek.com>

* Update ci.yml (#437)

* [Ascend] Fix auto-scheduler memory detector ODR violation (#435)

Co-authored-by: silentCoder-dev <silentcoder@foxmail.com>

* fix(ascend): analyze multi-buffer eligibility on scheduled IR (#422)

* fix(ascend): analyze multi-buffer eligibility on scheduled IR

* fix(ascend): align guarded broadcast fill analysis

* fix(ascend): rewrite multi-buffer scheduling guards

* fix(ascend): clarify multi-buffer scheduling constraints

* fix(ascend): keep assert failure handling out of line (#438)

Bisheng can fail to select a stacksave instruction when an inlined
device_assert_with_msg is followed by an Ascend sync wait. The failure
reproduces with a five-character message on CANN 9.2.0 targeting dav-3510.

Move printf and assert(false) into a noinline failure helper while retaining
the conditional check in the inline wrapper. Preserve the diagnostic text,
BlockIdx, and failure action for both AICore and SIMT callers.

Previously completed local validation: the standalone reproducer fails
before the change and compiles afterward with Bisheng at -O2. Representative
MoE and lossless quantization generated kernels also compile with the fixed
header, and pre-commit passes for the header. No runtime or performance
measurements were performed.

* fix(ascend): correct latency and initiation interval estimates (#436)

* fix(ascend): correct latency and initiation interval estimates

* fix(ascend): preserve HF32 mode across guarded tasks

* fix(ascend): preserve physical row widths for dynamic MTE copies

* feat(ascend): add assume conflict hints (#339)

* feat(ascend): add conflict hints

* fix(ascend): preserve dependency distances

* docs(ascend): fix conflict hint operand comment

* docs(ascend): align conflict hint entry types

* test(ascend): run UB fills in vector functions

* test(ascend): update assume conflict codegen assertions

---------

Co-authored-by: mengyuhao <mengyuhao@deepseek.com>

* [Ascend] Use native merging for validated SIMD overloads (#440)

* [Ascend] Restore native SIMD vadds merging

* fix(ascend): use native merging for validated SIMD overloads

Extend the native vadds merging path to the verified overloads of 30 more
SIMD operation families. Forcing zeroing plus select through C API wrappers
can introduce predicate spills in masked updates, as seen in SwiGLU.

Gate native dispatch by operation and element type, and map vector vdupv
to the corresponding CCE vdup overload. Keep scalar BF16 vdup on software
merging because CANN 9.2 merges through an uninitialized temporary; select
that fallback from the destination vector type. Preserve precision-specific
SFU algorithms and the existing dispatch for unvalidated overloads.

Add codegen and NPU regression coverage for native calls, dtype fallbacks,
precision wrappers, broadcast/reduction layouts, and preservation of inactive
destination lanes. Document the merging behavior and exceptions.

Validation: native build with CUDA disabled, scoped format.sh, and diff checks
passed. On Ascend 950DT with CANN 9.2, 80 codegen/legalization/precision checks
and 44 NPU operation/type cases passed; the latter cover 132 empty, full, and
sparse mask checks. Broader per-operation performance is not inferred from
the earlier SwiGLU vadds measurements.

---------

Co-authored-by: silentCoder-dev <silentcoder@foxmail.com>

* [Refactor] Add backend namespaces for auto_schedule (#441)

* feat(ascend): support effective K on padded L0 tiles (#439)

* feat(ascend): support compact L0 GEMM regions

Guard transposed L1->L0 copies at the region level: the transposed load
issues whole 16-row groups on the L0 M/N axis, so a compact source region
whose MN extent is not 16-aligned emitted a misaligned MTE instruction and
trapped the device at runtime.  Also require the L0 destination K region to
cover the whole-group rounding of the source K region, which otherwise
writes past a compact L0 allocation.

Make the tl.tileop.gemm M/N/K metadata follow static operation region
extents for Ascend L0 GEMMs, so the latency/cost model sees the real MAD
geometry instead of the padded buffer shapes and stops overestimating
region-restricted compressions.

* [Ascend][GEMM] Use effective K with padded L0 storage

* fix(ascend): validate transpose alignment for nonempty copies

* fix(ascend): keep effective L0 GEMM regions in backend

* test(ascend): trim padded L0 GEMM regressions

* revert(ascend): restore blockscaled GEMM behavior

* [Ascend] Point the cherry-picked Ascend tests at the Ascend dialect

The asc commits just picked wrote four new test files against the Ascend-only
facade (import tilelang.language as T, then T.simd / T.Kernel). This branch keeps
upstream's default facade (CUDA), so they have to name the dialect explicitly,
like the rest of the fork-owned Ascend files.

Fixes 114 failures: test_tilelang_ascend_simdvf_merging (101, "module
'tilelang.language' has no attribute 'simd'"), test_ascend_assume_conflict (9),
test_tilelang_ascend_l0_gemm_effective_k (2), and two in
test_ascend_annotate_multi_buffer_eligible that #422 rewrote back onto the
facade.

testing/ascend/ on 8x Ascend950DT (CANN 9.2.0): 921 passed / 304 failed before
this commit, 1035 passed / 190 failed after. Against the pre-refactor baseline
(804/250) the branch now has 231 more passing tests and no failure that the
baseline does not also have.

* [Ascend] Require a usable CUDA device before auto-detecting the CUDA target

Target auto-detection returns the first registered backend that reports itself
available, and _detect_cuda_target() only asked check_cuda_availability(), which
just looks for a CUDA toolkit/nvcc path. A host that merely has the CUDA toolkit
installed was therefore auto-detected as CUDA even with no CUDA device at all --
which is exactly this NPU box, where /usr/local/cuda exists next to the NPU
stack. Every Ascend example and test that omitted target= then lowered through
tilelang/cuda/pipeline.py and died, e.g.:

  AttributeError: module 'tilelang.cuda._ffi_api' has no attribute
  'AnnotateDeviceBoundTmaCopies'
  (examples/ascend/flash_attention/test_gqa.py)

The asc branch happened to be immune because its `tilelang.language` facade
imported tilelang.ascend.target early, registering the Ascend detector ahead of
CUDA's. That is import-order luck: the facade here is the CUDA dialect, so CUDA
registers first and wins. Fix the detector instead of the order -- the same
predicate _detect_torch_cuda_arch() already uses -- so auto-detection no longer
depends on which dialect a facade happened to pull in. An explicit
target="cuda" still resolves everywhere.

testing/ascend/ on 8x Ascend950DT (CANN 9.2.0): 1035 passed / 190 failed before,
1225 passed / 0 failed after. examples/ascend/flash_attention/test_gqa.py and
test_mha.py both compile and pass.

* [Ascend] Resolve nested thread scopes in the shared thread accessors

T.rng_init() raised "The thread extent is not known at trace time" for Ascend
(examples/ascend/test_simtvf_rng.py, both the ascend and pto targets), even
though the call sits inside a `with T.SimtVF(threads=...)` block.

tilelang/cuda/language/random.py binds the backend-neutral surface:

    import tilelang.language.common as T
    ...
    ex = T.get_thread_extent()

so it reaches tilelang.language.kernel.get_thread_extent by name. The earlier
refactor moved the SimtVF-aware resolution out of those shared accessors and into
tilelang/ascend/language/kernel.py; a dialect-local override only shadows the
dialect namespace and cannot intercept T.get_thread_extent() there. The call
therefore fell through to the launch frame, and Ascend's T.Kernel declares no
threads= (the NPU launch has no threadIdx at kernel scope), so it raised.

Move the mechanism back to where the accessors live. The stack itself is
backend-neutral -- "the innermost nested thread scope" -- and only the frame that
pushes onto it is Ascend-owned (T.SimtVF via SimtVFContext). The Ascend dialect
module now re-exports the accessors instead of redefining them, so
tilelang.ascend.language.kernel stays a complete view of the thread API.

This is the same class of miss as tilelang/profiler/bench.py earlier: fork-owned
behaviour that has to be reachable from backend-neutral callers cannot live in a
dialect namespace alone.

Verified on 8x Ascend950DT (CANN 9.2.0):
  testing/ascend/  1225 passed, 0 failed
  examples/ascend/  186 passed, 0 failed
plus the SimtVF thread-binding contract tests (20 passed) and
examples/ascend/test_simtvf_rng.py (2 passed).

* [Ascend] Remove the PTO and VMI content

PTO was a second Ascend codegen backend (a "pto" key on the ascend target
selecting PTODSL source plus the ptoas toolchain) and VMI was its virtual
vector intrinsic surface (T.vmi.*). Both are dropped from this branch, leaving
the AscendC path as the only Ascend codegen.

111 files changed, +294/-12810, 21 files deleted.

Deleted outright:
- src/ascend/codegen/{codegen_pto.cc,codegen_pto.h,rt_mod_pto.cc}
- tilelang/contrib/ptodsl/ (the PTO DSL: gemm, rng, simt, dcache_bypass)
- tilelang/ascend/language/vmi.py
- 11 PTO-only tests/examples, plus example_rmsnorm_persistent_simtvf.py and its
  test, which existed to exercise PTO's persistent-fragment lowering

Core removals:
- target.py: _make_pto_target, target_is_pto, normalize_pto_target and the "pto"
  normalizer registration. target_is_plain_ascend is gone too -- with PTO removed
  it was identical to target_is_ascend, so its two call sites use that directly.
- backend.py / codegen.py / execution_backend.py: the PTO BackendModule and its
  DeviceCodegen and cython-only ExecutionBackendSpec.
- wrapper.py: the whole PTO host-wrapper family (_PTOKernelDescriptor,
  _PTOKernelCallSite, _PTOHostKernelCall, _PTOHostCallCollector,
  TLPTOSourceWrapper, ~560 lines) and the TLWrapper dispatch arm.
- libgen.py: the PTODSL->ptoas compile path, update_pto_kernels and the PTO temp
  dir helpers (~230 lines).
- utils.py/tvm_ffi.py: is_pto_target and its branches; the tvm_ffi param-shape
  path and the two "plain ascend" predicates collapse to is_ascend_target.
- engine/param.py: the PTO-only packed storage ABI. storage_packing_factor and
  _is_pto_target go away (it always returned 1 without PTO); storage_shape stays
  as the named hook but is now the logical shape unchanged.
- src/op/builtin.{h,cc}: the whole tl.vmi.* op family -- the three registration
  macros and 52 op registrations -- plus codegen's IsVmiOp rejection path and the
  tl.vmi.* clauses in ascend_pipe.h's pipe classification.
- src/transform/make_packed_api.cc: the is_pto key test.
- CI: the PTOAS wheel install and both `-m pto` test steps.

Tests and examples (66 files): every pto parametrization dropped while keeping
its ascend/asc variant, PTO-only helpers and imports removed, and 10 examples'
`--target` choices reduced to ["ascend"].

Verified on 8x Ascend950DT (CANN 9.2.0):
  testing/ascend/   1010 passed, 0 failed   (was 1225; ~215 PTO-variant tests gone)
  examples/ascend/    96 passed, 0 failed   (was 186; the rest were pto variants)
  collection: 1118 tests, 0 errors
and the built library no longer contains any "tl.vmi." or
"target.build.tilelang_pto" string.

* [Backend] Move AllReduce backend policy out of the shared reduce lowerers

src/backend/common/op/reduce.h is backend-neutral, but it had grown two
TargetIsAscend branches:

  - CheckAllReduceWidth skipped its power-of-two requirement on Ascend
  - NeedsWorkspace restated AscendAllReduce's constexpr dispatch to decide
    whether the lowering needs a shared-memory workspace

Both are properties of the backend's all-reduce algorithm, not of the IR, so
they belong to the backend `Impl` -- the same place the neighbouring policy
already lives (Impl::SupportsFp16Bf16NanReduce, Impl::GetPreferredVectorizedSize,
Impl::MakeBatchAllReduce). The shared lowerers now ask the Impl and no longer
know that Ascend exists.

- CheckAllReduceWidth(reducing_threads, scale, op_name,
  requires_power_of_two_width = true): keeps the three universal checks
  (positive threads, positive scale, scale divides threads) and gates the
  XOR-butterfly power-of-two check on the flag. The Target parameter is gone.
- NeedsWorkspace is deleted from common; the two shared call sites use
  Impl::AllReduceNeedsWorkspace(...).

Each backend Impl answers both:
- cuda/rocm Reduce and FinalizeReducer: power-of-two required, workspace only
  when reducing_threads > 32.
- ascend Reduce forwards to src/ascend/op/ascend_allreduce_policy.h, which
  states the two answers once for the Ascend backend. The standalone Ascend
  finalize lowerer (which is not a CRTP Impl) uses that header directly.

The workspace predicate keeps its original semantics; it now lives next to
tl::AscendAllReduce, whose dispatch it mirrors, instead of being restated in
shared code.

Verified on 8x Ascend950DT (CANN 9.2.0): testing/ascend/ 1010 passed, 0 failed,
including the non-power-of-two reduce widths (6,5) and (7,3) that motivated the
opt-out. src/backend/common/ is now free of TargetIsAscend.

* [Backend] Keep the AllReduce XOR-butterfly rule in the algorithm, not in the backend

Follow-up to 9a928180, which answered the shared lowerer's questions with one
per-assertion backend predicate (Impl::AllReduceWidthRequiresPowerOfTwo). That
named a check rather than a capability, so every further check would have added
another such hook. Reshape it:

- CheckAllReduceWidth(reducing_threads, scale, op_name) keeps only the three
  checks that hold for every all-reduce (positive threads, positive scale, scale
  divides threads). Its signature and body are now identical to upstream's.
- The power-of-two requirement moves to CheckXorButterflyWidth(reducing_threads,
  scale), named after the algorithm that actually imposes it. It is called from
  the backends that emit an XOR-butterfly intrinsic -- cuda/rocm in
  MakeBatchAllReduce and MakeScalarAllReduce -- and never from shared code.
  Ascend does not call it: AscendAllReduce falls back to a shared-memory tree
  reduction (ub_reduce) for arbitrary widths.
- Impl::AllReduceWidthRequiresPowerOfTwo is gone from all five Impls, and the
  Make{Batch,Scalar}AllReduce signatures are untouched.

The remaining backend hook is Impl::AllReduceNeedsWorkspace, which asks a real
capability question ("does your lowering need a scratch buffer") rather than
restating an assertion.

The rule's diagnostic loses its "tl.reduce:" / "tl.finalize_reducer:" prefix
because the emit helpers do not know which tile op drove them. Threading that
label through would have been the only alternative, and it would have put a
caller-supplied string into the intrinsic-naming helpers; the message names the
algorithm instead, which is the actionable part.

src/backend/common/ is free of TargetIsAscend, and this is now true of the
AllReduce path without any target-shaped hook.

Verified on 8x Ascend950DT (CANN 9.2.0): testing/ascend/ 1010 passed, 0 failed;
examples/ascend/ 96 passed, 0 failed; the non-power-of-two reduce widths (6,5)
and (7,3) still lower.

* [Layout] Rename MapRegion to MapRegionBounds

* [Ascend] Inline AllReduce workspace policy

* [Ascend] Drop the vestigial CPU codegen hooks

src/cpu/codegen/codegen_c.{h,cc} carried two hunks that the snapshot inherited
from the pre-rebase Ascend commit f64773d3 ("feat(ascend): Add __ubuf__ shared
memory support for SimtVF"). Neither has a consumer today:

- Init() took a prelude_include parameter whose default was exactly the string
  it replaced ("#include <tl_templates/cpp/common.h>\n"), so every 5-argument
  call produced byte-identical output. No caller anywhere passed a sixth
  argument -- not at this HEAD, and not at f64773d3 either. The parameter was
  dead on arrival.
- PrintType, VisitExpr_(BroadcastNode) and VisitStmt_(AllocBufferNode) were
  relaxed from final to override so that CodeGenTileLangAscend could subclass
  CodeGenTileLangC. CodeGenTileLangAscend now derives from tvm::codegen::CodeGenC
  directly and has its own Init(bool) that writes the Ascend preludes
  (codegen_ascend.cc:327), so nothing in the tree derives from
  CodeGenTileLangC at all.

The one caller of CodeGenTileLangC::Init is src/cpu/codegen/rt_mod_c.cc, which
uses the default. (codegen_c_host.cc calls CodeGenCHost::Init, a different
class.) src/backend/common/codegen/ and src/cuda/codegen/ are otherwise
unchanged from upstream; this was the only codegen file the branch touched
outside src/ascend/.

Reverting both files to upstream compiles cleanly (2 TUs rebuilt, exit 0) and
keeps testing/python/cpu/ at 88 passed, 0 failed.

* [Ascend] Restore upstream VerifyParallelLoop

Restore src/transform/verify_parallel_loop.cc byte-for-byte from upstream-main at 85fd8fc2d3. Leave all other Ascend changes untouched.

* [Ascend] Drop the re-added common target_utils include

src/cuda/op/copy.cc and src/metal/op/copy.cc each carried one line the
snapshot should not have:

  #include "backend/common/target_utils.h"

Upstream PR #2361 (2913ad3f, "Cleanup Metal Codegen, split AsyncCopy lowering
and common target utils") deleted exactly that line from both files and
replaced it with the backend-specific cuda/target_utils.h and
metal/target_utils.h. The old asc history replayed a pre-#2361 diff on top of
the split, so the deleted line and its replacement ended up side by side.

Neither file uses anything from it. backend/common/target_utils.h declares only
TargetHasAsyncCopy, which appears zero times in cuda/op/copy.cc. Every helper
the file calls -- TargetIsCuda, TargetIsCuTeDSL, TargetIsHopper, TargetIsSm100,
TargetHasBulkCopy, TargetHasStmatrix, TargetHasTmem -- is declared in
cuda/target_utils.h, included on the next line. metal/op/copy.cc is the same
story with TargetIsMetal and TargetMetalGetWarpSize, both from
metal/target_utils.h.

Checked that this is not a header-chain repair: none of the changed headers in
copy.cc's include chain (op/copy.h, op/utils.h,
transform/common/loop_fusion_utils.h) adds or removes an include, and the six
*/target_utils.h headers declare no colliding names.

With the line gone both files are byte-identical to upstream-main, so the
argument is identity with content upstream CI compiles rather than a test
result. Local rebuild recompiled only src/metal/op/copy.cc (exit 0).
src/cuda/op/copy.cc is not compiled on the Ascend box (USE_CUDA=OFF).

* [Backend] Route AllReduce width checks through the backend Impl

4d40c190 moved the XOR-butterfly power-of-two rule out of CheckAllReduceWidth and
into the backends' Make{Batch,Scalar}AllReduce emit helpers. That put the rule
after ResolveAllReduceThreadRange, which reports a different problem first: for
a reducing width that is both a non-power-of-two and warp-misaligned (48 with
warp_size 32), the "partial scalar AllReduce requires a warp-aligned thread
range" ICHECK fires and the power-of-two diagnostic never runs. Three CUDA tests
regressed on the message, and 96-thread cases still passed because 96 % 32 == 0,
so the ordering hole only shows up at the intersection of the two rules.

The emit helpers cannot be reordered instead: MakeScalarAllReduce takes the
thread offset and extent that ResolveAllReduceThreadRange produces. The check has
to move earlier, which means the shared lowerer must know which width
preconditions its backend imposes.

So the shared call sites now ask the backend:

    Impl::CheckAllReduceWidth(reducing_threads, scale, op_name, target);

replacing the three direct calls to the free function of the same name, at the
same points in the lowering (both tl.reduce paths and the tl.finalize_reducer
path). The rule text still lives in exactly one place, backend/common/op/reduce.h
-- CheckAllReduceWidth for the checks every all-reduce lowering owes, and
CheckXorButterflyWidth for the extra requirement of the XOR-butterfly shuffle
all-reduce. Each Impl states which apply to it:

  - cuda/rocm call both, since their all-reduce has no arbitrary-width fallback;
  - ascend calls only the universal one. AscendAllReduce does have XOR-butterfly
    (asc_shfl_xor) and hardware-reduce paths, but run() gates both of them on
    is_pow2(threads), so they can only ever see a power-of-two logical width.
    Every other width goes to the ub_reduce fallback, which folds threads/scale
    slots along a scale-stride loop and needs no power of two. The rule is
    therefore vacuous on Ascend, not absent from it.

src/ascend/op/finalize_reducer.cc implements LowerFinalizeReducer directly
rather than through FinalizeReducerLowerer, so it keeps calling the universal
rule and now says why. No target branching appears in shared code: it only
forwards to Impl, which is the same shape as the existing
Impl::AllReduceNeedsWorkspace hook.

The eight CheckXorButterflyWidth calls scattered through the four emit helpers
are gone, so Make{Batch,Scalar}AllReduce name intrinsics and nothing else. A
future width constraint needs no new hook: it becomes one more line in each
backend's CheckAllReduceWidth.

Verified on 8x Ascend950DT (CANN 9.2.0): testing/ascend/ 1010 passed, 12
skipped, 0 failed. On 8x H100 (CUDA 13.1), the three regressed tests plus the
four that share their parametrization: 7 passed. Full testing/python/ is 23
failed, 3551 passed -- all 15 failures of the pre-change baseline are still
there, and the 8 extras are the known flaky group (sm75 x4, tma_dsmem x2,
func_attrs x2), so nothing new regressed.

src/rocm/** is compiled by no build available here. Both ROCm TUs were checked
with c++ -fsyntax-only using the compile flags of a sibling TU (exit 0), after
confirming with a negative control that the check catches a deliberate error.

* [Ascend] Register the Ascend intrinsic Ops from the Ascend backend

src/op/builtin.{h,cc} carried 111 Ascend intrinsic Ops -- 82 tl.simd.* and 29
tl.ascend_* -- plus the tl.enable_auto_schedule pass option, with a comment
explaining that "the SIMD ops stay here since the Ascend backend keeps its
builtins in the common file". Upstream #2855 had already given CUDA, Metal and
ROCm their own op/builtin.{h,cc}; Ascend was the one backend still squatting on
the neutral header, which is why a backend-neutral file listed asc_shfl_xor
masks and Cube MAD instructions.

Move them to src/ascend/op/builtin.{h,cc}, mirroring cuda/op/builtin.{h,cc}:
declarations and the pass-config key in the header, TIR_DEFINE_TL_BUILTIN
registrations in the .cc. The T.simd.* ops need a macro of their own because
they register as "tl.simd.<name>" while their accessors are simd_<name>; that
macro is now local to ascend/op/builtin.cc instead of sitting in the neutral
one.

The 14 Ascend files that name these ops now include "ascend/op/builtin.h"
instead of "op/builtin.h". src/ascend/CMakeLists.txt needs no change: the
src/ascend/op/*.cc glob picks the new file up.

src/op/builtin.cc is now byte-identical to upstream, and src/op/builtin.h
differs only by the rng_* and device_assert* redeclarations. Those stay: a
neutral pass (transform/loop_vectorize) and the Ascend codegen both name those
ops, while #2855 registers them under the CUDA dialect, so the neutral header is
the only place both can see. Two other neutral edits the branch had made are
dropped with the move -- a deleted blank line before kLocalVarInit, and a
"// Pointer access metadata op" comment on access_ptr.

Verified: tl.simd.*, tl.ascend_*, tl.rng_init, tl.device_assert, tl.access_ptr
and tl.any_sync all resolve through Op::Get at runtime. On 8x Ascend950DT (CANN
9.2.0) testing/ascend/ is 1010 passed, 12 skipped, 0 failed. On 8x H100 (CUDA
13.1) the USE_CUDA=ON build is clean and testing/python/ is 23 failed, 3551
passed -- the same 23 as before this change, with the 15-failure stable baseline
a strict subset and the other 8 the known flaky group (sm75 x4, tma_dsmem x2,
func_attrs x2).

* [Ascend] Drop the PTO-era codegen_py.cc append from the Ascend CMake

src/ascend/CMakeLists.txt added src/cuda/codegen/codegen_py.cc to the Ascend
source list unconditionally, with a comment explaining that PTO reuses
CodeGenTileLangPY. PTO is gone from this branch, and nothing else needs that
append:

  - codegen_py.h is included only by codegen_cutedsl.h, so the sole consumer of
    CodeGenTileLangPY is CodeGenTileLangCuTeDSL;
  - codegen_cutedsl.cc and rt_mod_cutedsl.cc both live in
    src/cuda/CMakeLists.txt's TILE_LANG_CUDA_ACTIVE_SRCS, which lists
    codegen_py.cc as well.

So with USE_CUDA=ON the append was a duplicate entry for a file the CUDA list
already owns, and with USE_CUDA=OFF nothing in the build referenced the class at
all. codegen_py.cc registers no op, global, or pass of its own.

Verified: a USE_CUDA=OFF rebuild no longer produces codegen_py.cc.o and still
links; testing/ascend/ is 1010 passed, 12 skipped, 0 failed. On the USE_CUDA=ON
box the reconfigure leaves the CUDA source set unchanged (ninja reports no work
to do) and codegen_py.cc.o is still produced by the CUDA list, so CuTeDSL keeps
its base class.

* [Ascend] Drop the PTO reference from the Ascend skill

Section 8 (SIMD MicroAPI) pointed at the PTO Micro-Instruction Spec for lane
widths, rounding modes, saturation behavior and mask predicates. The PTO backend
is gone from this branch, so the skill should not send a reader to it.

The section's own introduction already names the right source: the T.simd.* ops
map directly to Ascend CCE MicroAPI instructions.

This was the last PTO or VMI mention under .agents/.

* [Ascend] Lazy-load libascendcl via an ascendcl stub and gate the backend behind USE_ASCEND

Formalize the ad-hoc dlopen loader in ascend_module.cc into a CUDA/ROCm-style
stub library (libstub_ascendcl.so): the runtime module now calls the ACL
entrypoints directly and the stub resolves them lazily, preferring the copy
already loaded by torch_npu (RTLD_DEFAULT/RTLD_NEXT), then LD_LIBRARY_PATH,
then well-known CANN install roots. aclGetRecentErrMsg degrades to nullptr so
error reporting never throws; a missing required symbol is reported by name.

Add USE_ASCEND to TILELANG_BACKENDS with env-var handling and CANN
auto-detection (probed before the pip-provided CUDA toolkit). Only
target_utils.cc and layout/ascend_layouts.cc stay unconditional: common
transforms call TargetIsAscend, and tilelang.layout registers
tl.AscendFractalLayout at import time. TILELANG_USE_ASCEND_STUBS=OFF links the
real libascendcl instead.

Verified on NPU hardware (testing/ascend, 1010 passed) and on a CANN-less
GB200 host where a single USE_CUDA+USE_ROCM+USE_ASCEND wheel imports cleanly
and runs CUDA kernels.

* [Ascend] Move the copy and gemm hints into the Ascend dialect

66c003c3 (#3203) made each backend declare the op knobs it actually honors
instead of advertising another backend's vocabulary on the common surface. The
Ascend backend was still doing the latter: T.copy carried transpose,
l2_cache_ctrl, unit_flag_ctrl, sub_blockid, scale, pad_value and data_select in
tilelang/language/copy_op.py, and T.gemm carried unit_flag_ctrl in
tilelang/language/gemm_op.py, with the Ascend-only bodies inline in the common
implementations -- including a call from common code into
tilelang.ascend.language.dma for the pad-register write.

Both now follow the pattern the CUDA and ROCm dialects use:

- tilelang/language/copy_op.py keeps only the target-neutral core (src, dst,
  coalesced_width, annotations, loop_layout) and its signature is identical to
  upstream's again.
- tilelang/ascend/language/copy_op.py shadows copy() with the seven Ascend
  keywords plus _L2_CACHE_CTRL_MAP / _normalize_l2_cache_ctrl, packs them into
  the annotations dict, and delegates to the common implementation. The scale
  case still builds the three-region call itself, since the common copy only
  passes two regions. dual_copy in the same module now resolves to the Ascend
  copy and keeps working unchanged.
- tilelang/ascend/language/gemm_op.py shadows gemm() with unit_flag_ctrl and
  packs it the same way.
- tilelang/ascend/language/__init__.py exports both, so
  tilelang.ascend.language.copy/gemm resolve to the Ascend dialect while
  tilelang.language.copy/gemm stay on the CUDA dialect as #3203 intends.

Test placement follows the code: l2_cache_ctrl is an Ascend knob, so
test_copy_l2_cache_ctrl_string_annotation_matches_keyword moves out of
testing/python/language/test_tilelang_language_copy.py (which runs the default
CUDA facade) into testing/ascend/language/test_tilelang_ascend_copy.py, as
#3203 did for the ROCm k_pack regression test. The Ascend dialect test gains
assertions that the seven hints are present on the Ascend signature and absent
from the common one.

Verified on 8x Ascend950DT (CANN 9.2.0): testing/ascend/ 1012 passed, 12
skipped, 0 failed; examples/ascend/ 96 passed, 0 failed. On 8x H100 (CUDA 13.1)
testing/python/ is back to its 23-failure baseline with the moved test gone from
the generic suite.

* [Testing] Add a requires_ascend feature mark

testing/python/backend/test_tilelang_backend_auto_schedule.py exercises Ascend's
boolean scheduling flag, but it lives in the suite the CUDA-only CI job runs.
Since 9b1b8799 gated the Ascend sources behind USE_ASCEND, a build without it
does not register tl.enable_auto_schedule, so those four tests fail with
"Invalid config option" instead of skipping.

tilelang.testing re-exported TVM's requires_cuda/requires_rocm/requires_metal/
requires_llvm but defined none of its own. TVM's Feature constructor publishes
itself into Feature._all_features, so a backend can register its own mark on the
same registry; requires_ascend does that, which also makes `pytest -m ascend`
work.

Two fields are deliberately left unset, unlike requires_cuda:

- target_kind_enabled gates on TVM_TEST_TARGETS, and its DEFAULT_TEST_TARGETS
  fallback has no ascend entry, so setting it would skip every Ascend test
  unless the environment named ascend.
- target_kind_hardware would call tvm.device("ascend").exist, but ascend is a
  TileLang target kind and not a registered TVM runtime:
  tvm.runtime.enabled("ascend") raises "Unknown optional runtime ascend".
  Device availability is checked through the existing
  tilelang.ascend.target.check_ascend_availability() (torch_npu) instead.

That split matches the two meanings in play: compile_time_check is "this build
has USE_ASCEND", run_time_check is "this machine has an NPU". The scheduling
tests only need the first, so they take the compile-only marks and keep running
on a host with the backend built but no device attached.

Verified: with USE_ASCEND=ON (8x Ascend950DT) 3 passed, 2 skipped -- the skips
are the CUDA-gated cases -- and `pytest -m ascend` collects all five. With the
CUDA-only build on 8x H100 all five report "Compile-time support for Ascend not
present".

* [Testing] Gate the Ascend suite on a build that has the backend

testing/ascend/ holds 83 files and had no backend gating at all. On a build
without USE_ASCEND, running it produced a wall of failures -- 4 failed out of 5
in testing/ascend/language/test_tilelang_ascend_copy.py -- instead of saying
that the backend is not in the library.

The gate has to act at collection time, not through a per-item mark. Four
modules do backend work at import scope and raise while pytest imports them,
before any mark could apply:

  simtvf/test_simtvf_fragment_narrow_checker.py  Target kind "ascend" is not defined
  test_tilelang_ascend_multibuffer_alignment.py  No module named 'torch_npu'
  layout/test_tilelang_ascend_norm_fn.py         Operator tl.ascend_set_hf32_mode is not registered
  transform/test_ascend_legalize_simd_merging.py Operator tl.simd.pset is not registered

So testing/ascend/conftest.py ignores its directory when
tilelang.testing.ascend_backend_compiled() is false, which keeps those imports
from happening. That predicate becomes public for the conftest to use; the
requires_ascend feature keeps it as its compile-time check.

testing/conftest.py already errors out when nothing was collected, to catch
misconfigured runs. An entirely gated Ascend suite is a legitimate outcome, so
it now reports one line instead: "Skipped the Ascend test suite: the Ascend
backend is not built into this library."

The gate is only the compile-time half of requires_ascend. Device availability
stays a per-test concern, so source-only lowering tests keep running on a host
with the backend built but no NPU attached.

Verified: with USE_ASCEND=OFF (8x H100, CUDA-only build) `pytest testing/ascend/`
reports the skip line and exits 0, with no collection errors. With USE_ASCEND=ON
(8x Ascend950DT) the suite is unchanged at 1012 passed, 12 skipped, 0 failed.

* [Ascend] Register tl.conflict_hint from the Ascend backend

src/op/schedule_hint.cc registered tl.conflict_hint, the marker op for
user-declared auto-schedule conflict facts. Nothing about it is neutral:

- it is emitted only by tilelang/ascend/language/schedule_hint.py, through
  assume_conflict / assume_no_conflict;
- it is consumed only by src/ascend/transform/normalize_conflict_hints.cc,
  which rewrites the marker into Fors' "conflict_hint" annotation and the
  tilelang_root "tl.root_conflict_hints" attribute before AutoSchedule reads
  them.

AutoSchedule is the Ascend scheduler, so the op belongs with the rest of the
Ascend registrations. It moves into src/ascend/op/builtin.{h,cc} and
src/op/schedule_hint.cc is deleted.

Using TIR_DEFINE_TL_BUILTIN here is exact, not approximate: it registers under
"tl.conflict_hint" and sets TScriptPrinterName to "conflict_hint", which is
precisely what the original hand-written TVM_REGISTER_OP did. num_inputs stays
-1 and the effect kind stays kPure, so RemoveNoOp still drops an unconsumed
marker when auto-schedule is off.

Two neighbours in src/op/ were audited and deliberately left alone:

- simt_vf.{h,cc} and simd_vf.{h,cc} define SimtVFOp and SimdVFOp, which neutral
  passes must dispatch on (src/transform/lower_tile_op.cc, storage_rewrite.cc,
  layout_inference.cc, lower_device_kernel_launch.cc, loop_fusion_utils.h).
  Moving them would make those neutral files depend on an Ascend header.
- the sf / sf_range fields the branch added to CopyNode in src/op/copy.h are
  shared-struct members read by the Ascend L1->L0 lowering, the same shape as
  GemmNode's sfa/sfb, and are not op registrations.

Verified: tl.conflict_hint resolves with num_inputs == -1; testing/ascend/ is
1012 passed, 12 skipped, 0 failed on 8x Ascend950DT (CANN 9.2.0).

* [Ascend] Move the Ascend scheduler attributes into the Ascend tree

src/transform/common/attr.h carried six constants the branch had added. Four of
them name IR that only exists between MaterializeScheduleUnits and
LowerScheduledTIR, which are both Ascend passes:

  kAscendPerCoreTask, kAscendTask, kAscendStage   read only under src/ascend/
  kScheduleUnit                                   read only under src/ascend/
                                                  (materialize_schedule_units,
                                                  auto_schedule/*)

They move to a new src/ascend/transform/attr.h, next to the other Ascend
transform headers (ascend_pipe.h, buffer_version.h, core_mask.h,
estimate_latency.h). The constants keep the tvm::tl::attr namespace and their
names, so the nine consumers only needed the include.

Two more constants in that hunk, tilelang_simt_vf_captures and
tilelang_simd_vf_captures, are deleted outright: grepping both the symbols and
their string values leaves only their own declarations, so nothing has read them
since the kernel-launch frame stopped using per-backend marker annotations.

kBufferVersion stays in the neutral header. Two of its three readers are Ascend
(normalize_buffer_version.cc, rewrite_buffer_version_layout.cc) but the third is
src/transform/lower_tile_op.cc, which is backend-neutral, so moving it would
mean a neutral pass including an Ascend header.

src/transform/common/attr.h is now free of Ascend-only names.

Verified: the build is clean and testing/ascend/ is 1012 passed, 12 skipped,
0 failed on 8x Ascend950DT (CANN 9.2.0).

* [Ascend] Give Ascend its own on-chip utils and leading-dim access ptr

src/op/utils.cc had gained a leading_dims_offset_only parameter on
MakeAccessPtrFromRegion. Only Ascend used it -- the three call sites are all in
src/ascend/op/copy.cc -- so the neutral function goes back to upstream verbatim
and the behaviour moves to src/ascend/op/utils.h as
MakeAscendLeadingDimAccessPtr.

The helper exists because the Ascend L1->L0 loads pass the M/K start positions
as their own intrinsic arguments (ascend_load_cbuf_to_ca(dst, src, mStartPosition,
kStartPosition, ...) and the asc_copy_l12l0a_mx scale companion). The pointer
must therefore carry only the leading dims -- the pipeline version, say -- and
counting the inner two dims into its offset would place the position twice.

It is built on the neutral MakeAccessPtrFromRegion rather than duplicating the
stride arithmetic: the innermost two dims are re-based to zero before
delegating, which leaves their extents (and so the reported extent) alone and
drops them from the offset. ffi::Array is copy-on-write, so the caller's region
is untouched. Regions with fewer than two dims are left as they are, because a
1-D region has no inner 2-D position to separate and the neutral path uses its
min as the offset.

The five on-chip scope predicates the branch had also put in the neutral header
-- IsL1Buffer, IsL0ABuffer, IsL0BBuffer, IsL0CBuffer, IsAscendOnChipBuffer --
move to the same new header. Their twelve consumers are all under src/ascend/;
the scopes they name (shared.l1, shared.l0a/b/c and the .dyn variants) exist only
on Ascend. They compose the neutral IsSharedBuffer, so the include direction
stays Ascend -> neutral.

src/op/utils.{h,cc} are now byte-identical to upstream.

Verified: the build is clean and testing/ascend/ is 1012 passed, 12 skipped,
0 failed on 8x Ascend950DT (CANN 9.2.0).

* [Ascend] Drop an unreachable VF-block guard from InjectSoftwarePipeline

src/transform/inject_pipeline.cc's VisitStmt_(const SBlockNode *) early-returned
for blocks named SIMT_VF, VECTOR or CUBE, skipping the pass's buffer bookkeeping
for them.

The guard cannot fire. Those block names are produced only by the Ascend
SimtVF / Cube / Vector frames (src/ir.cc, and lower_scheduled_tir.cc for the
scheduled path), while InjectSoftwarePipeline is invoked by the cuda, rocm,
metal and webgpu pipelines and by nothing else -- tilelang/ascend/pipeline.py
never calls it. So the pass never sees those names, and on the targets that do
run it they have no producer.

That also explains why removing it changes nothing: testing/ascend/ is 1012
passed, 12 skipped, 0 failed and examples/ascend/ is 96 passed, 0 failed both
with and without it. It was not merely untested, it was unreachable, and it made
a backend-neutral pass recognise Ascend block names.

src/transform/inject_pipeline.cc is now byte-identical to upstream.

The same SIMT_VF name check exists in four other neutral passes
(common/loop_fusion_utils.h, layout_inference.cc, lower_tile_op.cc,
thread_storage_sync.cc). Those passes *are* in the Ascend pipeline, so they are
not automatically dead and are left alone here.

* [Ascend] Restore the CPU pipeline's MaterializeKernelLaunch call

00a1cee0 taught MaterializeKernelLaunch to separate the program-index space
from the thread space, adding lower_grid_binding as a new leading parameter that
defaults to None (follow lower_thread_binding), and the Ascend pipeline passes
it explicitly: lower_grid_binding=True with lower_thread_binding=False, because
Ascend has a real 1-D core grid but no threads at kernel scope.

The same commit also expanded the CPU pipeline's call to spell out
lower_grid_binding=False and default_threads=None. Neither is needed: with
lower_thread_binding=False the parameter already derives False, and
default_threads is documented as ignored when lower_thread_binding is False. The
expanded call therefore restates the defaults and nothing else.

The CUDA, ROCm, Metal and WebGPU pipelines were left on the old call and keep
working off the same default, which is the point of the None fallback. CPU can
do the same, so it goes back to upstream's one-liner and
tilelang/cpu/pipeline.py is byte-identical to upstream again.

Verified: testing/python/cpu/ is 88 passed, 0 failed.

* [Ascend] Own allow_autoschedule in the Ascend pipeline

tilelang/backend/pass_pipeline/pipeline_utils.py had gained
allow_autoschedule(), a sibling of the backend-neutral helpers that live there
(allow_vectorize, allow_global_thread_synchronization, should_*). It reads
TL_ENABLE_AUTO_SCHEDULE, which exists only when the Ascend backend is compiled
in, and it has exactly one consumer: tilelang/ascend/pipeline.py, which uses it
to decide whether to run the AutoSchedule chain.

It moves into tilelang/ascend/pipeline.py, which is where the pass order it
gates already lives (per the backend layout rule that a backend keeps its own
pipeline). The name and behaviour are unchanged, so the call site is untouched;
only the import moves, in that module and in the Ascend-gated
test_tilelang_backend_auto_schedule.

tilelang/backend/pass_pipeline/pipeline_utils.py is now byte-identical to
upstream.

Verified: allow_autoschedule resolves from tilelang.ascend.pipeline;
testing/python/backend/test_tilelang_backend_auto_schedule.py is 3 passed,
2 skipped (the skips are its CUDA cases); testing/ascend/ is 1012 passed,
12 skipped, 0 failed on 8x Ascend950DT (CANN 9.2.0).

* [Ascend] Move the NPU frontend frames out of the shared src/ir.cc

src/ir.cc carried four Ascend-only TIR frames plus the Ascend mixed launch
interleaved with the shared #3186 launch machinery. Only the shared frames
remain there now:

  kept in src/ir.cc       tl.Parallel, tl.Pipelined, tl.Persistent,
                          tl.KernelLaunch, tl.WarpSpecialize, tl.SideEffect
  moved to src/ascend/ir.cc
                          tl.SimtVFFrame, tl.SimdVFFrame, tl.CubeFrame,
                          tl.VectorFrame, tl.MixedKernelLaunch, and the
                          buffer-stealing helpers their Enter/ExitWithScope
                          use

src/ir.cc loses 576 lines and gains 3: the new include, plus the two .def
chains that now terminate where the Ascend registrations were removed.

A dialect launch variant still has to build a tl.KernelLaunchFrame --
tilelang.ascend.language.kernel binds its launch surface to that object type --
so KernelLaunchFrameNode, KernelLaunchFrame and the two frame factories they
are assembled from (MakeThreadBindingFrame, MakeLaunchThreadFrame) move
verbatim into a new neutral src/ir.h, static -> inline. That is the only
structural change outside the move.

src/ascend/ir.cc goes into TILE_LANG_ASCEND_ALWAYS_SRCS rather than the
Ascend-only source list. tilelang/__init__.py imports tilelang.ascend
unconditionally and tilelang.ascend.language.frame binds tl.CubeFrame and
friends with @register_object at import time; a missing type key raises
"Cannot find object type index", so keeping the translation unit Ascend-only
would break "import tilelang" on every USE_ASCEND=OFF build. This mirrors why
ascend_layouts.cc is already in that list. The frames carry no CANN
dependency.

Verified: clang-format v23.1.0 clean on all three C++ files; git diff --check
clean; format.sh clang-format hook passes.
Verified: USE_ASCEND=ON build clean. testing/ascend/ 1012 passed, 12 skipped;
examples/ascend/ 96 passed; testing/python/cpu/ 88 passed.
Verified: USE_ASCEND=OFF configure compiles exactly ir.cc, target_utils.cc and
ascend_layouts.cc from src/ascend/, with one object file each; both ir.cc
translation units also compile standalone under that configuration.
Verified: on a CUDA host (USE_ASCEND=OFF) the library links, "import tilelang"
succeeds with all 9 FFI entrypoints and the 4 Python frame classes present,
and the full testing/python/ suite reports 23 failed / 3554 passed / 481
skipped, matching the pre-change baseline failure for failure.

* [Ascend] Declare NPU codegen capabilities as target kind attributes

Three shared passes branched on TargetIsAscend to reach NPU-specific
decisions. Each of those is a property of the target's codegen rather than
of the target's name, so the Ascend target kind now declares them and the
passes query the declaration:

  supports_kernel_status_return  (default false)
      AscendC kernel entries return void, so the host cannot read an int32
      status code from the device the way CPU-codegen targets do.
  supports_vector_predicate      (default true)
      AscendC emits lane-wise vector predicates, so a vectorized Select
      does not need a uniform condition.
  max_vector_bits                (default 64)
      Widest vector load/store the vectorizer may issue.

split_host_device.cc now reads supports_kernel_status_return before falling
back to its device-type heuristic, which keeps CPU / ext_dev / Hexagon on
the status-code path. loop_vectorize.cc reads the other two: the
Select-condition uniformity rule and MaxVectorLoadBits.

The attributes are declared only on the ascend kind, so every other target
sees them absent and keeps its previous behaviour unchanged.

The three MaxVectorLoadBits call sites all receive a target that cannot be
undefined: loop_vectorize.cc uses Target::Current(false), which checks
rather than returning a null Target, and the layout cost model is reached
from LayoutInference, which ICHECKs the target attribute.

Verified: Target("ascend") resolves the three to false/true/64, while cuda,
llvm, ext_dev, hexagon and rocm report all three absent, so the fallbacks
are untouched.
Verified: USE_ASCEND=ON build clean; clang-format v23.1.0 clean on all
three files. Cold-cache runs: testing/ascend/ 1012 passed 12 skipped,
examples/ascend/ 96 passed, testing/python/cpu/ 88 passed,
testing/python/transform/test_tilelang_transform_loop_vectorize.py
16 passed.
Verified: on a CUDA host (USE_ASCEND=OFF) the build is clean and the full
testing/python/ suite reports 23 failed / 3554 passed / 481 skipped,
matching the pre-change baseline failure for failure.

* [Ascend] Declare the grid-only launch model as a target attribute

Replace the TargetIsAscend branches in LowerDeviceKernelLaunch with a
`launch_grid_only` attribute on the ascend target kind: launch-param
collection and the thread_extent metadata keep only the blockIdx.* grid
axes when the attribute is set, since threadIdx domains are SimtVF
region-local lanes baked into codegen rather than runtime launch
dimensions.

Also drop the forced call_packed guard and its launch_params ICHECK
exemption: MakePackedAPI runs before this pass and rewrites the host
function's target to the host target, so an ascend callee can never
compare same-device-type with its caller here (verified empirically:
531/531 launch sites in the full suite take the packed path already).

* [Ascend] Restore MakePackedAPI's strict undefined-variable check

The local/local.fragment exemption papered over device-internal handles
leaking into host functions during early SIMD-VF bring-up. The leak no
longer exists in the current pipeline (the full Ascend suite passes with
the strict upstream check), and a genuinely leaked variable should fail
here with an actionable message instead of crashing later inside host
codegen.

* [Ascend] Replace shared-pass VF dispatch with registries and move the VF ops under src/ascend

LowerTileOp and LayoutInference no longer hard-code Ascend behavior;
backends contribute through three registries:

- LoweredParallelLoopHook (transform/common/lower_hooks.h): CUDA
  registers PTX cp.async injection from its own translation unit,
  removing the TargetIsCuda branch and the cuda includes from the
  shared pass.
- AttrNodeRemapHook: Ascend registers the `tl.buffer_version` key remap
  from src/ascend/transform/buffer_version.cc. The map stays Var-keyed
  deliberately: distinct versioned storages may share one name (see
  test_auto_schedule_distinguishes_same_name_buffer_version_annotations),
  so names cannot serve as stable keys.
- RegionOpImpl (op/region_op.h): block-name keyed {make, enter_scope}.
  LayoutInference wraps region blocks via `make` for its inference
  worklist; both passes push the region execution scope via
  `enter_scope`, replacing two duplicated copies of the SimtVF
  thread-context machinery. The SimtVFOp/SimdVFOp lowering branches in
  LowerTileOp were elaborate no-ops (their Lower() is a passthrough)
  and are removed outright.

simt_vf/simd_vf move to src/ascend/op/ and register their RegionOpImpl
there. They stay in TILE_LANG_ASCEND_ALWAYS_SRCS, matching the
always-available VF frontend frames, so every build configuration
behaves exactly as before.

* [Ascend] Isolate ThreadSync and restore upstream pass

Keep the existing Ascend behavior in a private implementation with isolated symbols and a separate pass registration. Route the Ascend pipeline through the private pass.

Restore the common ThreadSync implementation and regression tests byte-for-byte to upstream-main at 66c003c3.

Validation: both implementations compiled and linked in an audit DSO; 64 IR tests passed and 20 original-versus-private IR comparisons matched. Formatting and lint checks passed.

* [Ascend] Rebase constr_visitor onto the upstream text, keeping the two load-bearing deltas

The fork variant of constr_visitor.h diverged from upstream by 284 lines,
of which only two hunks are functionally required (upstream #3220):

- Constr::Populate simplifies a predicate before EnterConstraint so a
  bound boolean guard expands into the condition it represents (five
  flag-allocation tests fail without it).
- ConstrSet::Populate installs binds before predicates so expansion
  works on Merge/RenameFrom-scrambled sequences (four multibuffer
  counter tests fail without it), at the documented cost of wide
  bind-time bounds.

Everything else (Guard visibility, Merge comparison details, comments,
style) adopts the upstream text. TaskAwareConstrVisitor no longer names
the now-private Guard type; it follows the base class's own SeqStmt
idiom instead.

* [Transform] Restore bind-before-use order when merging constraint sets (#3221)

ConstrSet::Merge concatenates two lexically-collected sets, so an entry
of the first set can read a bind var the second set defines. Replaying
that order digests the predicate before its definition exists --
EnterConstraint evaluates eagerly, exactly like Bind -- and the guard
fact is silently lost, which downstream reports as unprovable guard
disjointness on merged sets (#3220).

Two changes:
- Merge now normalizes the concatenated list: entries keep their order,
  except that a bind is hoisted (with the binds its definition reads,
  recursively) directly before the first entry reading its var. Within a
  single lexically-collected set nothing reads a later bind, so the
  program-order tightening guarantee of #2805 is untouched.
- Constr::Populate simplifies a predicate against the installed binds
  before entering it, so a guard held in a bound boolean var expands
  into the condition it represents.

tl.analysis.ConstrSetsProve exposes merge+populate+CanProve to Python so
the merged-set scenarios stay testable without staging a pass pipeline.

(cherry picked from commit 825bfd7ead)

* [Transform] Preserve MaterializeKernelLaunch positional compatibility

* [Ascend] Isolate LayoutInference/LowerTileOp and restore the upstream passes

Following the ThreadSync isolation precedent: Ascend owns full forks of
the two lowering orchestrators, and the shared passes go back to the
upstream text byte-for-byte.

- src/ascend/transform/{layout_inference,lower_tile_op}.cc (namespace
  tvm::tl::ascend, registered as tl.transform.AscendLayoutInference /
  AscendLowerTileOp) embed the Ascend behavior directly: VF region
  scopes (helpers in vf_regions.h), the Var-keyed tl.buffer_version
  remap, int32 layout-index coercion, and symbolic thread-extent
  tolerance. The CUDA cp.async post-processing is dropped from the
  Ascend fork. VF regions no longer enter the inference worklist: the
  entries were placeholders (InferLayout returned nothing), so the
  hollow SimtVFOp/SimdVFOp wrapper classes are deleted outright.

- src/transform/{lower_tile_op.cc,layout_inference/layout_inference.cc}
  and src/cuda/transform/ptx_async_copy_injector.cc are restored to
  tileai/main. The registry indirection they consumed (lower_hooks,
  region_op, the buffer-version remap hook) is removed.

The forks start at upstream #3221 content; upstream changes to the
shared passes now need manual porting, and the shared/forked pair
doubles as a cross-check of lowering behavior.

* [Op] Promote the rng and device_assert builtins to common

rng_init/rng_rand/rng_rand_float and device_assert/device_assert_with_msg
are consumed by the CUDA and Ascend codegens and by the shared vectorizer,
so their home is src/op/builtin.{h,cc}. The shared header already carried
duplicate declarations as a layering shim for the Ascend side; move the
definitions over and drop the CUDA-side declarations and definitions.
Registration names are unchanged; this is a pure relocation.

* [Frontend] Split the launch frame factory declarations from their definitions

MakeThreadBindingFrame and MakeLaunchThreadFrame move from inline
definitions in ir.h to definitions in ir.cc; the header keeps the
declarations and the materialization contract comments. Pure layout
change: both frontend TUs (src/ir.cc, src/ascend/ir.cc) keep linking
against the single shared implementation.

* [Transform] Census dialect copy refinements in ReducerPlanAndMaterialize

The per-buffer use census matched copies by op identity, so a copy
spelled with a dialect op would be invisible to reducer planning.
Recognize any op whose builder yields a strict CopyNode refinement;
spellings that build a plain CopyNode (async_copy, tma_copy) keep
their historical exclusion, so existing targets census identically.

* [Ascend] Type the Ascend copy as a tl.tileop.ascend_copy dialect op

The Ascend frontend now emits tl.tileop.ascend_copy, whose TLOpBuilder
constructs AscendCopyNode : CopyNode — the MX scale-factor companion
region plus decode-once typed views of the Ascend copy hint annotations
(nd2nz, dual_dst_ctl, transpose, data_select, unit_flag_ctl,
sub_blockid, l2_cache_ctrl, pad_value). Lowering and layout inference
move onto the node's virtual methods; the CopyImpl registry entry stays
as a bridge that upgrades plain tl.tileop.copy calls (e.g. copies
synthesized by ReducerPlanAndMaterialize) onto the same implementation,
and Ascend matchers accept both spellings via IsAscendCopyCall.

This restores shared src/op/copy.cc to upstream byte-for-byte and
shrinks the src/op/copy.h delta to un-finalizing CopyNode so dialects
can extend it.

* [Transform] Restore common StorageRewrite to upstream

Sync src/transform/storage_rewrite.cc exactly with tile-ai/tilelang main at 030556de97. Ascend uses its private MergeUBAllocations pass and no longer invokes StorageRewrite in its default pipeline.

Validation: rebuilt the restored pass in isolation against the matching TVM runtime; all 10 mixed-dtype/control cases match upstream, and all 13 lexical allocation and inplace-toggle tests pass.

* [Ascend] Sink the VF loop-fusion and validator variants into the dialect

ParallelLoopFuserSkipSimdVF and the validator's SIMD_VF skip carried
Ascend VF vocabulary in shared headers while serving exactly one
caller, AscendLayoutInference. Both move next to it as subclasses.

src/transform/layout_inference/ is byte-identical to upstream again;
loop_fusion_utils.h keeps only a target-neutral extension seam (a
virtual PreserveParallelLoopNest hook, required because the ForNode
visitor is final). The SIMD_VF skip itself stays load-bearing: SIMD_VF
T.Parallel loops are lowered by AscendSimdVFLowerParallel and never
carry loop-layout annotations, so the post-inference validator must
not descend into those regions.

* [Ascend] Own the dialect's print and device_assert surfaces

tilelang/cuda/language/print.py decided the shared-buffer print gating
by probing the host for an NPU at import time (torch.npu.is_available),
and tilelang/cuda/debug.py extended its assert gate the same way — the
compile target never entered the decision. On an NPU host that meant
CUDA shared prints lost their single-thread gating; on a CUDA-only host
an Ascend kernel's print would call get_thread_bindings inside a
threads-less kernel.

The Ascend dialect now owns both policies: ascend/language/print.py
prints without CUDA-style thread gating (a kernel body runs once per AI
core and UB is core-local), reusing the target-neutral debug_print_*
macro emitters; ascend/debug.py gates device_assert on the NPU toolkit
the way the CUDA dialect gates on nvcc. Both CUDA modules and
cuda/transform/__init__.py return to upstream byte-for-byte — the
latter's LowerPTXAsyncCopy wrapper had no callers and its FFI pass no
longer exists.

* [Ascend] Hold the target scope at lower call sites, not inside lower

tilelang.lower expects its caller to hold the target scope: the
vectorize planner consults Target::Current(false) (fail-loud), the JIT
path and tools/compile_only already enter it, and upstream's own tests
wrap bare lower calls the same way. The fork's with-target inside
lower_to_host_device_ir papered over that contract for every caller;
drop it and restore tilelang/engine/lower.py to upstream byte-for-byte.

The thirteen Ascend tests whose lowering actually reaches the planner
now enter the scope explicitly at their call sites, as does the
annotate_unlimit_memory example. Also remove a stray module-scope test
invocation in test_simtvf_fragment_narrow_checker.py that ran the test
at import time, and conftest's unused pytest import.

* [Ascend] Move the Ascend postproc callback into the dialect

register_ascend_postproc(_callback) was appended to the shared
tilelang/engine/callback.py; its FFI consumer is Ascend-owned
(rt_mod_ascend.cc looks up tilelang_callback_ascend_postproc before
emitting the final AscendC source). Move the registration API to
tilelang/ascend/callback.py and restore engine/callback.py to upstream
byte-for-byte.

* [Ascend] Sink the SimtVF nested thread scope into the dialect

The SimtVF context stack and the SimtVF-first preamble on the four
thread accessors lived in shared tilelang/language/kernel.py, justified
by backend-neutral callers reaching the accessors by name. Only one
such caller remained — rng_init's default sequence-id derivation — so
the justification no longer holds:

- tilelang/ascend/language/kernel.py now owns the context stack and
  dialect accessors that consult the innermost SimtVF scope before
  falling back to the shared launch-frame accessors; the SimtVF frame
  already pushed through this module.
- tilelang/ascend/language/random.py derives rng_init's default seq
  (lane id + core index * lane count) through the dialect accessors and
  delegates the emission; rng_rand/rng_rand_float re-export unchanged.
- tilelang/language/kernel.py returns to upstream byte-for-byte.

* [Language] Shrink the shared frontend's Ascend deltas

Three fossil reductions, no behavior change:

- gemm_op.py: drop the fork-added sfa/sfb/sf_k_start plumbing from
  _gemm_impl. The Ascend blockscaled_gemm now appends the scale-factor
  regions as trailing tl.tileop.gemm args itself, mirroring upstream's
  tcgen05_gemm_blockscaled wire format (GemmNode's sfaRegion/sfbRegion
  are upstream-native).
- copy_op.py: fold _normalize_copy_regions_with_extents back into
  _normalize_copy_regions. The extents returns and copy()'s dst_orig
  had no consumer anywhere (leftovers of an early nd2nz-scatter
  frontend); the live parts — the element-count relaxation for
  differently-major'd fractal copies and the helper being importable
  by the Ascend dialect's scale= path — are kept.
- eager/builder.py: drop "local.simd.var" from is_var. The scope never
  had a producer in any revision; SimdVF variables use plain local.var.

* [Testing] Cover bare-name lower callers with the target scope

The caller-holds-the-target-scope migration enumerated call sites by
grepping tilelang.lower(, which missed the twenty-four Ascend test
files importing lower by bare name from tilelang.engine.lower; the
twenty-seven of their tests whose lowering reaches the vectorize
planner failed with "Target context required". Enter the scope
explicitly at those call sites — inside the shared _device_script /
_pass_snapshots helpers where the tests funnel through one, merged
into the pytest.raises / PassContext with-statements elsewhere.

* [Ascend] Own alloc_shared and drop the orphaned major-layout helper

tilelang/language/allocate.py carried two Ascend intrusions:

- _annotate_ascend_major_layout had no caller anywhere: it served the
  #149-era alloc_l1(major=...) surface, orphaned when #244 moved L0
  layout selection into the consuming gemm. Deleted.
- The bool workaround (bool buffers downgrade to the static "shared"
  scope for the smem-merge pass) was gated on an NPU host probe
  (#264) — the wrong axis: an NPU host compiling CUDA kernels skipped
  the hack too. The common alloc_shared is unconditional again, and
  the Ascend dialect shadows alloc_shared without the workaround (UB
  takes bool buffers in the requested scope), matching the dialect's
  Kernel/copy/print ownership pattern.

The shared file is byte-identical to upstream again.

* [IR] Align PersistentFor traversal with upstream

Restore upstream grouping, coordinate decoding, and tail/break handling. Keep the num_stages and annotations extensions required by Ascend pipelining.

Use full-last-dimension groups in the two Ascend FP8 examples to preserve their row-major assignment. Add host-only TIR regressions for static and symbolic domains, ordering, tail coverage, and pipeline annotations.

* [Ascend] Clear the tileop/utils layers' Ascend residue

- gemm_base.py returns to upstream byte-for-byte: the dead block_scaled
  alias is gone, and GemmMAD now owns the two Ascend decodes — the
  unit_flag_ctrl annotation and the is_blockscaled widening (an Ascend
  L0-input block-scaled gemm carries no SF regions on the node, so the
  base's structural predicate cannot see it).
- utils/tensor.py: map_torch_type had no caller — added by the FP8
  codegen work, orphaned when the KernelParam refactor took over dtype
  mapping, absent upstream.
- utils/language.py: drop "shared.l1" from is_shared. GemmMAD never
  uses the is_gemm_* predicates and the only shared-pass consumer
  (DecoupleTypeCast) cannot reach an L1 cast — L1 is only touched by
  DMA copies and cast-carrying DMA copies are rejected up front.

* [Ascend] Sink the SimdVF reduce path into the dialect

The shared reduce macro branched on the Ascend SimdVF builder state
through a deferred dialect import (module scope would cycle). The
dialect now owns that path: tilelang/ascend/language/reduce_op.py
emits shared-to-shared reductions inside SimdVF directly on the UB
regions (no fragment exists there to round-trip through) and delegates
everything else; the thin reduce_* wrappers are redeclared since the
common ones call the common reduce by module-level binding.

The remaining shared reduce_op.py delta is the target-neutral region
normalization (retrieve_shape/_get_buffer), which lets reduce accept
BufferRegion/BufferLoad operands like copy and gemm do — an upstream
candidate.

* [Misc] Restore .gitignore and .pymarkdown to upstream

The pymarkdown front-matter extension protected the YAML-front-matter
skill files back when the hook still covered them; they live under
.agents/ now, which upstream's own hook config excludes, and a full
--all-files run passes with the upstream config. The .gitignore
additions were personal-environment entries (CLAUDE.md, .clangd, a
local scratch script) and Ascend artifact dirs; the latter belong in
.git/info/exclude on the machines that produce them.

* [Lint] Apply format.sh fixes across recently touched files

clang-format reflows in ascend/op/builtin.h, ascend lower_tile_op.cc
(include order) and materialize_kernel_launch.cc (lambda wrap); ruff
drops unused imports in jit/adapter/{libgen,wrapper}.py and fixes
spacing in testing/__init__.py. No behavior change; build and smoke
tests pass.

* fix(ascend): honor the element width named by a BRC_* vld dist (#446)

T.simd.vld derives the vector type -- and with it the emitted C API and its
load granule -- from the source buffer's element type, so a BRC_B16/B32 dist
over a narrower buffer silently degraded to the buffer's width. Over a uint8
buffer, BRC_B16 broadcast a single byte instead of the 16-bit element the
suffix names, dropping the upper byte of a packed pair.

Honour the suffix for BRC_*: widen the result past the buffer element width and
size the access footprint as one element of that width. An unknown suffix, or
one too narrow to be a whole element of the source buffer, is now rejected
instead of loading the wrong granule.

Other distributions keep taking their width from the source buffer: UNPK_*
legitimately widens, so its suffix describes the loaded data rather than the
result and cannot follow the same rule.

(cherry picked from commit 177321c0e93205a048c2cc6efa3ac5322b9d7803)

* fix(ascend): treat out-of-VF fills as scalar (#448)

Two related defects made a T.fill outside a VF block misbehave.

1. Pipe classification. AutoSchedule and InsertSync take a task's hardware pipe
from GetAscendFillPipeMask, which reported PIPE_V for every non-L1 fill. A fill
only runs on the vector unit inside a VF block, where the enclosing block
already contributes PIPE_V through GetAscendBlockPipe; outside one it lowers to
element-wise BufferStores issued by the scalar unit. InsertSync therefore saw a
PIPE_V producer feeding a PIPE_V consumer and elided the S->V handshake between
them, letting the consumer read stale UB contents: a kernel that zero-fills the
invalid tail of a partially valid UB tile before reducing it in a T.SimtVF
returned garbage for the tail tile. Restore the pre-#426 classification, so a
non-L1 fill is PIPE_S and an L1 fill is MTE2.

2. Codegen. Such a fill still emitted a vectorized broadcast store, and the
vector constructors it needs (make_float2/make_int2) are simt_callee-only, so
bisheng rejects them outside a VF body: a static T.fill over a 64x384 UB tile
did not compile at all. Emit one scalar store per lane for a vector broadcast
store issued outside a VF body, which is the element-wise scalar work the
scalar unit would run anyway.

Validation: the reported reproducer now returns [64.0, 64.0, 1.0] instead of
[64.0, 64.0, 3.0]; new coverage in
testing/ascend/language/test_tilelang_ascend_dynamic_ub.py and
testing/ascend/language/test_tilelang_ascend_fill.py (device) and
testing/ascend/analysis/test_ascend_task.py (lowered sync events); the full
testing/ascend suite passes (1217 passed, 12 skipped).

(cherry picked from commit f192961587b6493ff9fcf307b0f7584c2af5770b)

* [Examples] Restore deepseek_mhc's example_mhc_post to upstream

* [JIT] Drop the identity storage_shape hook

KernelParam.storage_shape came in with the packed-FP4 work, when Torch
allocation needed a storage shape distinct from the logical one. Torch
2.8's float4_e2m1fn_x2 dtype made the two agree element for element, a
cleanup reduced the hook to `return list(self.shape)`, and a dedicated
test pinned the identity. A live-called identity with a test asserting
nothing happens is speculation, not API; the wrapper call site returns
to param.shape with the adapter changes in the next commit.

* [Ascend][JIT] Move the Torch NPU integration out of the shared wrapper

The DLPack exchange patch (~150 lines) lived inside the shared
cython_wrapper.pyx only because that was the one pyx the build
compiled; it patches torch.Tensor's global exchange table and has no
coupling to the wrapper. It now builds as its own extension,
tilelang_ascend_npu_exchange, from tilelang/ascend/torch_npu_exchange.pyx
(pure CPython/DLPack, torch_npu imported lazily), and the
tilelang/ascend/torch_exchange.py facade points at it — callers are
unchanged.

The wrapper's three torch.npu.is_available() host probes are gone:
device and raw-stream selection are injected by the adapter, which
knows the compile target (_device_providers), with defaults preserving
the upstream CUDA behavior. An NPU host compiling CUDA kernels no
longer allocates outputs on the wrong device. The vestigial
__cinit__ target parameter and the storage_shape call site are cleaned
up along the way.

* [Ascend] Assemble L0 MAD gemm calls in the dialect

The shared _gemm_impl carried an is_ascend_l0_gemm branch that
substituted allocation-shape M/N/K before the static-tile check (L0
region extents give the effective, possibly symbolic MAD geometry) and
an assert allowing higher-order C regions only for L0C. Both sink into
the Ascend dialect: _l0_gemm_call assembles the tl.tileop.gemm call
itself for L0A/L0B/L0C operands, mirroring the common assembly's
positional contract the same way upstream's tcgen05_gemm_blockscaled
does; everything else keeps delegating to _gemm_impl. The remaining
shared delta is target-neutral (symbolic shape proofs, the clear_accum
write-only C access mask).

* [Language] Drop copy's no-op empty-annotations normalization

annotations=ann and annotations=None produce structurally and
textually identical Calls when ann is empty (both normalize to an
empty Map), so the conditional was pure noise; restore the upstream
spelling.

* [Ascend] Fix two tests missed by the dialect and target-scope moves

Two test files were missed when the Ascend surfaces moved out of the shared
frontend and the target scope moved to lower()'s call sites:

- testing/ascend/language/test_tilelang_ascend_simd_vld_brc.py still imported
  ``tilelang.language as T``, which no longer re-exports the Ascend dialect, so
  ``T.SimdVF`` raised AttributeError in the three tests that build a SimdVF body.
  Use ``import tilelang.ascend.language as T`` like the other Ascend tests.
- testing/ascend/analysis/test_ascend_task.py called
  ``lower(main, target="ascend")`` without holding the target scope in two
  tests, so the vectorize planner's ``Target::Current(false)`` check failed with
  "Target context required". Wrap those calls in
  ``with tvm.target.Target("ascend")``, matching the thirteen tests and the
  annotate_unlimit_memory example already updated in df459687.

Both are test-side misses from the isolation work; no library behavior changes.

Validation (Ascend NPU host, USE_CUDA=ON + USE_ASCEND=ON build):
- pytest -n 32 --maxfail=3 ascend            -> 1037 passed, 12 skipped, 0 failed
- pytest -n 4 --maxfail=3 ../examples/ascend -> 96 passed

* [Profiler] Split timing helpers and add wall-clock benchmarking (#3234)

* [Profiler] Separate runtime device backends

* [Profiler] Limit backend split to code relocation

* [Profiler][Backend] Declare profiling backend compatibility

* [Profiler] Keep do_bench in its original module

* [Profiler] Split timing helpers and add wall-clock benchmarking

* [Profiler] Share CUDA and MPS event and synchronization helpers

* [Profiler] Remove native MPS event and CUDA forwarding mock tests

(cherry picked from commit db5584f2d3)

* [CI] Include Ascend unit tests in the local CI script

* [Profiler] Move msprof timing into its own module

* [IR] Guard PersistentFor body instead of breaking

* [Runtime] Defer adapter device detection to Torch runtime

* [Backend] Align reducer finalization with upstream

* [Lint] Fix Ascend reducer namespace formatting

* [Cache] Remove Python pipeline cache stamp

* [CI] Explicitly enable Ascend in build entry points

* [CI][Ascend] Keep performance regression snapshots outside the checkout (#471)

* [CI][Ascend] Use the PR base for performance regression (#474)

* fix(ascend): preserve cross-iteration WAR dependencies on shared buffers (#466)

Condition snapshots outside the constraint Bind stack were not renamed
between iterations. Opposite branches in different iterations could
therefore be incorrectly treated as mutually exclusive, causing shared
buffer WAR dependencies to be omitted.

Seed producer and consumer substitutions with variables written in the
compared iteration scope before renaming constraints. This prevents MTE2
from overwriting shared UB before an earlier Vector read completes.

Fixes the MHC norm backward regression exposed by the dependency ordering
change in 7b580f05. Add regression coverage for cross-iteration branch
transitions.

Co-authored-by: caojian5 <caojian5@huawei.com>

* cherry-pick from #3209 to #3238 (#467)

* [BugFix][Metal] Respect GEMM buffer region offsets (#3209)

(cherry picked from commit 030556de97)
(cherry picked from commit 52456a97da0c9c4eca63ebe4d30dfb788a89e0a0)

* [Runtime] Fix library loading for symlink installs (#3227)

(cherry picked from commit 4c9cf5c7ea)
(cherry picked from commit e4b2117e9614ee549ecc617ec1af5545fb69cc4a)

* [CUDA] Fix boolean bitwise negation codegen (#3228)

* [CUDA] Fix boolean bitwise negation codegen

* [CUDA] Fix vectorized integer bitwise negation

Co-authored-by: Карим <kareemm@yandex-team.ru>

* [Test] Minimize CUDA bitwise negation regression coverage

---------

Co-authored-by: Карим <kareemm@yandex-team.ru>
(cherry picked from commit 96815790f5)
(cherry picked from commit cfe44e0e77361c658e3688a243859fe85e14c0da)

* [Example] FP8 sparse MLA forward for DeepSeek V3.2 on Hopper (#3224)

(cherry picked from commit e350730319)
(cherry picked from commit e6f4a344fc36cf992e619c8e550095fa570f2499)

* [CPU] Import the CPU dialect in CPU tests (#3236)

* [CPU] Import the CPU dialect in CPU tests

* [CPU] Move the use_tma rejection test to the dialect op-hints tests

With the CPU tests written in the CPU dialect, use_tma is not expressible
there any more (it is a CUDA-dialect keyword), so test_tilelang_cpu_atomic.py
had to import tilelang.cuda.language just for this one case. The behavior
under test - a CUDA-dialect atomic_add(use_tma=True) compiled for target="c"
is rejected by the CPU backend - is a cross-dialect property, so it now lives
next to the other cross-dialect hint tests and the CPU test file only uses
the CPU dialect. The match is also tightened to the backend error text so a
Python TypeError cannot satisfy it.

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit 80299bf1c5)
(cherry picked from commit 9da48685f2c37652ad9129f8266900e8c7b2c97e)

* [BugFix][Vectorize] Keep atomic_add scalar for invariant/non-contiguous destinations (#3219)

* [Vectorize] Keep atomic_add scalar for invariant/non-contiguous destinations

Auto-vectorization widened a scalar atomic_add to AtomicAddx2/x4 based only on
dtype and scope, ignoring whether the destination lanes form a contiguous,
aligned run. An invariant/broadcast destination (B[i//2], B[(i//2)*2], B[0])
was emitted as a contiguous wide atomic at the base address, silently
corrupting neighbouring elements, while an odd base faulted with
'CUDA error: misaligned address'.

Add CanVectorizeAtomicTarget: the vectorized loop variable must advance exactly
one element per lane in the innermost destination index and the base must be
provably aligned to the atomic width. Otherwise fall back to scalar atomics.
The check runs on both the original and the var->0 substituted destination, so
it covers address_of / tl.access_ptr / tvm_access_ptr uniformly and preserves
the trailing memory_order operand.

Testing: python -m pytest testing/python/language/test_tilelang_language_atomic.py -q
(75 passed, 9 skipped)

* [Vectorize] Validate every lane transition in atomic widening

CanVectorizeAtomicTarget only compared lanes 0 and 1, so a destination that
agrees on the first transition but repeats at larger widths (e.g. B[i % 2] at
four lanes) was not rejected by the predicate itself. Check every lane
transition up to the vector width instead: exactly one innermost index must
advance one element per lane and the others must stay constant across all lanes.

Cover B[i % 2] in the existing invariant-destination test by bounding the
emitted vector width by the destination's valid run.

* [Vectorize] Reuse the vectorizer's Ramp propagation for the atomic target check

`CanVectorizeAtomicTarget` re-derived lane contiguity by substituting every
lane value into every index and running `CanProveEqual` twice per lane
(indices x (lanes-1) x 2 Simplify+prove calls), and needed both the original
and the visited destination to recover the base.

The `TLVectorizer` already carries this information: visiting an index turns
`var_` into `Ramp(0, 1, lanes)`, Add/Sub/Mul keep the Ramp with a scaled
stride, and FloorDiv/FloorMod/Select fold into a generic vector. So the
destination is widenable iff every leading index visits to a scalar and the
innermost one visits to a unit-stride Ramp whose base is provably aligned to
the atomic width -- the same Ramp/stride test `IndicesCanVectorize` applies
when planning ordinary loads and stores.

Replace the free function with `TLVectorizer::AtomicTargetIsContiguous`,
which is one visit per index plus a single alignment proof. Behaviour is
unchanged: the PR's atomic test file passes (84), 25 probe kernels (split-K
2D fp32/fp16/bf16, region shared->global, runtime offset, `B[i//2]`,
`B[i%2]`, `B[i%4]`, `B[2*i]`, `B[0]`, odd base, 2-D constant last index)
generate byte-identical `kernel_source` against the previous revision, and
hand-written `T.vectorized` atomics keep the same widths and numerics.

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit c6ece483b1)
(cherry picked from commit 8ff7afced3bbc15dadfcde418fb708c0d751d1c6)

* [Transform][CUDA] Plan atomic vector widths from destination addresses (#3238)

* Plan atomic vector widths from destination addresses

* Avoid synthetic buffers in atomic vectorization analysis

* Assert atomic vectorization IR invariants

(cherry picked from commit e688a439a8)
(cherry picked from commit 07a76f8829dc555a461310eb6492ee8a853bf64b)

---------

Co-authored-by: Anders <anders@magnitude.dev>
Co-authored-by: Sepcnt <30561671+sepcnt@users.noreply.github.com>
Co-authored-by: Карим <kareemm@yandex-team.ru>
Co-authored-by: xuebozhang525-alt <xuebozhang525@gmail.com>
Co-authored-by: penguin_wwy <940375606@qq.com>
Co-authored-by: Alfred <166222074+Dino1844@users.noreply.github.com>

* feat(ascend): merge auto-scheduled on-chip buffers (#379)

* [Ascend] Separate buffer lifetime proofs from allocation

Generate a positive storage alias contract from automatic or manual synchronization ordering and pack compatible on-chip allocations. Keep physical lifetime footprints separate from logical synchronization regions.

Use the Ascend dialect and private lowering pass on asc-on-upstream-main, including typed AscendCopy padding and storage identity remapping.

* [Ascend][Test] Hold target scope when inspecting buffer aliases

Enter the Ascend target context in the alias-inspection helper and pass that target to lower(). This follows the caller-owned target scope contract and lets the ND2NZ and SIMT lifetime regressions run without changing the private compiler passes.

* [Ascend][Tests] Consolidate compiler and runtime regression coverage (#476)

* [Ascend][Tests] Consolidate scheduling and transform regressions

* [Ascend][Tests] Reduce application-derived regression fixtures

* [Ascend][Tests] Simplify DMA, layout, SIMD and runtime regressions

* [Ascend][Tests] Fix isolated imports and share neutral IR builders

* [Ascend][Tests] Check broadcast metadata cleanup and snapshot boundaries

* [Ascend][Tests] Document disabled latency snapshot boundary

* fix(ascend): avoid double args in debug printf (#465)

* fix(ascend): avoid double args in debug printf

- Convert double debug values to float before calling Ascend SIMT printf.
- Replace %lf with %f in scalar and buffer debug paths.
- Preserve the existing frontend double interface and block-index handling.
- Keep the fix compatible with CANN versions that reject double printf parameters.

* fix(ascend): implement float32 precision handling for double values in printf

* fix(ascend): Reject float64 debug printing

* fix(ascend): retain AICore float64 debug printing

* fix(ascend): reject unsupported AICore float64 printing

* refactor(ascend): clarify float64 print diagnostic

* refactor(ascend): simplify float64 print rejection

---------

Co-authored-by: caojian <caojian5@huawei.com>

* fix(ascend): uniquify cloned VF helper names (#488)

* fix(ascend): scalarize fp8 vector arithmetic lane by lane (#494)

* fix(ascend): scalarize fp8 vector arithmetic lane by lane

Elementwise fp8 arithmetic inside a SimtVF region produced wrong values
whenever a thread owned more than one element. For `C = A + B` over fp8 e4m3
with n=256 and 128 threads, `-128 + 64` yielded ~+16 instead of -64; only the
one-element-per-thread case (n=128) was correct.

Root cause: Ascend carries small fp8 vectors as plain integer typedefs
(`fp8_e4_2_t` is `uint16_t`, see tl_templates/ascend/ascend_fp8.h), unlike the
CUDA/HIP backends where these are structs. `PrintVecBinaryOp` therefore emitted
`*(fp8_e4_2_t *)a + *(fp8_e4_2_t *)b`, which C++ resolves as *integer*
addition of two packed byte pairs rather than as floating-point arithmetic.
The scalar path was unaffected because a lone fp8 goes through
`tl::float_e4m3_t`, whose implicit `operator float()` makes `+` a float add.

f16 and bf16 were already listed in `PrintVecBinaryOp`'s scalarization
condition for exactly this reason (packed 16-bit carriers with no matching
small-vector operation); fp8 was simply missing from the list.

Fix:
- Add e4m3/e4m3fn/e5m2 vectors to the scalarization set. e8m0fnu is excluded:
  it is a scale-factor encoding with no scalar arithmetic type, and
  `GetAscendFP8ScalarValueType` rejects it.
- Teach `PrintVecElemLoad`/`PrintVecElemStore` to address one byte per fp8 lane
  and round-trip through the existing `tl::float_e4m3_t`/`float_e5m2_t`
  conversions, so the vector path gets the same semantics the scalar path has.
  These helpers are shared with vector comparisons, Select, unary ops and calls,
  so fp8 in those contexts is corrected too; comparisons were already safe
  because `PrintBinaryExpr` dispatches on the bool result dtype.

Validation on Ascend950DT (dav-3510):
- examples/ascend/test_fp8_vecadd.py is repaired from a 0-collecting harness
  into a parametrized test over {e4m3, e5m2} x {16, 8, 4, 2, 1} elements per
  thread. All 5 e4m3 widths failed before this change and all 10 cases pass
  after it; the generated code now shows per-lane
  `tl::float_e4m3_t(...)` round-trips instead of a packed integer add.
- `pytest examples/ascend`: 106 passed (96 pre-existing + 10 new).
- `pytest testing/ascend`: 949 passed.

* refactor(ascend): reuse the shared fp8 vector predicate

Use IsAscendVectorizableFP8 at the three vector codegen sites instead of
maintaining a duplicate FP8 flavor check. Lane load/store already return
for scalars, and binary vector emission is reached only for vector values,
so the helper's extra lane-count check is redundant.

Keep the existing byte-based lane conversion and arithmetic scalarization.
Clarify that arithmetic results round to FP8 while copying an FP8 value
preserves its encoded bits. This follow-up changes only the C++ codegen
file and does not modify tests.

Validation: incremental Ascend build and formatting checks passed. All 23
focused FP8, vector-codegen, and vector-select tests passed on Ascend 950.
Generated source for 24 cases is byte-for-byte identical before and after
the simplification; five dtype lowering probes retain the same outcomes.

* feat(ascend): add MMAD traversal direction control (#495)

* [Ascend] Align L0C-to-UB fixpipe transfers (#459)

* fix(ascend): align L0C-to-UB fixpipe transfers

* fix(ascend): keep fixpipe padding bounds analyzable

* [Ascend] Add owner-exclusion dependencies to AutoSchedule (#444)

* [Ascend] Add owner-exclusion dependencies to AutoSchedule

* fix(ascend): harden owner-exclusion scheduling edge cases

Preserve owner-external fill metadata for single-version plans, reject malformed Z3 dependency tuples, improve unsupported-distance diagnostics, and share forward/reverse dependency latency calculation.

* fix(ascend): honor single-version eligibility opt-out

* [CI] Retrigger CI after rebase

* [Ascend] Clarify owner diagnostics and validate counter elision

* [Ascend] Drop extra storage diagnostics from InsertSync

* [Ascend] Warn when single-version opt-out overrides explicit eligibility

* [Ascend][AutoSchedule] Preserve iteration order at equal issue times (#512)

* [Ascend][AutoSchedule] Preserve iteration order at equal issue times

* [Ascend][AutoSchedule] Cover manual-stage issue-time ties

* [Docs][Skills] Expand Ascend kernel programming and performance guidance (#442)

* docs(skills): add TileLang review and Ascend workflows

* docs(skills): separate kernel and backend development guidance

* [Docs][Skills] Preserve existing skill names

* [Docs][Skills] Shorten Ascend auto-schedule skill name

* [Docs][Skills] Keep repository guidance focused on Ascend kernels

* [Docs][Skills] Move kernel reproducer guidance to shared skills

* [Docs][Skills] Preserve shared skill discovery descriptions

* [Docs][Skills] Clarify kernel setup and skill navigation

* Refine Ascend skill text for public use

* Document T.Stage performance tuning for Ascend kernels

* Remove redundant skill catalog

* [Skills] Consolidate core Ascend guidance into skill entrypoints

* [Skills] Integrate porting guidance into relevant sections

* [Skills] Move VF measurement into Ascend performance guidance

* [Skills] Keep semantic guidance backend-neutral

* [Skills] Merge performance guidance into tilelang-ascend

* [Skills] Align Ascend guidance with current APIs and lowering

* [Docs][Skills] Document Ascend multibuffer contracts and limits

* [Docs][Skills] Define SIMD mask in introductory example

---------

Co-authored-by: Denver Jin <denverjin@outlook.com>
Co-authored-by: mengyuhao <mengyuhao@deepseek.com>

* [Ascend] Disable Bisheng argument data cache preloading by default (#511)

* [Ascend] Disable Bisheng argument data cache preloading by default

* [Ascend] Preserve paired Bisheng options in Cython compilation

* [Ascend] Avoid forwarding explicit Cython compiler flags twice

* [JIT] Remove Cython compiler flag handoff comments

* fix(ascend): compact sparse sync event ids (#515)

Co-authored-by: caojian5 <caojian5@huawei.com>

* [Ascend] Initialize NZ transfer parameters for L0C-to-UB copies (#520)

* fix(ascend): initialize NZ parameters for L0C-to-UB copies

* test(ascend): remove L0C-to-UB NZ parameter regression test

* [Ascend] Add optimized GQA FlashAttention backward (#373)

* feat(ascend): add optimized GQA FlashAttention backward

Add an optional LSE output to the shared FlashAttention forward builder and
implement the GQA backward path as a Delta preprocessing launch followed by a
frontend-staged AutoSchedule KV-centric mixed kernel. The fused kernel
pipelines five GEMMs across AIC and AIV, accumulates dK/dV privately, and
reduces BF16 dQ through atomic adds; per-task T.Stage annotations pin the
overlapped schedule while AutoSchedule derives the buffer versioning and
inserts local and cross-core synchronization.

Add focused correctness coverage plus a cold-L2 Torch SDPA benchmark that
reports FFTS latency and pipe utilization. Document the fixed 128x128 BF16
constraints, explicit pass configuration, cannsim VF latencies, and the
measured full-shape result.

Use the tilelang.ascend.language entry point, which is the only dialect
registered on asc-on-upstream-main.

* test(ascend): track GQA backward performance

Register the GQA backward example in the Ascend perf regression script so the
fused kernel is measured on every regression run.

* perf(ascend): optimize GQA backward FixPipe stores and buffering

Improve the KV-centric GQA backward kernel while retaining T.Stage(0/1)
constraints and AutoSchedule-managed buffer indices and synchronization.
Pair UnitFlag controls on temporary L0C producers and consumers, and write
BF16 dQ contributions atomically from FixPipe to GM without the intermediate
UB tile or Vector cast.

Rebalance Q/dO L1 buffers from three to four versions and P/dS and Vector
intermediates from three to two. L1 usage remains 448 KiB while UB usage
drops from 243.5 to 162.5 KiB. Preserve the minimum of three query tiles and
the existing BF16 gradient tolerances. Extend correctness coverage to short
loops and non-power-of-two tile counts, and refresh the implementation and
performance documentation.

On Ascend950DT with CANN 9.2.0 and TileLang 8121f415, five interleaved
cold-L2 rounds reduce dQ-zeroing plus fused latency from 6834.581 us to
6418.273 us (402.19 to 428.27 effective TFLOPS). Raw FFTS records confirm
20 measured invocations per profile. Forward and Delta are outside timing.

Validation: three GQA backward cases and the existing GQA/MHA forward tests
passed using the local TileLang build; 55 additional launches across six
shapes/scales passed against FP32 autograd or Torch SDPA. Formatting for all
PR files and git diff --check passed. The compiler and runtime are unchanged.

* [Ascend][Language] Cherry-pick #3237 and load MX scale factors through L0 SF handles; retire T.copy(scale=) (#475)

* [Do not review][Op][Language][CUDA] Add common block-scaled GEMM semantics and backend dispatch (#3237)

* [Language][CUDA] Add T.gemm_blockscaled and dispatch block-scaled GEMM by target and accumulator scope

Block-scaled GEMM had two explicit entry points (T.tcgen05_gemm_blockscaled
for SM100 TCGEN5MMA, T.mma_gemm_blockscaled for SM120 mma.sync) with
duplicated bodies and no auto-dispatching tier like T.gemm. Instruction
selection also checked the SFA/SFB branch before AllowTcgen5Mma and
required SM120 there, so every 1-CTA SM100 block-scaled kernel failed with
"requires an SM120 CUDA target" (only use_2cta=True callers survived).

- SelectInst now picks the block-scaled instruction from the target and
  the accumulator scope: SM100 with C in tensor memory lowers to TCGEN05,
  SM120 with C in a fragment lowers to cuda.mma.blockscaled, anything else
  is a compile error. There is deliberately no dense fallback.
- Add T.gemm_blockscaled as the auto-dispatching entry (CUDA dialect).
  The explicit variants keep their signatures and share one
  _gemm_blockscaled_impl; T.tcgen05_gemm_blockscaled now carries
  is_tcgen05 and uses tl.tileop.tcgen05_gemm so it really pins that path.
- sf_layout rides the annotations as a StringImm; unwrap it in the SM120
  lowering.
- Tests: hardware-free call-protocol checks, an SM100 1-CTA MXFP8
  regression test through both entry points, a no-fallback negative test,
  an SM100 example test, and a gemm_api parametrization of the SM120
  NVF4 tests.

* [Op] Declare block-scaled GEMM support as a GemmImpl capability

The ROCm, Metal and CPU SelectInst implementations never look at the
SFA/SFB regions, so a block-scaled GEMM compiled for those targets was
lowered as a dense GEMM and silently dropped the scale factors.

Add GemmImpl::supports_blockscaled (true only for cuda.Gemm) and reject
block-scaled GEMMs in GemmNode::GetGemmInstructionKey before delegating
to a backend that does not declare it. Backends no longer need to
remember to reject SFA/SFB individually. Add a hardware-free test that a
CPU-bound block-scaled GEMM fails at layout inference.

* [Op] Give block-scaled GEMM its own tl.tileop.gemm_blockscaled op key

Block-scaled GEMM keeps sharing GemmNode with the dense op, but the three
frontend entry points now emit a dedicated tl.tileop.gemm_blockscaled op
(builder reuses Gemm(args, ann), mirroring tl.tileop.tcgen05_gemm) instead
of a 16-argument tl.tileop.gemm / tl.tileop.tcgen05_gemm. The printed IR is
self-describing and passes can match block-scaled GEMMs by name rather
than by counting call arguments. The explicit TCGEN05 variant is
distinguished by the is_tcgen05 annotation, as T.tcgen05_gemm is from
T.gemm. Teach the FLOP analyzer the new op name.

* [Op] Make block-scaled GEMM a first-class tile op (GemmBlockScaledNode)

tl.tileop.gemm_blockscaled now builds GemmBlockScaledNode, a GemmNode
subclass that owns the SFA/SFB regions and k_start instead of carrying
them as optional trailing slots on the dense node. Passes that only care
about "a GEMM" keep matching through GemmNode (IsInstance respects the
hierarchy); code that needs the scale factors matches the subclass via
AsGemmBlockScaled. The dense op rejects the 16-slot protocol so scale
factors can never be silently ignored.

- GemmNode: drop the SF fields, un-final the type, make Clone /
  GetGemmInstructionKey virtual, route Lower/InferLayout through
  overridable global-function names, share the 13-slot parse.
- GemmBlockScaledNode: SF fields + reflection, access regions, Clone,
  the backend capability check, tl.gemm_blockscaled.{infer_layout,lower}.
- producer_consumer_ws / cuda SelectInst read SF fields through the
  subclass; materialize_ws_schedule keeps the block-scaled op when it
  marks a tcgen05 atom instead of rewriting it to tl.tileop.tcgen05_gemm.
- Python: register tl.GemmBlockScaled as GemmBlockScaled(Gemm) with the
  matching global functions; backend impls are unchanged.

* [CUDA] Drop redundant comment on block-scaled instruction selection

* [Op][CUDA] Separate block-scaled GEMM semantics and dispatch

Move the block-scaled tile operator into dedicated common files and give CUDA instruction selection its own implementation. Reuse the GEMM backend registry through a typed block-scaled selector instead of the supports_blockscaled flag.

Reject TMEM A and incompatible explicit instruction requests in the block-scaled selector. Add CUDA-gated dispatch regression coverage while retaining shared GEMM layout and scheduling behavior.

* [Language] Expose block-scaled GEMM in the common dialect

* [Test] Remove block-scaled GEMM example tests

* [CUDA] Preserve GEMM semantics in warp-specialized schedules

* Restore explicit TCGEN05 GEMM ops in warp specialization

* Separate Python block-scaled GEMM tile-op registration

* [Language] Share dense GEMM slot construction across GEMM frontends

_gemm_impl and _gemm_blockscaled_impl each carried their own copy of the
operand legalization, region normalization, rank/shape checks, static-dim
requirement, mbar validation and 13-slot call construction. Factor that
into _gemm_dense_slots(api_name, ...) so the block-scaled frontend only
appends the SFA/SFB regions and k_start. Error messages keep their
per-entry-point prefix (T.gemm / T.gemm_blockscaled), which
examples/autodd relies on.

* [CUDA] Make T.gemm_blockscaled synchronous like T.gemm

T.gemm_blockscaled is the auto-dispatching tier of the block-scaled GEMM
family, so it should follow T.gemm's completion contract: the result is
complete when the call returns. On the SM100 TCGEN05 path the lowering
now posts completion to `mbar` and inserts the matching
mbarrier_wait_parity with the phase LowerTileOp derives from the
enclosing loop, exactly as the dense _gemm_ss lowering does; the
explicit T.tcgen05_gemm_blockscaled (is_tcgen05) still never waits.
No C++ change is needed: InjectPipeline's IsMbarPhaseConsumer and the
loop-phase plumbing already match GemmBlockScaledNode through GemmNode.

The SM100 correctness kernel now follows each entry point's contract.
The synchronous entry waits on a single per-iteration barrier and the
MMA warp releases the smem stage itself afterwards, so a missing wait
would race the producer's TMA against the in-flight MMA. A hardware-free
lowering test pins the wait (with the loop phase threaded through) for
T.gemm_blockscaled and its absence for T.tcgen05_gemm_blockscaled.

* [Language][CUDA] Move explicit CUDA GEMM variants into the CUDA dialect

wgmma_gemm, tcgen05_gemm, tcgen05_gemm_blockscaled, mma_gemm_blockscaled
and make_blockscaled_gemm_layout pin a CUDA instruction family and have
no target-neutral meaning, but were defined in the common
tilelang/language/gemm_op.py and only re-exported from
tilelang/cuda/language/intrinsics.py. Define them in
tilelang/cuda/language/gemm_op.py next to the CUDA gemm /
gemm_blockscaled shadows and import them from there. The common module
now holds only gemm, gemm_blockscaled and the shared _impl helpers, and
no longer reaches into tilelang.cuda.intrinsics for the TMEM layout
helper. Function bodies are relocated verbatim; the dialect export
surfaces are unchanged.

* [Test] Drop redundant gemm_api parametrization from SM120 NVF4 tests

T.mma_gemm_blockscaled and T.gemm_blockscaled build structurally
identical tl.tileop.gemm_blockscaled calls, so parametrizing the NVF4
codegen and correctness tests over both entry points compiled the same
IR twice on SM120 hardware. The dispatch contract for the unified entry
on SM120 is already covered hardware-free by
test_blockscaled_instruction_selection. Restore the file to main.

* [Op] Require exactly 13 positional slots for tl.tileop.gemm

tl.tileop.gemm is registered with variable arity and Gemm::Gemm only
bounded the argument count from above, so a call with fewer than 13
slots reached InitFromDenseArgs and indexed past the end of the array.
Check for exactly 13 slots up front and report the actual count; the
block-scaled hint in the message is kept for the 16-slot case. Cover
both the long and the short call in the parse test.

* [Op] Lower block-scaled GEMM through the shared tl.gemm entry points

GemmBlockScaledNode routed layout inference and lowering through its own
tl.gemm_blockscaled.{infer_layout,lower} global functions, selected by
virtual InferLayoutGlobalFunc/LowerGlobalFunc hooks on GemmNode. Those
functions were copies of tl.gemm.{infer_layout,lower}: both only call
infer_layout()/lower() on the Python object, and the object handed over
the FFI boundary is already the GemmBlockScaled Python class, so the
block-scaled behaviour comes from that class and from the instruction
key the C++ selector returns. Drop the hooks and the duplicate global
functions; a future GEMM flavour that needs different Python lowering
overrides lower() on its class instead.

* [CUDA] Split the TCGEN05 block-scaled implementation from the dense class

The Python implementation layer still told dense and block-scaled GEMM
apart by sniffing scale-factor fields: GemmBase carried SFARegion /
SFBRegion / sf_k_start / is_blockscaled via getattr defaults, and
GemmTCGEN5 branched on is_blockscaled at five points before diverting to
an 81-line _lower_blockscaled. SM120 already had the right shape, with
cuda.mma.blockscaled mapped by the registry to its own class.

Give TCGEN05 the same shape. The C++ selector returns
cuda.tcgen05.blockscaled, which the registry maps to
GemmTCGEN5BlockScaled(GemmBlockScaledMixin, GemmTCGEN5). The dense
GemmTCGEN5 exposes _warp_partition / _make_mma_emitter /
_assign_layouts hooks and a tcgen05_allow_ws flag, and no longer knows
about scale factors. GemmBlockScaledMixin
(tilelang/tileop/gemm_blockscaled/gemm_blockscaled_base.py) owns the
SFA/SFB regions, k_start and the sf_* annotation parsing that the TCGEN05
and SM120 lowerings each duplicated; which scale layouts a backend
supports stays with that backend. GemmBase drops its block-scaled
properties, matching the dense C++ node. Also drops an unused
_FLOAT8_DTYPES constant from gemm_tcgen05.py.

The instruction-selection test now also checks that each block-scaled
key resolves to a GemmBlockScaledMixin implementation with the expected
scale operands and granularities.

* [Op] Register block-scaled GEMM backends separately from GemmImpl

The block-scaled instruction selector was a slot on the dense GemmImpl
struct, so the dense GEMM header forward-declared GemmBlockScaled and
src/cuda/op/gemm.cc had to include the block-scaled selector just to
fill that slot. That is the same dependency direction the rest of this
change removed from GemmNode, GemmBase and GemmTCGEN5.

Give block-scaled GEMM its own GemmBlockScaledImpl registry in
src/op/gemm_blockscaled.{h,cc}. GemmBlockScaledNode::GetGemmInstructionKey
resolves it and, when no backend is registered for the target, still
names the dense backend that owns the target in the error. CUDA registers
itself from src/cuda/op/gemm_blockscaled.cc with the same target
predicate as its dense GemmImpl; the selector becomes internal and the
cuda/op/gemm_blockscaled.h header goes away. src/op/gemm.h now has no
block-scaled references.

Ascend adaptation for asc-on-upstream-main: T.blockscaled_gemm with
explicit sfa/sfb operands now emits tl.tileop.gemm_blockscaled (the L0
flavor without operands stays a dense gemm with the "blockscaled"
annotation, its scales living in the MX registers).
src/ascend/op/gemm_blockscaled.cc registers the block-scaled selector
("ascend.mad.blockscaled"), which the Python registry maps to
GemmMADBlockScaled (GemmBlockScaledMixin + GemmMAD); the dense GemmMAD
drops its SF-region knowledge like the CUDA classes. insert_oob_padding,
ascend_pipe and estimate_latency match both GEMM op spellings.

(cherry picked from commit d4787e9bb1)
(cherry picked from commit 11daa03f253dff19e97b45015ee46a4c2b048898)

* [Ascend] Pin the MX scale-factor slot semantics with a hardware probe

Three probes on Ascend 950DT (2026-09-18) establish the contract of the
L0A/L0B MX shadow register file that asc_copy_l12l0a_mx/asc_mmad_mx use:

- The SF destination is strictly the data tile's L0 address in 16-byte
  units: offsetting the emitted dst by one fractal (+32 in /16 units)
  yields wrong products (rel err 1.24 vs 8.8e-9), so the slot association
  has no addressing slack.
- SF slots are STICKY: after overwriting the data at the same L0 address
  with a plain (scale-less) load, a block-scaled MAD still applies the
  previously loaded scales (rel err 7.0e-9 against new-data*old-scales,
  and clearly not unscaled or cleared). Hoisting scale loads out of data
  reload loops is therefore legal.
- The SF load and the data load are order-free on the MTE1 queue:
  emitting asc_copy_l12l0a_mx before asc_copy_l12l0a is bit-identical.

Keep the stickiness probe as a regression test; it also documents the
findings for the upcoming split of the scale companion load out of
T.copy(scale=...) into standalone SF-view copies.

* [Ascend] Add L0 SF views and the standalone MX scale-factor copy

alloc_l0a_sf(a_l0)/alloc_l0b_sf(b_l0) return the MX scale-factor view of an
L0 data tile: a Buffer aliasing the tile's storage Var with SF dtype/shape,
never allocated. The aliasing is the point - the hardware keys a tile's MX
scale slots to the data tile's own address (asc_copy_l12l0a_mx dst = data
pointer / 16; asc_mmad_mx reads the slots via its A/B data pointers), so the
view inherits the tile's multi-buffer versioning, dependence edges and
address derivation for free. T.copy(sf_l1, view) emits tl.tileop.ascend_copy
with the mx_sf_load annotation (typed AscendCopyNode::is_sf_load) and lowers
to the new 8-arg tl.ascend_load_ca_sf/cb_sf intrinsics, emitting exactly the
asc_copy_l12l0a_mx/b_mx call the fused T.copy(scale=...) companion form
produces; the fused form is unchanged and coexists until its users migrate.

The SF destination pointer travels over the shared Var with its offset
converted to DATA elements: AscendLowerTileOp's access-ptr rewriting resolves
the Var to the allocation-site buffer and remaps the offset through the
tile's fractal layout exactly as it does for the data copy's dst pointer.
Only the leading (version) dims of the view region feed that pointer; the
x/y slot coordinates stay source-driven, as in the fused form. The multi-
buffer expansion expresses each alias's leading stride as the allocation
pitch in its own elements, so the version byte offsets agree by construction.

Aliased partial-footprint views need three pass exemptions, all keyed on the
same criterion (total storage bits differ => not a data alias):
- layout_inference: alias layout propagation reshapes by the bit ratio,
  which requires equal footprints; skip partial views at all three
  propagation sites instead of failing Layout::Reshape.
- normalize_fractal_storage: SF views are not canonical Cube data views;
  their accesses stay in SF coordinates for the scale-load lowering.
- auto_schedule storage sizing: size a storage Var by its largest alias, so
  a tiny SF view seen first can never undersize the L0 allocation.

Verified on Ascend 950DT: single-tile and pipelined (NUM_STAGES=2, sliced SF
sources) kernels produce bit-identical results to the fused form, and the
standalone _mx call carries the same stage offset expression as the data
load. testing/ascend regression batch (mx tail fill, blockscaled matrix,
alignment guards, estimate latency, multibuffer storage/counter): 247
passed; examples blockscaled/padding suites: 25 passed.

* [Ascend] Route L0 block-scaled GEMM through SF views and delete T.copy(scale=)

T.blockscaled_gemm now requires sfa/sfb in both flavors: L1 inputs keep
passing the L1 scale buffers, and L0 inputs pass the MX scale-factor views
of the operand tiles (alloc_l0a_sf/alloc_l0b_sf), loaded by standalone
T.copy(sf_l1, view) scale loads. Every Ascend block-scaled GEMM is now a
structural tl.tileop.gemm_blockscaled: the "blockscaled" annotation flavor
dies, GemmMAD is purely dense (is_blockscaled False; _lower_l0_blockscaled
moves to GemmMADBlockScaled, whose infer_layout assigns the SF_K layout to
L1-scoped scale operands only), and insert_oob_padding matches block-scaled
GEMMs by node type alone.

The fused scale-companion form is deleted end to end: the T.copy scale=
kwarg and its 3-region wire format, AscendCopyNode::sf/sf_range (+
reflection), the 16-arg ascend_load_cbuf_to_ca/cb protocol and its codegen
branch, the op.sf-driven SF layout assignment, and the k_align=64 L1-extent
check in copy InferLayout. The pinned "divisible by 64" diagnostic stays
covered: GemmMAD's operation-level K%64 assert fires for misaligned GEMMs,
and alloc_l0a_sf's default shape derivation rejects a misaligned K at trace
time with the same wording.

Pass migration:
- insert_oob_padding: MxL1KAxisCollector keys MX tail-fill propagation on
  "L1->L0 data copy whose dst Var feeds a block-scaled GEMM" and skips
  standalone SF loads (their dst Var IS the data Var, but their source is
  the differently laid out L1 scale buffer, not MX data).
- estimate_latency: SF loads are excluded from the copy cost features,
  keeping the accounting identical to the fused era when the scale bytes
  rode the data copy invisibly.
- dependency analysis: Var-shared aliases with differing bit footprints are
  disjoint storage (shared SameAliasFootprint in ascend/op/utils.h, also
  adopted by normalize_fractal_storage). This removes the false WAW between
  a data load and its SF load that materialized an asc_sync_pipe(PIPE_MTE1)
  between the two MTE1 instructions; ordering against the consuming MAD
  still flows through the view-to-view region pair, since the GEMM reads
  the same view buffer the SF copy writes. Equal-footprint aliases
  (T.view/T.reshape reinterprets) keep the conservative conflict.

All scale= users migrate to the split form: example_blockscaled_gemm_l0
(pipelined, uint16/uint8 scales), test_blockscaled_gemm,
test_gemm_padding_irregular_shapes (transpose + uint8 scales), the
blockscaled matrix/alignment/tail-fill tests and the MX slot probes; the
dialect-surface test drops the scale kwarg from the expected signature.

Verified on Ascend 950DT: testing/ascend 1040 passed / 12 skipped
(including the tail-fill geometry goldens and their no-PIPE_MTE1-sync
guard); examples blockscaled/padding/padded-copy suites 25 passed.

* [Ascend][Language] Rename blockscaled_gemm to gemm_blockscaled

The Ascend dialect's block-scaled GEMM entry kept a legacy spelling; now
that it emits the first-class tl.tileop.gemm_blockscaled op, name it after
the common surface it shadows (upstream #3237's T.gemm_blockscaled),
following the dialect-shadowing pattern already used for copy/gemm/unroll.
The SFA/SFB parameter names align with the common signature; the Ascend
shadow still differs deliberately in that k_start and the sf_*_granularity
knobs are implicit (the lowering derives the scale K offset from the SFA
region slice, and MX scales cover 32 K elements per factor).

The shadow is registered in the dialect-surface test's owned-symbols map;
all call sites in tests and examples migrate. No compatibility alias is
kept.

* [Ascend] Give MX scale-factor handles a dedicated scope

alloc_l0a_sf/alloc_l0b_sf now return a buffer in its own storage scope
(shared.l0a.sf / shared.l0b.sf) with its own Var, instead of a
partial-footprint alias of the data tile. The binding to the tile rides the
scale-load op as the mx_sf_of annotation (a Var, attached by the dialect
copy from the alloc-time registry); the scope itself classifies the op
everywhere else, so the mx_sf_load annotation, AscendCopyNode::is_sf_load
and the frontend WeakSet die.

The alias-era footprint heuristic is deleted wholesale: SameAliasFootprint
and its exemptions in layout inference (three propagation sites), fractal
storage normalization, and dependency analysis (the conservative
different-buffer rule is restored), plus the storage-sizing representative
workaround in AutoSchedule - with separate Vars none of those cases can
arise. Scope-keyed classification replaces the op-marker in the tail-fill
collector, the latency features and the pipe mask (L1->l0*.sf is MTE1).

Version lockstep, the one thing aliasing provided structurally, is now
explicit: GetOnChipStorages includes SF handles so the multi-buffer
eligibility planner claims an owner loop for them, and AutoSchedule copies
the bound tile's selected version count onto the handle after selection
(PropagateMxSfHandleVersions), so ResolveCore/InsertSync/Materialize ring
handle and tile together. The scale-load lowering resolves the bound tile
by Var, computes the stage offset as the handle's leading version index
times the tile's (padded) allocation footprint, and emits the destination
pointer over the tile's Var exactly like the data copy; handle allocations
are dropped in AscendLowerTileOp before storage planning ever sees the
scope. During bring-up the unclassified scope fell through to the
Vector-pipe default and the scheduler happily hoisted the scale loads past
the MAD - the pipe-mask entry is load-bearing.

Verified on Ascend 950DT: SF view/slot tests including the pipelined
stage-offset identity, the 254-test regression batch (tail-fill geometry
goldens, blockscaled matrix, alignment guards, dialect surface, latency
goldens, multibuffer suites, cross-iteration snapshots) and the 24
blockscaled/padding example tests all pass; the scale-load scheduling
microbenchmark reproduces the alias design's numbers exactly
(14.28 / 14.29 / 9.03 us for pinned per-iter / pinned hoisted / default
ping-pong).

* [Ascend] Derive MX scale-factor bindings from their consuming GEMMs

The binding between a scale-factor handle and its data tile was carried by
a frontend WeakKeyDictionary (alloc_l0a_sf recorded it, the dialect copy
attached it to the scale-load op as mx_sf_of). But the binding is already
structural in the IR: the tl.tileop.gemm_blockscaled call that consumes a
handle as SFA/SFB names the tile as its A/B operand. Derive it from there
instead (CollectMxSfBindings): AutoSchedule's version propagation reads the
GEMMs directly, and AscendLowerTileOp resolves the binding onto the parsed
copy node before lowering. A handle consumed with two different tiles is
now a diagnosed error, and a scale load whose handle feeds no GEMM fails
with a message naming the missing consumer.

This deletes the registry, the frontend-attached annotation, and the
dialect copy's special branch entirely - a scale load is emitted through
the ordinary copy path and classified purely by its destination scope, so
hand-written or synthesized IR needs no frontend cooperation at all.
alloc_l0a_sf's data-tile argument now only drives the default shape
derivation and a scope sanity check.

Verified on Ascend 950DT: SF view/slot/tail-fill, blockscaled matrix,
alignment guards and dialect-surface tests (53 passed) plus the 24
blockscaled/padding example tests.

* [Ascend] Polish MX scale-factor handle plumbing after review

- Delete the dead "mx_sf_of" annotation decode: nothing emits the
  annotation; AscendLowerTileOp resolves the binding structurally from
  the consuming gemm_blockscaled and sets the field directly.
- Inject that binding through CopyOnWrite (AscendCopy now declares the
  COW method) instead of writing through a const_cast.
- Align the scale-load empty-copy guard with the DMA path (CanProve
  short-circuit), explain the single-version-dimension limit in the
  handle-rank check, and refresh stale fused-copy-era comments.
- alloc_l0a_sf/alloc_l0b_sf: refuse the default sf_shape derivation for
  tiles carrying leading version dims and document that `buf` is
  advisory - the consuming gemm defines the binding.

* [Ascend][Tests] Drop the SF-view codegen regression file

The MX scale-factor handle path keeps its numerical coverage through the
pipelined block-scaled examples (a wrong stage offset or missing scale
load produces wrong products there) and the slot-semantics probe; the
per-instruction codegen pins are not worth a dedicated file.

* [Ascend] Reject MX scale-factor slot plans that ring out of lockstep

From PR475 review: with the scale load in an outer pipelined loop and
the data reload in an inner one, both storages ring two versions but the
scale load's slot follows the outer loop while the tile and its
consuming MAD follow the inner one, so the MAD reads slots whose scales
were never loaded - silent wrong results on 950DT (the hardware keys the
scale slots to the tile's address). PropagateMxSfHandleVersions
equalizes only the version counts; the version index is driven by each
storage's own owner loop.

Validate in the multi-buffer plan builder, where the owners are known: a
handle ringing more than one version must share its data tile's owner
loops. The diagnostic names the two workarounds (load the scales in the
same pipelined loop as the data tile, or pin the tile to a single
version and let the sticky slots serve every reload) and also covers the
fully hoisted form, which previously died on the generic "no annotated
owner loop" ICHECK. Broadcasting hoisted scale loads into every version
slot instead of rejecting them is left as a TODO.

* [Ascend] Broadcast hoisted MX scale loads into every version slot

Supports the schedule shape the previous commit rejected: a scale load
hoisted outside its data tile's pipelined loop (into an outer pipelined
loop, or above every loop) while the tile keeps ping-ponging.

Plan the handle from its bound tile instead of independently.
PropagateMxSfHandleVersions compares owner loops (CollectMultiBufferOwners)
and copies the tile's version count only when the scale load rings under
the same owners - the lockstep behavior, unchanged. Otherwise the handle
stays single-version and the copy lowering broadcasts the load:
AscendLowerTileOp injects the tile's slot count (mx_sf_versions) next to
the mx_sf_of binding, and LowerMxSfLoad emits one asc_copy_l12l0*_mx per
version slot. The scales are invariant across the tile's ring period,
slots are sticky, and the plain storage dependence on the unversioned
handle orders each broadcast after every consumer of the previous
scales, so the inner data ping-pong survives intact and no sync rules
change.

With unsupported plans no longer constructible, the plan-time rejection
(ValidateMxSfHandleOwners) is removed and its regression file now pins
the supported behavior instead: broadcast emission for both hoisted
forms, the pinned single-slot form, and the review reproducer's numerics
on 950DT (rel 5.5e-08; the broadcast pair is guarded by the
PIPE_M->PIPE_MTE1 wait on the previous iteration's last MAD).

* [Docs][Ascend] Announce Ascend 950 backend support (#519)

* [Docs][Ascend] Announce Ascend 950 backend support

* Update README.md

---------

Co-authored-by: xuruifan <silentcoder@foxmail.com>

* [Ascend][Cache] Publish immutable binary cache directories (#524)

Port the immutable directory scheme from the CUDA binary cache
(tile-ai/tilelang#3177) to AscendBinaryCache, which still published
flat mutable files with no integrity verification: a crash could
publish a truncated .aibin that load() would then serve forever, and
os.replace let later writers overwrite entries that concurrent readers
may hold open.

Factor the shared behavior into a BinaryCache base
(tilelang/cache/binary_cache.py): each key is an immutable directory
holding metadata.json (format tag, size, sha256) and the raw binary,
staged privately, fsynced, and published with a single rename so the
first valid writer wins. Any anomaly on load is a plain cache miss and
never deletes shared entries. CUDABinaryCache and AscendBinaryCache now
only define their cache roots, format tags, and make_key.

* [Docs][Ascend] Clarify kernel programming responsibilities (#527)

* [Docs][Ascend] Clarify kernel programming responsibilities

* [Docs][Ascend] Clarify programming comparison headings

* [Ascend] Bind L0 scale factors at allocation and unify storage planning (#529)

* [Ascend] Bind L0 scale factors at allocation and unify storage planning

* [Ascend] Preserve SF group owners and prove copy capacity

* [Ascend] Reject only proven SF copy overflows

* [Ascend] Require positive proof of SF copy capacity

* [Ascend] Sync upstream through tile-ai/tilelang#3284 (#533)

* [BugFix][Metal] Respect GEMM buffer region offsets (#3209)

(cherry picked from commit 030556de97)

* [Runtime] Fix library loading for symlink installs (#3227)

(cherry picked from commit 4c9cf5c7ea)

* [CUDA] Fix boolean bitwise negation codegen (#3228)

* [CUDA] Fix boolean bitwise negation codegen

* [CUDA] Fix vectorized integer bitwise negation

Co-authored-by: Карим <kareemm@yandex-team.ru>

* [Test] Minimize CUDA bitwise negation regression coverage

---------

Co-authored-by: Карим <kareemm@yandex-team.ru>
(cherry picked from commit 96815790f5)

* [Example] FP8 sparse MLA forward for DeepSeek V3.2 on Hopper (#3224)

(cherry picked from commit e350730319)

* [CPU] Import the CPU dialect in CPU tests (#3236)

* [CPU] Import the CPU dialect in CPU tests

* [CPU] Move the use_tma rejection test to the dialect op-hints tests

With the CPU tests written in the CPU dialect, use_tma is not expressible
there any more (it is a CUDA-dialect keyword), so test_tilelang_cpu_atomic.py
had to import tilelang.cuda.language just for this one case. The behavior
under test - a CUDA-dialect atomic_add(use_tma=True) compiled for target="c"
is rejected by the CPU backend - is a cross-dialect property, so it now lives
next to the other cross-dialect hint tests and the CPU test file only uses
the CPU dialect. The match is also tightened to the backend error text so a
Python TypeError cannot satisfy it.

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit 80299bf1c5)

* [BugFix][Vectorize] Keep atomic_add scalar for invariant/non-contiguous destinations (#3219)

* [Vectorize] Keep atomic_add scalar for invariant/non-contiguous destinations

Auto-vectorization widened a scalar atomic_add to AtomicAddx2/x4 based only on
dtype and scope, ignoring whether the destination lanes form a contiguous,
aligned run. An invariant/broadcast destination (B[i//2], B[(i//2)*2], B[0])
was emitted as a contiguous wide atomic at the base address, silently
corrupting neighbouring elements, while an odd base faulted with
'CUDA error: misaligned address'.

Add CanVectorizeAtomicTarget: the vectorized loop variable must advance exactly
one element per lane in the innermost destination index and the base must be
provably aligned to the atomic width. Otherwise fall back to scalar atomics.
The check runs on both the original and the var->0 substituted destination, so
it covers address_of / tl.access_ptr / tvm_access_ptr uniformly and preserves
the trailing memory_order operand.

Testing: python -m pytest testing/python/language/test_tilelang_language_atomic.py -q
(75 passed, 9 skipped)

* [Vectorize] Validate every lane transition in atomic widening

CanVectorizeAtomicTarget only compared lanes 0 and 1, so a destination that
agrees on the first transition but repeats at larger widths (e.g. B[i % 2] at
four lanes) was not rejected by the predicate itself. Check every lane
transition up to the vector width instead: exactly one innermost index must
advance one element per lane and the others must stay constant across all lanes.

Cover B[i % 2] in the existing invariant-destination test by bounding the
emitted vector width by the destination's valid run.

* [Vectorize] Reuse the vectorizer's Ramp propagation for the atomic target check

`CanVectorizeAtomicTarget` re-derived lane contiguity by substituting every
lane value into every index and running `CanProveEqual` twice per lane
(indices x (lanes-1) x 2 Simplify+prove calls), and needed both the original
and the visited destination to recover the base.

The `TLVectorizer` already carries this information: visiting an index turns
`var_` into `Ramp(0, 1, lanes)`, Add/Sub/Mul keep the Ramp with a scaled
stride, and FloorDiv/FloorMod/Select fold into a generic vector. So the
destination is widenable iff every leading index visits to a scalar and the
innermost one visits to a unit-stride Ramp whose base is provably aligned to
the atomic width -- the same Ramp/stride test `IndicesCanVectorize` applies
when planning ordinary loads and stores.

Replace the free function with `TLVectorizer::AtomicTargetIsContiguous`,
which is one visit per index plus a single alignment proof. Behaviour is
unchanged: the PR's atomic test file passes (84), 25 probe kernels (split-K
2D fp32/fp16/bf16, region shared->global, runtime offset, `B[i//2]`,
`B[i%2]`, `B[i%4]`, `B[2*i]`, `B[0]`, odd base, 2-D constant last index)
generate byte-identical `kernel_source` against the previous revision, and
hand-written `T.vectorized` atomics keep the same widths and numerics.

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit c6ece483b1)

* [Transform][CUDA] Plan atomic vector widths from destination addresses (#3238)

* Plan atomic vector widths from destination addresses

* Avoid synthetic buffers in atomic vectorization analysis

* Assert atomic vectorization IR invariants

(cherry picked from commit e688a439a8)

* [Do not review][Op][Language][CUDA] Add common block-scaled GEMM semantics and backend dispatch (#3237)

* [Language][CUDA] Add T.gemm_blockscaled and dispatch block-scaled GEMM by target and accumulator scope

Block-scaled GEMM had two explicit entry points (T.tcgen05_gemm_blockscaled
for SM100 TCGEN5MMA, T.mma_gemm_blockscaled for SM120 mma.sync) with
duplicated bodies and no auto-dispatching tier like T.gemm. Instruction
selection also checked the SFA/SFB branch before AllowTcgen5Mma and
required SM120 there, so every 1-CTA SM100 block-scaled kernel failed with
"requires an SM120 CUDA target" (only use_2cta=True callers survived).

- SelectInst now picks the block-scaled instruction from the target and
  the accumulator scope: SM100 with C in tensor memory lowers to TCGEN05,
  SM120 with C in a fragment lowers to cuda.mma.blockscaled, anything else
  is a compile error. There is deliberately no dense fallback.
- Add T.gemm_blockscaled as the auto-dispatching entry (CUDA dialect).
  The explicit variants keep their signatures and share one
  _gemm_blockscaled_impl; T.tcgen05_gemm_blockscaled now carries
  is_tcgen05 and uses tl.tileop.tcgen05_gemm so it really pins that path.
- sf_layout rides the annotations as a StringImm; unwrap it in the SM120
  lowering.
- Tests: hardware-free call-protocol checks, an SM100 1-CTA MXFP8
  regression test through both entry points, a no-fallback negative test,
  an SM100 example test, and a gemm_api parametrization of the SM120
  NVF4 tests.

* [Op] Declare block-scaled GEMM support as a GemmImpl capability

The ROCm, Metal and CPU SelectInst implementations never look at the
SFA/SFB regions, so a block-scaled GEMM compiled for those targets was
lowered as a dense GEMM and silently dropped the scale factors.

Add GemmImpl::supports_blockscaled (true only for cuda.Gemm) and reject
block-scaled GEMMs in GemmNode::GetGemmInstructionKey before delegating
to a backend that does not declare it. Backends no longer need to
remember to reject SFA/SFB individually. Add a hardware-free test that a
CPU-bound block-scaled GEMM fails at layout inference.

* [Op] Give block-scaled GEMM its own tl.tileop.gemm_blockscaled op key

Block-scaled GEMM keeps sharing GemmNode with the dense op, but the three
frontend entry points now emit a dedicated tl.tileop.gemm_blockscaled op
(builder reuses Gemm(args, ann), mirroring tl.tileop.tcgen05_gemm) instead
of a 16-argument tl.tileop.gemm / tl.tileop.tcgen05_gemm. The printed IR is
self-describing and passes can match block-scaled GEMMs by name rather
than by counting call arguments. The explicit TCGEN05 variant is
distinguished by the is_tcgen05 annotation, as T.tcgen05_gemm is from
T.gemm. Teach the FLOP analyzer the new op name.

* [Op] Make block-scaled GEMM a first-class tile op (GemmBlockScaledNode)

tl.tileop.gemm_blockscaled now builds GemmBlockScaledNode, a GemmNode
subclass that owns the SFA/SFB regions and k_start instead of carrying
them as optional trailing slots on the dense node. Passes that only care
about "a GEMM" keep matching through GemmNode (IsInstance respects the
hierarchy); code that needs the scale factors matches the subclass via
AsGemmBlockScaled. The dense op rejects the 16-slot protocol so scale
factors can never be silently ignored.

- GemmNode: drop the SF fields, un-final the type, make Clone /
  GetGemmInstructionKey virtual, route Lower/InferLayout through
  overridable global-function names, share the 13-slot parse.
- GemmBlockScaledNode: SF fields + reflection, access regions, Clone,
  the backend capability check, tl.gemm_blockscaled.{infer_layout,lower}.
- producer_consumer_ws / cuda SelectInst read SF fields through the
  subclass; materialize_ws_schedule keeps the block-scaled op when it
  marks a tcgen05 atom instead of rewriting it to tl.tileop.tcgen05_gemm.
- Python: register tl.GemmBlockScaled as GemmBlockScaled(Gemm) with the
  matching global functions; backend impls are unchanged.

* [CUDA] Drop redundant comment on block-scaled instruction selection

* [Op][CUDA] Separate block-scaled GEMM semantics and dispatch

Move the block-scaled tile operator into dedicated common files and give CUDA instruction selection its own implementation. Reuse the GEMM backend registry through a typed block-scaled selector instead of the supports_blockscaled flag.

Reject TMEM A and incompatible explicit instruction requests in the block-scaled selector. Add CUDA-gated dispatch regression coverage while retaining shared GEMM layout and scheduling behavior.

* [Language] Expose block-scaled GEMM in the common dialect

* [Test] Remove block-scaled GEMM example tests

* [CUDA] Preserve GEMM semantics in warp-specialized schedules

* Restore explicit TCGEN05 GEMM ops in warp specialization

* Separate Python block-scaled GEMM tile-op registration

* [Language] Share dense GEMM slot construction across GEMM frontends

_gemm_impl and _gemm_blockscaled_impl each carried their own copy of the
operand legalization, region normalization, rank/shape checks, static-dim
requirement, mbar validation and 13-slot call construction. Factor that
into _gemm_dense_slots(api_name, ...) so the block-scaled frontend only
appends the SFA/SFB regions and k_start. Error messages keep their
per-entry-point prefix (T.gemm / T.gemm_blockscaled), which
examples/autodd relies on.

* [CUDA] Make T.gemm_blockscaled synchronous like T.gemm

T.gemm_blockscaled is the auto-dispatching tier of the block-scaled GEMM
family, so it should follow T.gemm's completion contract: the result is
complete when the call returns. On the SM100 TCGEN05 path the lowering
now posts completion to `mbar` and inserts the matching
mbarrier_wait_parity with the phase LowerTileOp derives from the
enclosing loop, exactly as the dense _gemm_ss lowering does; the
explicit T.tcgen05_gemm_blockscaled (is_tcgen05) still never waits.
No C++ change is needed: InjectPipeline's IsMbarPhaseConsumer and the
loop-phase plumbing already match GemmBlockScaledNode through GemmNode.

The SM100 correctness kernel now follows each entry point's contract.
The synchronous entry waits on a single per-iteration barrier and the
MMA warp releases the smem stage itself afterwards, so a missing wait
would race the producer's TMA against the in-flight MMA. A hardware-free
lowering test pins the wait (with the loop phase threaded through) for
T.gemm_blockscaled and its absence for T.tcgen05_gemm_blockscaled.

* [Language][CUDA] Move explicit CUDA GEMM variants into the CUDA dialect

wgmma_gemm, tcgen05_gemm, tcgen05_gemm_blockscaled, mma_gemm_blockscaled
and make_blockscaled_gemm_layout pin a CUDA instruction family and have
no target-neutral meaning, but were defined in the common
tilelang/language/gemm_op.py and only re-exported from
tilelang/cuda/language/intrinsics.py. Define them in
tilelang/cuda/language/gemm_op.py next to the CUDA gemm /
gemm_blockscaled shadows and import them from there. The common module
now holds only gemm, gemm_blockscaled and the shared _impl helpers, and
no longer reaches into tilelang.cuda.intrinsics for the TMEM layout
helper. Function bodies are relocated verbatim; the dialect export
surfaces are unchanged.

* [Test] Drop redundant gemm_api parametrization from SM120 NVF4 tests

T.mma_gemm_blockscaled and T.gemm_blockscaled build structurally
identical tl.tileop.gemm_blockscaled calls, so parametrizing the NVF4
codegen and correctness tests over both entry points compiled the same
IR twice on SM120 hardware. The dispatch contract for the unified entry
on SM120 is already covered hardware-free by
test_blockscaled_instruction_selection. Restore the file to main.

* [Op] Require exactly 13 positional slots for tl.tileop.gemm

tl.tileop.gemm is registered with variable arity and Gemm::Gemm only
bounded the argument count from above, so a call with fewer than 13
slots reached InitFromDenseArgs and indexed past the end of the array.
Check for exactly 13 slots up front and report the actual count; the
block-scaled hint in the message is kept for the 16-slot case. Cover
both the long and the short call in the parse test.

* [Op] Lower block-scaled GEMM through the shared tl.gemm entry points

GemmBlockScaledNode routed layout inference and lowering through its own
tl.gemm_blockscaled.{infer_layout,lower} global functions, selected by
virtual InferLayoutGlobalFunc/LowerGlobalFunc hooks on GemmNode. Those
functions were copies of tl.gemm.{infer_layout,lower}: both only call
infer_layout()/lower() on the Python object, and the object handed over
the FFI boundary is already the GemmBlockScaled Python class, so the
block-scaled behaviour comes from that class and from the instruction
key the C++ selector returns. Drop the hooks and the duplicate global
functions; a future GEMM flavour that needs different Python lowering
overrides lower() on its class instead.

* [CUDA] Split the TCGEN05 block-scaled implementation from the dense class

The Python implementation layer still told dense and block-scaled GEMM
apart by sniffing scale-factor fields: GemmBase carried SFARegion /
SFBRegion / sf_k_start / is_blockscaled via getattr defaults, and
GemmTCGEN5 branched on is_blockscaled at five points before diverting to
an 81-line _lower_blockscaled. SM120 already had the right shape, with
cuda.mma.blockscaled mapped by the registry to its own class.

Give TCGEN05 the same shape. The C++ selector returns
cuda.tcgen05.blockscaled, which the registry maps to
GemmTCGEN5BlockScaled(GemmBlockScaledMixin, GemmTCGEN5). The dense
GemmTCGEN5 exposes _warp_partition / _make_mma_emitter /
_assign_layouts hooks and a tcgen05_allow_ws flag, and no longer knows
about scale factors. GemmBlockScaledMixin
(tilelang/tileop/gemm_blockscaled/gemm_blockscaled_base.py) owns the
SFA/SFB regions, k_start and the sf_* annotation parsing that the TCGEN05
and SM120 lowerings each duplicated; which scale layouts a backend
supports stays with that backend. GemmBase drops its block-scaled
properties, matching the dense C++ node. Also drops an unused
_FLOAT8_DTYPES constant from gemm_tcgen05.py.

The instruction-selection test now also checks that each block-scaled
key resolves to a GemmBlockScaledMixin implementation with the expected
scale operands and granularities.

* [Op] Register block-scaled GEMM backends separately from GemmImpl

The block-scaled instruction selector was a slot on the dense GemmImpl
struct, so the dense GEMM header forward-declared GemmBlockScaled and
src/cuda/op/gemm.cc had to include the block-scaled selector just to
fill that slot. That is the same dependency direction the rest of this
change removed from GemmNode, GemmBase and GemmTCGEN5.

Give block-scaled GEMM its own GemmBlockScaledImpl registry in
src/op/gemm_blockscaled.{h,cc}. GemmBlockScaledNode::GetGemmInstructionKey
resolves it and, when no backend is registered for the target, still
names the dense backend that owns the target in the error. CUDA registers
itself from src/cuda/op/gemm_blockscaled.cc with the same target
predicate as its dense GemmImpl; the selector becomes internal and the
cuda/op/gemm_blockscaled.h header goes away. src/op/gemm.h now has no
block-scaled references.

(cherry picked from commit d4787e9bb1)

* [JIT] Add missing uint64 argument type mappings (#3229)

(cherry picked from commit e40e09e3e6)

* [BugFix] Restore symbolic loop-layout injectivity proof; reject layouts on symbolic shared tiles (#3233)

* [BugFix] Support symbolic shared layouts with mixed fixed and dynamic extents

Related to tile-ai/tilelang#2906 and tile-ai/tilelang#2909. Credit Soham Panda for the original diagnosis and proposed defensive fix.

Co-authored-by: Soham Panda <sohampanda1@gmail.com>

* [Test] Minimize symbolic shared-layout regression coverage

* [Layout] Document why both symbolic injectivity proofs are kept

CanProveInjective and CanProveLeftInverse have disjoint blind spots:
equality propagation proves (i, i + j) but cannot invert floordiv/floormod
with a symbolic divisor, while the left-inverse decoder handles the padded
partition ((i*n+j)%128, (i*n+j)//128) but has no candidate for (i, i + j).
Record the counterexamples at the call site.

* [BugFix] Reject layouts on symbolic shared tiles instead of remapping them

makeBufferWithLayout dereferenced a null IntImm for a shared buffer with a
symbolic extent (#2906). Rather than deriving a symbolic replication factor,
keep the constant arithmetic and raise a ValueError naming the buffer and the
offending extent. The check sits at the remap site so that layouts inferred
by ops (scan) are covered as well as T.annotate_layout.

The left-inverse injectivity proof stays: it fixes T.Parallel loops over a
mixed static/symbolic space, which #2719 broke independently of shared
layouts. Its regression test is now a pure global-to-global loop, and the
#2906 test asserts the error for both the annotated and the inferred path.

---------

Co-authored-by: Soham Panda <sohampanda1@gmail.com>
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit efe9654bf4)

* [ROCm] Run portable example validation in CI (#3165)

* [ROCm] Run portable example validation in CI

Signed-off-by: andyluo7 <andy.luo@amd.com>

* [ROCm] Make portable example checks fail reliably

Signed-off-by: andyluo7 <andy.luo@amd.com>

---------

Signed-off-by: andyluo7 <andy.luo@amd.com>
(cherry picked from commit 48cd23e181)

* [BugFix][CUDA] Only emit 256-bit global load/store on SM100+ targets (#3248)

(cherry picked from commit 4bc6c32a63)

* [Metal] Support 32-bit integer atomic add (#3211)

(cherry picked from commit 901f941c0d)

* [Testing] Drop duplicated codegen-smoke tests; fix fastmath self-comparison assertions (#3252)

(cherry picked from commit 054bcf1de3)

* [ROCm] Add GLM-5.3 k-pool Top-K transform (#3254)

* [ROCm] Add GLM-5.3 k-pool cache writer

Signed-off-by: andyluo7 <andy.luo@amd.com>

* [ROCm] Match GLM-5.3 k-pool geometry

Signed-off-by: andyluo7 <andy.luo@amd.com>

* [ROCm] Add GLM-5.3 k-pool decode tail

Signed-off-by: andyluo7 <andy.luo@amd.com>

* [ROCm] Bind GLM-5.3 tail boundary count

Signed-off-by: andyluo7 <andy.luo@amd.com>

* [ROCm] Add GLM-5.3 paged k-pool logits

Signed-off-by: andyluo7 <andy.luo@amd.com>

* [ROCm] Guard GLM-5.3 k-pool page bounds

Signed-off-by: andyluo7 <andy.luo@amd.com>

* [ROCm] Add GLM-5.3 k-pool Top-K transform

Signed-off-by: andyluo7 <andy.luo@amd.com>

---------

Signed-off-by: andyluo7 <andy.luo@amd.com>
(cherry picked from commit eab74a4ae5)

* [Frontend] Support Python iterables and comprehensions (#3230)

* [Frontend] Support Python iterables and comprehensions

* [Frontend] Avoid excessive nesting in generated loop code

* [Frontend] Reject TIR values and non-iterables in compile-time for loops

The Python-iterable path in ctx_for consumed anything that was not a loop
frame. Var carries an __iter__ shim for single-binding unpacking, so
`for i in n` with a symbolic n expanded once with i aliased to n and renamed
the shape variable; Buffer.__getitem__ never raises IndexError, so
`for x in A` looped forever. Reject PrimExpr and Buffer up front and keep
the original diagnostic for objects that are not iterable at all.

* [Frontend] Validate Python iterable sources and comprehensions

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit 6ba187e20d)

* [Frontend] Remove warnings for immutable variable rebinding (#3262)

Remove warnings for immutable variable rebinding

(cherry picked from commit 195b6ebbf3)

* [BugFix] Reject non-power-of-two AllReduce thread strides (#3266)

The XOR butterfly in tl::AllReduce pairs thread t with t ^ (scale * k), which
only moves the reduce coordinate when `scale` (the thread stride between
consecutive reduce participants) is a power of two. #2611 added the check for
the logical width (threads / scale) but not for the stride, so reducing a
(32, 3) fragment along dim 0 with its default inferred layout lowered to
AllReduce<MaxOp, 96, 3> and silently mixed columns (and replicas); with exactly
96 or 192 threads the shared-memory exchange also read past the workspace.

- CheckAllReduceWidth, shared by tl.reduce and tl.finalize_reducer, now
  requires the stride to be a power of two and points the user at padding the
  non-reduced fragment extent.
- The CUDA AllReduce template carries the matching static_assert.
- The reducer v2 narrow-plan gate rejects such strides alongside the existing
  width check, so the planner falls back to the wide plan instead of emitting
  the broken collective.

(cherry picked from commit 1b908fc0eb)

* [NVRTC] Fix warp reduction compilation (#3260)

Fixes https://github.com/tile-ai/tilelang/issues/3259

(cherry picked from commit 261a9e4bb1)

* [BugFix][CUDA] Lower a kernel-body assert to a device-legal check (#3206)

* [BugFix][CUDA] Lower a kernel-body assert to a device-legal check

A plain assert in a kernel body is parsed into a tirx.AssertStmt. The CUDA
codegen did not override that visitor, so it inherited CodeGenC's, which
streams a host-only TVMFFIErrorSetRaisedFromCStrParts call plus `return -1`
into the emitted __global__ kernel. nvcc then rejects the file with
"identifier TVMFFIErrorSetRaisedFromCStrParts is undefined", so a valid
kernel refused to build.

Override VisitStmt_(AssertStmtNode) in CodeGenTileLangCUDA and lower to the
device helpers the T.device_assert intrinsic path already uses
(device_assert / device_assert_with_msg). Every function this codegen emits
is a __global__ kernel, so the override cannot misroute a host-side assert;
host and CPU code have their own codegens with their own visitors.

Fixes #3019

* [Test][CUDA] Gate kernel-body assert regressions on CUDA

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit b7bfe82807)

* [BugFix][JIT] Allocate a dynamic-shape output that precedes its sizing input (#3207)

* [BugFix][JIT] Allocate a dynamic-shape output that precedes its sizing input

A @tilelang.jit kernel whose output carries a symbolic dimension and appears
before the input supplying that dimension compiled fine but crashed at call
time with "IndexError: list index out of range" in the host wrapper.

Two things had to line up. _process_dynamic_symbolic walked the parameters in
signature order and recorded each symbolic dimension against the first
parameter mentioning it, outputs included, so for main(B(N,), A(N,)) with
out_idx=[0] the owner of N was B, the output being allocated. And the caller
assembled its tensor list in a single pass, so an output at parameter 0 would
have read tensor_list[1] before it was filled even with the owner pointing at
an input.

Visit inputs first when building the symbolic map, and place every input
tensor before allocating the outputs, in both the tvm_ffi and cython adapters.

Fixes #3017

* fix(jit/tvm_ffi): resolve scalar dimensions against the parameter-aligned list

`_process_dynamic_symbolic` records parameter indices taken from the PrimFunc
signature: a scalar variable parameter is stored as `(2, param_index, -1, 1)`.
The output-allocation loop resolved that branch as `inputs[ref_tensor_idx]`, but
`inputs` holds only the non-output parameters, so for a signature whose output
precedes its inputs the index points at the wrong slot and for
`[B_output, A_input, N_scalar]` it runs off the end of the list.

Resolve the scalar branch against `tensor_list` as well, which is sized and
indexed by parameter position, and report a clear error if the referenced scalar
slot is still unset instead of raising IndexError.

The non-scalar branch already gained the same treatment in the parent commit;
this makes the two branches consistent.

---------

Co-authored-by: xy200303 <xy200303@users.noreply.github.com>
(cherry picked from commit 6679f91387)

* [CUDA] Fuse exact FP4 to FP8 conversion through FP32 (#3204)

* [CUDA] Fuse exact FP4 to FP8 conversion through FP32

* [CUDA] Use PascalCase for the exact FP4 conversion helper

* [CUDA] Route the exact FP4 to FP8 transcode through __tl_cvt intrinsics

Separate the algebraic fold from the conversion intrinsic. The codegen now
rewrites Cast(e4m3, Cast(f32, e2m1)) with default rounding into a plain
Cast(e4m3, e2m1), and the direct pair is handled like every other cast in
VisitExpr_(CastNode): a scalar ::bitcast helper and PrintVectorizedCast with a
new chunk_lanes parameter so four lanes go through one __byte_perm pair.

The byte-level plumbing moves into cuda_fp4.h as __tl_cvt_e2m1x4_to_e4m3x4
with x2 and scalar wrappers, matching the header's existing helpers. A direct
T.Cast("float8_e4m3fn", fp4) now takes the same fast path instead of the
elementwise float fallback, and the pair is registered in
IsCudaVectorizableCast.

Tests cover both the folded chain and the direct cast. SASS for sm_90a is
unchanged in size with no stack frame.

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit 8cf0e7eff4)

* [Fix]Clamp T.copy/T.async_copy coalesced_width to achievable vector size instead of LOG(FATAL) (#3246)

* [Fix]Clamp T.copy/T.async_copy coalesced_width to achievable vector size instead of LOG(FATAL)

* add type cast

* format code and modify warning message

* [Fix] Clamp coalesced_width to the widest achievable width instead of gcd

Treat the hint as an upper bound on the per-thread vector width: take the
widest width the geometry supports without exceeding the request, stepping
down to a divisor of the geometry-derived width so alignment and contiguity
stay proven. gcd collapsed requests such as 257 or 3 to a scalar copy even
though float4 was achievable.

Pin the oversized test to the clamped width and cover the realistic
coalesced_width=8 case alongside 257.

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit 227515d27b)

* [BugFix] Bind loop targets independently of mutable scalar variables (#3232)

* [BugFix] Bind loop targets independently of mutable scalar variables

* [BugFix] Only store into a Ref/alloc_var whose binding region is still open

The store-vs-fresh decision in Builder.bind read the Python locals without
checking whether the TIR region that bound the name had already closed, so a
write to an expired alloc_var was silently emitted into the hoisted buffer
while a read of the same name was rejected. Treat an expired binding as absent
(the rule rval already applies to reads) and share the check between the two.

The loop_target flag keeps its remaining role: a live alloc_var shadowed by a
for target still becomes a fresh induction binding. Also stop resolving the
rewriter's `_` temporary against locals, which made tuple loop targets alias
an alloc_var named `_`, and skip the "re-bound" shadow warning for loop
targets, which stepped loops triggered with wrong advice.

Tests cover the expired-binding path structurally and at runtime, the live
nested-region store, the `_` tuple target, and the absent warning.

* [Frontend] Shorten the store-target comment in Builder.bind

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit e4e14156d7)

* [BugFix][Language] Lower NaN-propagating clamp through device templates (#3205)

* [BugFix][Language] Make T.clamp propagate NaN

T.clamp composed min(max(x, lo), hi). On CUDA those lower to the fmaxf/fminf
family, which returns the non-NaN operand, so clamp(NaN, lo, hi) silently
returned lo instead of NaN, unlike torch.clamp and numpy.clip.

Re-inject the input NaN through an if_then_else. tir.isnan is implemented for
only a subset of the float dtypes, so the predicate is evaluated on an fp32
cast, which preserves NaN; dtypes that cannot hold a NaN skip the guard
entirely. T.max and T.min keep their existing fmaxf/fminf behaviour.

Fixes #3024

* [BugFix][Language] Lower clamp through device templates

* [Test][ROCm] Use mcpu for the HIP clamp codegen target

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit 73c8fa0579)

* [Fix] Reject unsupported reduction NaN propagation dtypes (#3273)

* [Fix] Reject unsupported reduction NaN propagation dtypes

* [Refactor][Reduce] Name the min/max NaN guard condition

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit ea42ce56f4)

* [BugFix] Fix parallel loop lowering with let inlining disabled (#3269)

(cherry picked from commit 7763c88000)

* [Transform][CUDA] Fix async copy lowering with partitioned layouts (#3278)

[Transform] Refresh let bindings during loop partitioning

(cherry picked from commit 134ade7d81)

* [CUDA] Keep FP8 vector copies packed (#3276)

* [CUDA][Codegen] Preserve packed FP8 vector copies

* [CUDA] Keep packed FP8 copy change focused on codegen

* [CUDA] Define packed copy operations on FP8 vector types

(cherry picked from commit 356f309621)

* [CUDA] Support SM120 block-scaled GEMM fragments and odd warp atom grids (#3257)

* [CUDA] Support SM120 block-scaled GEMM fragments and odd warp atom grids

* [Test] Consolidate SM120 block-scaled GEMM regression coverage

* [CUDA] Consolidate SM120 block-scaled MMA selection conditions

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit 7a5f446fde)

* [CUDA] Reject conflicting SM120 scale fragment layouts (#3284)

(cherry picked from commit 74369d9e34)

* [JIT] Add missing uint64 argument type mappings (#3229)

(cherry picked from commit e40e09e3e6)

* [BugFix] Restore symbolic loop-layout injectivity proof; reject layouts on symbolic shared tiles (#3233)

* [BugFix] Support symbolic shared layouts with mixed fixed and dynamic extents

Related to tile-ai/tilelang#2906 and tile-ai/tilelang#2909. Credit Soham Panda for the original diagnosis and proposed defensive fix.

Co-authored-by: Soham Panda <sohampanda1@gmail.com>

* [Test] Minimize symbolic shared-layout regression coverage

* [Layout] Document why both symbolic injectivity proofs are kept

CanProveInjective and CanProveLeftInverse have disjoint blind spots:
equality propagation proves (i, i + j) but cannot invert floordiv/floormod
with a symbolic divisor, while the left-inverse decoder handles the padded
partition ((i*n+j)%128, (i*n+j)//128) but has no candidate for (i, i + j).
Record the counterexamples at the call site.

* [BugFix] Reject layouts on symbolic shared tiles instead of remapping them

makeBufferWithLayout dereferenced a null IntImm for a shared buffer with a
symbolic extent (#2906). Rather than deriving a symbolic replication factor,
keep the constant arithmetic and raise a ValueError naming the buffer and the
offending extent. The check sits at the remap site so that layouts inferred
by ops (scan) are covered as well as T.annotate_layout.

The left-inverse injectivity proof stays: it fixes T.Parallel loops over a
mixed static/symbolic space, which #2719 broke independently of shared
layouts. Its regression test is now a pure global-to-global loop, and the
#2906 test asserts the error for both the annotated and the inferred path.

---------

Co-authored-by: Soham Panda <sohampanda1@gmail.com>
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit efe9654bf4)

* [ROCm] Run portable example validation in CI (#3165)

* [ROCm] Run portable example validation in CI

Signed-off-by: andyluo7 <andy.luo@amd.com>

* [ROCm] Make portable example checks fail reliably

Signed-off-by: andyluo7 <andy.luo@amd.com>

---------

Signed-off-by: andyluo7 <andy.luo@amd.com>
(cherry picked from commit 48cd23e181)

* [BugFix][CUDA] Only emit 256-bit global load/store on SM100+ targets (#3248)

(cherry picked from commit 4bc6c32a63)

* [Metal] Support 32-bit integer atomic add (#3211)

(cherry picked from commit 901f941c0d)

* [Testing] Drop duplicated codegen-smoke tests; fix fastmath self-comparison assertions (#3252)

(cherry picked from commit 054bcf1de3)

* [ROCm] Add GLM-5.3 k-pool Top-K transform (#3254)

* [ROCm] Add GLM-5.3 k-pool cache writer

Signed-off-by: andyluo7 <andy.luo@amd.com>

* [ROCm] Match GLM-5.3 k-pool geometry

Signed-off-by: andyluo7 <andy.luo@amd.com>

* [ROCm] Add GLM-5.3 k-pool decode tail

Signed-off-by: andyluo7 <andy.luo@amd.com>

* [ROCm] Bind GLM-5.3 tail boundary count

Signed-off-by: andyluo7 <andy.luo@amd.com>

* [ROCm] Add GLM-5.3 paged k-pool logits

Signed-off-by: andyluo7 <andy.luo@amd.com>

* [ROCm] Guard GLM-5.3 k-pool page bounds

Signed-off-by: andyluo7 <andy.luo@amd.com>

* [ROCm] Add GLM-5.3 k-pool Top-K transform

Signed-off-by: andyluo7 <andy.luo@amd.com>

---------

Signed-off-by: andyluo7 <andy.luo@amd.com>
(cherry picked from commit eab74a4ae5)

* [Frontend] Support Python iterables and comprehensions (#3230)

* [Frontend] Support Python iterables and comprehensions

* [Frontend] Avoid excessive nesting in generated loop code

* [Frontend] Reject TIR values and non-iterables in compile-time for loops

The Python-iterable path in ctx_for consumed anything that was not a loop
frame. Var carries an __iter__ shim for single-binding unpacking, so
`for i in n` with a symbolic n expanded once with i aliased to n and renamed
the shape variable; Buffer.__getitem__ never raises IndexError, so
`for x in A` looped forever. Reject PrimExpr and Buffer up front and keep
the original diagnostic for objects that are not iterable at all.

* [Frontend] Validate Python iterable sources and comprehensions

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit 6ba187e20d)

* [Frontend] Remove warnings for immutable variable rebinding (#3262)

Remove warnings for immutable variable rebinding

(cherry picked from commit 195b6ebbf3)

* [BugFix] Reject non-power-of-two AllReduce thread strides (#3266)

The XOR butterfly in tl::AllReduce pairs thread t with t ^ (scale * k), which
only moves the reduce coordinate when `scale` (the thread stride between
consecutive reduce participants) is a power of two. #2611 added the check for
the logical width (threads / scale) but not for the stride, so reducing a
(32, 3) fragment along dim 0 with its default inferred layout lowered to
AllReduce<MaxOp, 96, 3> and silently mixed columns (and replicas); with exactly
96 or 192 threads the shared-memory exchange also read past the workspace.

- CheckAllReduceWidth, shared by tl.reduce and tl.finalize_reducer, now
  requires the stride to be a power of two and points the user at padding the
  non-reduced fragment extent.
- The CUDA AllReduce template carries the matching static_assert.
- The reducer v2 narrow-plan gate rejects such strides alongside the existing
  width check, so the planner falls back to the wide plan instead of emitting
  the broken collective.

(cherry picked from commit 1b908fc0eb)

* [NVRTC] Fix warp reduction compilation (#3260)

Fixes https://github.com/tile-ai/tilelang/issues/3259

(cherry picked from commit 261a9e4bb1)

* [BugFix][CUDA] Lower a kernel-body assert to a device-legal check (#3206)

* [BugFix][CUDA] Lower a kernel-body assert to a device-legal check

A plain assert in a kernel body is parsed into a tirx.AssertStmt. The CUDA
codegen did not override that visitor, so it inherited CodeGenC's, which
streams a host-only TVMFFIErrorSetRaisedFromCStrParts call plus `return -1`
into the emitted __global__ kernel. nvcc then rejects the file with
"identifier TVMFFIErrorSetRaisedFromCStrParts is undefined", so a valid
kernel refused to build.

Override VisitStmt_(AssertStmtNode) in CodeGenTileLangCUDA and lower to the
device helpers the T.device_assert intrinsic path already uses
(device_assert / device_assert_with_msg). Every function this codegen emits
is a __global__ kernel, so the override cannot misroute a host-side assert;
host and CPU code have their own codegens with their own visitors.

Fixes #3019

* [Test][CUDA] Gate kernel-body assert regressions on CUDA

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit b7bfe82807)

* [BugFix][JIT] Allocate a dynamic-shape output that precedes its sizing input (#3207)

* [BugFix][JIT] Allocate a dynamic-shape output that precedes its sizing input

A @tilelang.jit kernel whose output carries a symbolic dimension and appears
before the input supplying that dimension compiled fine but crashed at call
time with "IndexError: list index out of range" in the host wrapper.

Two things had to line up. _process_dynamic_symbolic walked the parameters in
signature order and recorded each symbolic dimension against the first
parameter mentioning it, outputs included, so for main(B(N,), A(N,)) with
out_idx=[0] the owner of N was B, the output being allocated. And the caller
assembled its tensor list in a single pass, so an output at parameter 0 would
have read tensor_list[1] before it was filled even with the owner pointing at
an input.

Visit inputs first when building the symbolic map, and place every input
tensor before allocating the outputs, in both the tvm_ffi and cython adapters.

Fixes #3017

* fix(jit/tvm_ffi): resolve scalar dimensions against the parameter-aligned list

`_process_dynamic_symbolic` records parameter indices taken from the PrimFunc
signature: a scalar variable parameter is stored as `(2, param_index, -1, 1)`.
The output-allocation loop resolved that branch as `inputs[ref_tensor_idx]`, but
`inputs` holds only the non-output parameters, so for a signature whose output
precedes its inputs the index points at the wrong slot and for
`[B_output, A_input, N_scalar]` it runs off the end of the list.

Resolve the scalar branch against `tensor_list` as well, which is sized and
indexed by parameter position, and report a clear error if the referenced scalar
slot is still unset instead of raising IndexError.

The non-scalar branch already gained the same treatment in the parent commit;
this makes the two branches consistent.

---------

Co-authored-by: xy200303 <xy200303@users.noreply.github.com>
(cherry picked from commit 6679f91387)

* [CUDA] Fuse exact FP4 to FP8 conversion through FP32 (#3204)

* [CUDA] Fuse exact FP4 to FP8 conversion through FP32

* [CUDA] Use PascalCase for the exact FP4 conversion helper

* [CUDA] Route the exact FP4 to FP8 transcode through __tl_cvt intrinsics

Separate the algebraic fold from the conversion intrinsic. The codegen now
rewrites Cast(e4m3, Cast(f32, e2m1)) with default rounding into a plain
Cast(e4m3, e2m1), and the direct pair is handled like every other cast in
VisitExpr_(CastNode): a scalar ::bitcast helper and PrintVectorizedCast with a
new chunk_lanes parameter so four lanes go through one __byte_perm pair.

The byte-level plumbing moves into cuda_fp4.h as __tl_cvt_e2m1x4_to_e4m3x4
with x2 and scalar wrappers, matching the header's existing helpers. A direct
T.Cast("float8_e4m3fn", fp4) now takes the same fast path instead of the
elementwise float fallback, and the pair is registered in
IsCudaVectorizableCast.

Tests cover both the folded chain and the direct cast. SASS for sm_90a is
unchanged in size with no stack frame.

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit 8cf0e7eff4)

* [Fix]Clamp T.copy/T.async_copy coalesced_width to achievable vector size instead of LOG(FATAL) (#3246)

* [Fix]Clamp T.copy/T.async_copy coalesced_width to achievable vector size instead of LOG(FATAL)

* add type cast

* format code and modify warning message

* [Fix] Clamp coalesced_width to the widest achievable width instead of gcd

Treat the hint as an upper bound on the per-thread vector width: take the
widest width the geometry supports without exceeding the request, stepping
down to a divisor of the geometry-derived width so alignment and contiguity
stay proven. gcd collapsed requests such as 257 or 3 to a scalar copy even
though float4 was achievable.

Pin the oversized test to the clamped width and cover the realistic
coalesced_width=8 case alongside 257.

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit 227515d27b)

* [BugFix] Bind loop targets independently of mutable scalar variables (#3232)

* [BugFix] Bind loop targets independently of mutable scalar variables

* [BugFix] Only store into a Ref/alloc_var whose binding region is still open

The store-vs-fresh decision in Builder.bind read the Python locals without
checking whether the TIR region that bound the name had already closed, so a
write to an expired alloc_var was silently emitted into the hoisted buffer
while a read of the same name was rejected. Treat an expired binding as absent
(the rule rval already applies to reads) and share the check between the two.

The loop_target flag keeps its remaining role: a live alloc_var shadowed by a
for target still becomes a fresh induction binding. Also stop resolving the
rewriter's `_` temporary against locals, which made tuple loop targets alias
an alloc_var named `_`, and skip the "re-bound" shadow warning for loop
targets, which stepped loops triggered with wrong advice.

Tests cover the expired-binding path structurally and at runtime, the live
nested-region store, the `_` tuple target, and the absent warning.

* [Frontend] Shorten the store-target comment in Builder.bind

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit e4e14156d7)

* [BugFix][Language] Lower NaN-propagating clamp through device templates (#3205)

* [BugFix][Language] Make T.clamp propagate NaN

T.clamp composed min(max(x, lo), hi). On CUDA those lower to the fmaxf/fminf
family, which returns the non-NaN operand, so clamp(NaN, lo, hi) silently
returned lo instead of NaN, unlike torch.clamp and numpy.clip.

Re-inject the input NaN through an if_then_else. tir.isnan is implemented for
only a subset of the float dtypes, so the predicate is evaluated on an fp32
cast, which preserves NaN; dtypes that cannot hold a NaN skip the guard
entirely. T.max and T.min keep their existing fmaxf/fminf behaviour.

Fixes #3024

* [BugFix][Language] Lower clamp through device templates

* [Test][ROCm] Use mcpu for the HIP clamp codegen target

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit 73c8fa0579)

* [Fix] Reject unsupported reduction NaN propagation dtypes (#3273)

* [Fix] Reject unsupported reduction NaN propagation dtypes

* [Refactor][Reduce] Name the min/max NaN guard condition

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit ea42ce56f4)

* [BugFix] Fix parallel loop lowering with let inlining disabled (#3269)

(cherry picked from commit 7763c88000)

* [Transform][CUDA] Fix async copy lowering with partitioned layouts (#3278)

[Transform] Refresh let bindings during loop partitioning

(cherry picked from commit 134ade7d81)

* [CUDA] Keep FP8 vector copies packed (#3276)

* [CUDA][Codegen] Preserve packed FP8 vector copies

* [CUDA] Keep packed FP8 copy change focused on codegen

* [CUDA] Define packed copy operations on FP8 vector types

(cherry picked from commit 356f309621)

* [CUDA] Support SM120 block-scaled GEMM fragments and odd warp atom grids (#3257)

* [CUDA] Support SM120 block-scaled GEMM fragments and odd warp atom grids

* [Test] Consolidate SM120 block-scaled GEMM regression coverage

* [CUDA] Consolidate SM120 block-scaled MMA selection conditions

---------

Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
(cherry picked from commit 7a5f446fde)

* [CUDA] Reject conflicting SM120 scale fragment layouts (#3284)

(cherry picked from commit 74369d9e34)

* [Ascend] Adapt upstream clamp lowering and regression coverage

---------

Signed-off-by: andyluo7 <andy.luo@amd.com>
Co-authored-by: Anders <anders@magnitude.dev>
Co-authored-by: Sepcnt <30561671+sepcnt@users.noreply.github.com>
Co-authored-by: Карим <kareemm@yandex-team.ru>
Co-authored-by: xuebozhang525-alt <xuebozhang525@gmail.com>
Co-authored-by: penguin_wwy <940375606@qq.com>
Co-authored-by: Alfred <166222074+Dino1844@users.noreply.github.com>
Co-authored-by: Soham Panda <sohampanda1@gmail.com>
Co-authored-by: andyluo7 <43718156+andyluo7@users.noreply.github.com>
Co-authored-by: Xiangwen Wang <77378439+LJC00118@users.noreply.github.com>
Co-authored-by: 小云 <130276203+xy200303@users.noreply.github.com>
Co-authored-by: xy200303 <xy200303@users.noreply.github.com>
Co-authored-by: Ziming Wang <125807850+ZenAlexa@users.noreply.github.com>
Co-authored-by: edragain <3425282590@qq.com>

* [Ascend][Docs] Extend npusim cycle measurements beyond VFs (#526)

* [Ascend][Docs] Extend npusim cycle measurements beyond VFs

* [Ascend][Docs] Document task cost overrides and SF copy cycles

* [Ascend][Docs] Keep cycle guidance focused on kernel authoring

* [Fix][TVM] Update TVM to preserve dynamic reduction tail guards (#3294) (#538)

[Fix][TVM] Update modular remainder bounds and preserve tail guards

* [Ascend] Support post-update SIMD loads (#537)

* [Ascend] Support post-update SIMD loads

* [Ascend] Infer SIMD pointer types and validate post-update loads

* [Ascend] Trim redundant post-update checks and tests

* [Ascend] Rename SIMD load increment to post_inc

* [Ascend] Calibrate blockscaled GEMM and scale-load costs (#528)

* [Ascend] Calibrate blockscaled GEMM latency and II

* [Ascend][Tests] Focus blockscaled GEMM cost coverage

* [Ascend] Keep latency parameter comments focused on semantics

* [Ascend] Model MX scale-factor copy latency and II

* [Ascend] Add DeepGEMM-compatible GEMM kernels (#291)

* [Ascend] Add DeepGEMM frontend kernels

Implement dense, batched, M-grouped, and K-grouped BF16/FP8 APIs with
MNK and KMN kernels, tail handling, and explicit L0 scale-factor buffers.

* [Ascend] Add DeepGEMM tests and benchmarks

Add focused numerical coverage and benchmark entry points, replace the
various-shapes example, and update the Ascend performance regression runner.

* [Ascend] Limit DeepGEMM launches to active output tiles

* [Ascend] Reuse MNK L0 operands across output tiles

Reuse only when the K iteration count equals the L0 stage count. Preserve
read-only K owners and the unrolled producer/consumer order.

* [Ascend] Refactor DeepGEMM APIs and compile heuristics with Cython

Move host configuration selection into a Cython extension built and installed
with TileLang, preserving the upstream cost model and candidate ordering.
Keep hashable configuration metadata in Python and reuse descriptor validation
across GEMM families. Bind layout-specific APIs through factories while
preserving public signatures and defaults.

Validate layouts before copying C and reject malformed scale recipes and
unsupported output descriptors explicitly.

Validation:
- 2002 heuristic comparisons against Python and upstream C++ with the existing
  TileLang launch cap normalized
- 108 API metadata, scale, and launch-argument differential cases
- Public signature, default, name, and serialization compatibility checks
- 4 existing DeepGEMM NPU tests
- Scoped repository format and staged whitespace checks

* [Ascend] Remove DeepGEMM heuristic cache and refresh comments

Call the Cython selector directly for each configuration request. Describe
current layout and launch contracts, use DeepGEMM-Ascend naming, and clarify
that configuration cache hints refer to hardware L2 behavior.

Validation: staged API import and uncached-call check, scoped repository
pre-commit checks, and staged diff checks.

* [Ascend] Integrate SIMD scale conversion into DeepGEMM

* [Ascend] Refactor DeepGEMM scale transform configuration

Select scale tiles, launch size, and load/store cache policies in a dedicated host config. Keep the group count dynamic and bound grouped metadata storage by tile geometry.

* [Ascend] Simplify DeepGEMM scale-transform scheduling

* [Build] Enable ROCm and Ascend by default on Linux (#544)

* [Ascend] Curate examples and consolidate feature tests (#545)

* [Ascend] Curate examples and consolidate feature tests

* [Ascend] Report example throughput and remove Top-K gate

Enable example performance reporting by default, with TFLOPS for matrix
kernels and logical I/O bandwidth for vector kernels. Add measured FP8 VF
latencies and remove the Top-K gate example, test, and references.

* [Ascend] Batch high-level SIMT RMSNorm row tiles

Amortize VF launches and small RSTD transfers by processing adjacent rows
together. Keep the computation in Parallel loops, fragments, and reducers so
thread mapping and collective synchronization remain compiler-generated.

Use compact UB tiles with two input versions, reuse the input for output when
the UB budget requires it, and retain weights in registers across rows. Cache
RSTD tiles for a strided writeback, with a bounded-cache fallback for large
batches. Preserve the latest benchmark output and command-line behavior.

Validation on Ascend950 with the existing local TileLang runtime:
- The existing RMSNorm tests and additional shape, zero-row, tail, cache-limit,
  and repeated-launch checks passed: 18 cases, with unchanged tolerances.
- Scoped pre-commit checks and git diff --check passed.
- FP32 batch=4096 measured 50.5/61.0/85.1 us for d=4096/5120/7168 with msprof,
  30 warmups and 50 measured iterations per case (one full profiling round).

Only the RMSNorm kernel example is changed; tests and benchmark artifacts are
not included in this commit.

* [Ascend] Optimize SIMT per-token FP8 quantization

Process adjacent 128-element groups in contiguous DMA tiles and write scales
through UB instead of scattered scalar stores. Select 32, 64, or 128 groups
per tile while keeping up to 64 cores busy, and explicitly double-buffer the
input, quantized output, and scale buffers.

Use a float2 fragment layout for contiguous UB accesses, explicit abs plus
max reduction, and the saturating FP8 cast without a redundant clamp. Enable
TileLang fast math and remove the VF latency measured for the old layout.
The SIMD implementation and example harness remain unchanged.

Validation: all four existing FP8 example tests and changed-file formatting
hooks pass on Ascend 950 with the existing local TileLang build. Three
interleaved msprof rounds on 8192x8192 FP32 input show median kernel latency
falling from 308.32 us to 111.05 us (2.78x), with TileLang fast math enabled
for both the original implementation and the optimized kernel.

* [Ascend] Improve high-level SIMT RMSNorm layout

Use a unit coalescing width for the column-parallel input, reduction,
weight, and output loops. Keep the reducer and thread mapping expressed
through high-level TileLang operations.

Copy the row reciprocal standard deviations from the result fragment to
UB at the end of the VF. This lets code generation select the writing
thread and removes redundant stores and a synchronization point.

Validation: 18 RMSNorm correctness cases passed, including narrow and odd
dimensions, the bounded RSTD cache fallback, and repeated launches.
Scoped pre-commit checks passed using the existing local TileLang build.

On Ascend950DT_9581, three-round msprof medians for batch 4096 and widths
4096/5120/7168 were 45.67/53.60/76.08 us, versus 48.75/58.11/82.88 us for
the fastest explicit thread-level implementation in the same run. This
reduces latency by 6.3%/7.8%/8.2%; batch 128 remains within 0.10 us of it.

* [Ascend] Translate FlashAttention README to English

* Update README.md

---------

Co-authored-by: Denver Jin <denverjin@outlook.com>
Co-authored-by: silentCoder-dev <silentcoder@foxmail.com>

* fix(ascend): correct swapped ROUND_F/ROUND_C rounding descriptions (#547)

The ``round`` parameter documentation of ``T.simd.vcvt`` described
``"ROUND_F"`` as rounding toward +inf (ceiling) and ``"ROUND_C"`` as
rounding toward -inf (floor). The two descriptions were swapped.

CANN's vconv wrappers pair ``ROUND_F`` with ``CastFloor`` and ``ROUND_C``
with ``CastCeil``, i.e. ROUND_F rounds toward -inf (floor) and ROUND_C
rounds toward +inf (ceiling). Fix the docstring to match the hardware
modes.

Docstring-only change; no functional impact.

* [CI] Restore upstream multi-backend checks

Restore the upstream lint, CUDA, ROCm, Metal and CuTeDSL jobs in place of
the fork-specific Ascend workflow. Defer dedicated Ascend CI integration.

Exclude examples/ascend from the CUDA and CuTeDSL example sweeps so those
jobs do not collect tests that require an NPU.

* [CI] Add Ascend tests on the self-hosted runner

* [CI] Restore the upstream distribution workflow

* [CI] Restore the upstream PR reminder workflow

* [Fix] Keep relaxed copy shape checks in the Ascend dialect

* [Fix] Preserve shape and stride source kinds in Cython arguments

* [Maint] Restore upstream tooling and defer the Ascend performance bot

* [Fix] Avoid narrowing conversion in vector size constraint

* [Docs][Ascend] Update backend documentation and benchmarks

* [Docs][Ascend] Refine backend acknowledgements

* [Fix] Restore broadcast vectorization and scalar reductions

---------

Signed-off-by: andyluo7 <andy.luo@amd.com>
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
Co-authored-by: matt-tian <81284512+matt-tian@users.noreply.github.com>
Co-authored-by: xuruifan <xuruifan@deepseek.com>
Co-authored-by: silentCoder-dev <silentcoder@foxmail.com>
Co-authored-by: Denver Jin <denverjin@deepseek.com>
Co-authored-by: Yongqi Zhuo <Yongqi-Zhuo@users.noreply.github.com>
Co-authored-by: Meng Yuhao <42602106+AutumnKite@users.noreply.github.com>
Co-authored-by: Denver Jin <45557821+Denverjin@users.noreply.github.com>
Co-authored-by: mengyuhao <mengyuhao@deepseek.com>
Co-authored-by: Lei Wang <34334180+LeiWang1999@users.noreply.github.com>
Co-authored-by: Denver Jin <denverjin@outlook.com>
Co-authored-by: cj <erhsh_165@126.com>
Co-authored-by: caojian5 <caojian5@huawei.com>
Co-authored-by: Anders <anders@magnitude.dev>
Co-authored-by: Sepcnt <30561671+sepcnt@users.noreply.github.com>
Co-authored-by: Карим <kareemm@yandex-team.ru>
Co-authored-by: xuebozhang525-alt <xuebozhang525@gmail.com>
Co-authored-by: penguin_wwy <940375606@qq.com>
Co-authored-by: Alfred <166222074+Dino1844@users.noreply.github.com>
Co-authored-by: Soham Panda <sohampanda1@gmail.com>
Co-authored-by: andyluo7 <43718156+andyluo7@users.noreply.github.com>
Co-authored-by: Xiangwen Wang <77378439+LJC00118@users.noreply.github.com>
Co-authored-by: 小云 <130276203+xy200303@users.noreply.github.com>
Co-authored-by: xy200303 <xy200303@users.noreply.github.com>
Co-authored-by: Ziming Wang <125807850+ZenAlexa@users.noreply.github.com>
Co-authored-by: edragain <3425282590@qq.com>
2026-09-30 07:31:01 +08:00

39 lines
1.7 KiB
CMake
Raw Permalink Blame History

# Resolve each backend independently. Linux builds include ROCm and Ascend by
# default: vendored/local headers and runtime stubs avoid a build-time dependency
# on either SDK. CUDA still needs a toolkit, while macOS defaults to Metal.
set(TILELANG_BACKENDS CUDA ROCM METAL LLVM ASCEND)
set(TILELANG_BACKEND_DOC_CUDA "Enable CUDA backend (ON/OFF/or CUDA SDK path)")
set(TILELANG_BACKEND_DOC_ROCM "Enable ROCm backend (ON/OFF/or ROCm SDK path)")
set(TILELANG_BACKEND_DOC_METAL "Enable Metal backend")
set(TILELANG_BACKEND_DOC_LLVM "Enable LLVM backend")
set(TILELANG_BACKEND_DOC_ASCEND "Enable Ascend backend")
foreach(BACKEND IN LISTS TILELANG_BACKENDS)
set(_backend_var "USE_${BACKEND}")
set(_doc "${TILELANG_BACKEND_DOC_${BACKEND}}")
set(_default OFF)
if(BACKEND STREQUAL "CUDA" AND NOT APPLE AND TILELANG_CUDA_TOOLKIT_AVAILABLE)
set(_default ON)
elseif(BACKEND STREQUAL "METAL" AND APPLE)
set(_default ON)
elseif(CMAKE_SYSTEM_NAME STREQUAL "Linux" AND
(BACKEND STREQUAL "ROCM" OR BACKEND STREQUAL "ASCEND"))
set(_default ON)
endif()
# CMake variables (including cached values and -D arguments) take precedence
# over environment variables. Environment variables initialize a fresh cache.
if(DEFINED ${_backend_var})
set(_default "${${_backend_var}}")
elseif(DEFINED ENV{${_backend_var}})
set(_default "$ENV{${_backend_var}}")
endif()
# STRING preserves SDK paths as well as ON/OFF values.
set(${_backend_var} "${_default}" CACHE STRING "${_doc}")
# TVM's config.cmake redefines USE_* options later. Save the resolved value
# so the caller can restore it after including that configuration.
set(TILELANG_OPTION_${_backend_var} "${${_backend_var}}")
endforeach()
Reference in New Issue View Git Blame Copy Permalink