Programming models insights

Survey of existing tile-based and kernel DSL programming models relevant to Ascend NPU compiler design. Companion document to design-overview.md.

1. AscendC 

AscendC is the official C++ kernel language for Ascend NPU and the compilation target for PyAsc2. Understanding how it handles the three key challenges defines the baseline that PyAsc2 must improve upon.

AscendC is not a single programming model — it has evolved with the hardware. On A2/A3 (910B/C) the only level is basic_api (TPipe/TQue, memory-centric). On A5 (950) two additional levels appeared — MicroAPI (register-tensor SIMD with predication) and SIMT-API (CUDA-like per-thread scalar) — neither of which compiles for A2. We therefore split this section along the A2/A3 vs A5 axis: §1.1 covers basic_api behavior that applies to both targets (content written against A2’s c220, unchanged on c310); §1.2 covers what is new and A5-exclusive.

1.1 AscendC on A2/A3 

On A2/A3 there is one level: basic_api via build-mode c220. The three key challenges are analyzed below in this level.

Sync Insertion 

AscendC exposes synchronization directly to the user via TPipe and TQue. Every data transfer requires explicit EnQue/DeQue calls; every pipeline stage boundary requires explicit SetEventId/WaitEventId:

TPipe pipe;
TQue<QuePosition::VECIN, 2> inQueue;
TQue<QuePosition::VECOUT, 1> outQueue;
pipe.InitBuffer(inQueue, 2, TILE_SIZE);    // TPipe allocates UB memory for TQue buffers
pipe.InitBuffer(outQueue, 1, TILE_SIZE);

// MTE2 stage: load
LocalTensor<half> tile = inQueue.AllocTensor<half>();
DataCopy(tile, gm_src[offset], TILE_SIZE);
inQueue.EnQue(tile);              // signal: load done

// Vector stage: compute
LocalTensor<half> tile = inQueue.DeQue<half>();   // wait: load done
Add(out, tile, tile2, TILE_SIZE);
outQueue.EnQue(out);

// MTE3 stage: store
LocalTensor<half> out = outQueue.DeQue<half>();
DataCopy(gm_dst[offset], out, TILE_SIZE);
outQueue.FreeTensor(out);

// Cross-unit sync (e.g. between independent pipelines)
pipe.SetEventId(EVENT_ID0);       // producer signals
pipe.WaitEventId(EVENT_ID0);      // consumer waits

Source: Ascend C Operator Development Guide, CANN 8.0 https://www.hiascend.com/document/detail/en/canncommercial/800/opdevg/Ascendcopdevg/atlas_ascendc_10_0001.html

Missing or misplaced EnQue/DeQue / SetEventId/WaitEventId causes silent data hazards. AscendC does not insert any barriers automatically.

Conclusion: AscendC doesn’t solve the challenge — it exposes it. The user is responsible for every barrier manually. This is a source of bugs.

Ping-pong 

// Queue depth=2: two UB slots (ping + pong)
// Depth is set in TWO places — they must match:
constexpr int QUEUE_DEPTH = 2;

TQue<QuePosition::VECIN, QUEUE_DEPTH> inQueue;   // (1) template param
pipe.InitBuffer(inQueue, QUEUE_DEPTH, TILE_SIZE); // (2) runtime init

for (int i = 0; i < num_tiles; i++) {
    // MTE2: load tile i+1 while Vector computes tile i
    LocalTensor<half> tile = inQueue.AllocTensor<half>();
    DataCopy(tile, gm_src[(i+1) * TILE_SIZE], TILE_SIZE);
    inQueue.EnQue(tile);

    // Vector: compute tile i
    LocalTensor<half> cur = inQueue.DeQue<half>();
    Add(out, cur, cur2, TILE_SIZE);
    inQueue.FreeTensor(cur);
}

Depth 2 = two slots in UB. While Vector works on ping, MTE2 loads into pong. But: User chooses queue depth manually; User structures the loop manually to achieve overlap; Compiler doesn’t help

Conclusion: AscendC supports ping-pong, but requires manual orchestration. No automatic loop body partitioning.

UB Memory Allocation and Reuse 

TPipe pipe;
TQue<QuePosition::VECIN,  2> inQueueA;   // reserves 2 × TILE_SIZE bytes in UB
TQue<QuePosition::VECIN,  2> inQueueB;   // reserves 2 × TILE_SIZE bytes in UB
TQue<QuePosition::VECOUT, 1> outQueue;   // reserves 1 × TILE_SIZE bytes in UB
// Total: 5 × TILE_SIZE — must fit in 256 KB; no compiler check

pipe.InitBuffer(inQueueA,  2, TILE_SIZE);
pipe.InitBuffer(inQueueB,  2, TILE_SIZE);
pipe.InitBuffer(outQueue,  1, TILE_SIZE);

for (int i = 0; i < num_tiles; i++) {
    LocalTensor<half> a = inQueueA.AllocTensor<half>();  // acquire free UB slot
    // ... use a ...
    inQueueA.FreeTensor(a);  // return slot — reused next iteration
}

TPipe partitions UB at kernel initialization via InitBuffer — one call per queue, each reserving depth × tile_size bytes statically. Within the loop, AllocTensor acquires a free slot from the queue’s pool and FreeTensor returns it for reuse on the next iteration. Two queues never share a UB region, even if they are never live simultaneously — there is no liveness analysis.

Conclusion: AscendC requires the user to manually plan UB layout — sizing each buffer, choosing queue depths, and validating the total fit within 256 KB. Exceeding the limit causes silent memory corruption at runtime. The compiler provides no assistance.

1.2 AscendC on A5 

On A5, basic_api still exists — the same TPipe / TQue / LocalTensor programming style analyzed in §1.1 is available via build-mode c310, and the three-challenge analysis above carries over unchanged. What is new on A5 are (a) c310 deltas inside basic_api, and (b) two brand-new API levels (MicroAPI, SIMT-API) that are unavailable on A2. ^[58]

The three levels stack as follows:

Level	A2 (c220)	A5 (c310)	Style
basic_api	✓ 35 impl files	✓ 39 impl files	Hardware intrinsics, memory-centric (`__ubuf__` ptrs + explicit `repeatTime` / `BlkStride` / `RepStride`)
MicroAPI	✗ absent	✓ 20 interface files + `dav_c310/` backend	Register-tensor SIMD functional. `RegTensor<T>` values, `MaskReg` predication, `LoadAlign`/`StoreAlign`, arch-dispatched via `__NPU_ARCH__`
SIMT-API	✗ absent	✓ 21 files + `dav_c310/` backend	CUDA-like per-thread scalar. Per-thread values, atomics on `__gm__`/`__ubuf__`, warp-level primitives

Canonical example — same op (Relu) in each style:

// A2 basic_api (dav_c220)
template <typename T>
void ReluIntrinsicsImpl(__ubuf__ T* dst, __ubuf__ T* src,
                        uint8_t repeatTime, const UnaryRepeatParams& p) {
    vrelu(dst, src, repeatTime,
          p.dstBlkStride, p.srcBlkStride, p.dstRepStride, p.srcRepStride);
}

// A5 MicroAPI (micro_api/dav_c310)
namespace MicroAPI {
template <typename T, MaskMergeMode mode, typename U>
void Relu(U& dstReg, U& srcReg, MaskReg& mask) {      // U = RegTensor<T>
    ReluImpl<T, mode, U>(dstReg, srcReg, mask);
}
}

// A5 SIMT-API (simt_api/dav_c310)
template <typename T>
T ReluImpl(T x) { return (x < 0) ? 0 : x; }           // scalar on thread

A kernel written against MicroAPI or SIMT-API will not compile for A2 — both levels are A5-exclusive and both require the c310 build-mode. Choosing a level is therefore a portability decision a DSL must make explicit.

basic_api on c310 — what changed 

The programming model and the three-challenge baseline (manual sync, manual ping-pong, manual UB layout) are unchanged at the basic_api level on A5 — the primitive set widened but the programming style did not.

MicroAPI — register-tensor SIMD with predication (A5-exclusive)

MicroAPI operates on register tensors, not memory. The primitives take RegTensor<T> values, a first-class MaskReg for predication, and LoadAlign / StoreAlign to move data between UB and registers. The MicroAPI interface dispatches to per-arch backends at compile time via __NPU_ARCH__:

// micro_api/kernel_micro_vec_unary_intf_impl.h
#if   __NPU_ARCH__ == 3003
#include "micro_api/dav_l300/kernel_micro_vec_unary_impl.h"
#elif __NPU_ARCH__ == 3113
#include "micro_api/dav_l311/kernel_micro_vec_unary_impl.h"
#elif __NPU_ARCH__ == 5102
#include "micro_api/dav_m510/kernel_micro_vec_unary_impl.h"
#else
#include "micro_api/dav_c310/kernel_micro_vec_unary_impl.h"   // A5
#endif

namespace MicroAPI {
template <typename T, MaskMergeMode mode, typename U>
__simd_callee__ inline void Relu(U& dstReg, U& srcReg, MaskReg& mask) {
    ReluImpl<T, mode, U>(dstReg, srcReg, mask);
}
}

The caller pattern shipped in c310’s basic_api/dav_c310/kernel_operator_vec_unary_impl.h shows how MicroAPI ops compose into a kernel: a loop iterates over register-tile chunks, each chunk loaded via LoadAlign, processed under a predication mask, and stored via StoreAlign.

for (uint16_t i = 0; i < repeatTime; ++i) {
    mask = MicroAPI::UpdateMask<T, RegType::trait>(sreg);
    MicroAPI::LoadAlign(srcReg, src + i * repeatStride);
    func(dstReg, srcReg, mask);                        // e.g. MicroAPI::Relu
    MicroAPI::StoreAlign(dst + i * repeatStride, dstReg, mask);
}

A more substantial example — row-wise softmax over a UB tile — shows the MicroAPI style more clearly. On A5, softmax was rewritten from memory-based to register-based, taking advantage of the exposed register file:

// Simple softmax with MicroAPI — 3 phases over register-tile chunks:
// (1) row max, (2) exp(x - max) and accumulate sum, (3) divide by sum.
MicroAPI::MaskReg preg = MicroAPI::CreateMask<uint32_t, MicroAPI::MaskPattern::ALL>();
MicroAPI::RegTensor<float> srcVreg, maxVreg, sumVreg, tmpVreg;

Duplicate(maxVreg, F32_NEG_INF);
for (uint16_t j = 0; j < repeatTimes; ++j) {
    MicroAPI::LoadAlign(srcVreg, srcUb + j * FLOAT_REPEAT_SIZE);
    MicroAPI::Max(maxVreg, maxVreg, srcVreg, preg);
}
MicroAPI::ReduceMax(maxVreg, maxVreg, preg);            // register-level reduction

Duplicate(sumVreg, 0.0f);
for (uint16_t j = 0; j < repeatTimes; ++j) {
    MicroAPI::LoadAlign(srcVreg, srcUb + j * FLOAT_REPEAT_SIZE);
    MicroAPI::FusedExpSub(tmpVreg, srcVreg, maxVreg, preg);   // fused exp(x - max)
    MicroAPI::StoreAlign(dstUb + j * FLOAT_REPEAT_SIZE, tmpVreg, preg);
    MicroAPI::Add(sumVreg, sumVreg, tmpVreg, preg);
}
MicroAPI::ReduceSum(sumVreg, sumVreg, preg);

for (uint16_t j = 0; j < repeatTimes; ++j) {
    MicroAPI::LoadAlign(tmpVreg, dstUb + j * FLOAT_REPEAT_SIZE);
    MicroAPI::Div(tmpVreg, tmpVreg, sumVreg, preg);
    MicroAPI::StoreAlign(dstUb + j * FLOAT_REPEAT_SIZE, tmpVreg, preg);
}

Three loops over repeatTimes register-tile chunks; inside each loop the values stay in vector registers (RegTensor<float>), with only LoadAlign / StoreAlign crossing the UB boundary. Two distinctive A5 primitives are visible: ReduceMax / ReduceSum perform reduction within a single register (cross-register reduction still needs the outer loop), and FusedExpSub collapses exp(x - max) into one op — the register file’s capacity and the fused ISA make these chains cheap enough to stay register-resident, which is why softmax migrated off the memory-based path on A5. ^[58]

Sync Insertion. MicroAPI does not introduce automatic barriers. The MaskReg predication controls which lanes of a register tile participate in an op — it is not a barrier. Cross-unit synchronization between MTE/Vector/Cube pipelines still relies on set_flag / wait_flag and TPipe events from basic_api. Conclusion: sync is still the user’s responsibility, and MicroAPI kernels typically still layer over basic_api for data movement.

Ping-pong. Pipelining shifts from UB-tile granularity to register-tile granularity. The caller loop above runs over repeatTime register chunks per UB tile; overlap between load (MTE2) and compute (Vector) is still achieved via ping-pong at the UB level (basic_api), but within a UB tile, MicroAPI adds a second axis of latency hiding via chunked register operations. Conclusion: MicroAPI does not replace UB-level ping-pong; it adds a finer pipelining axis below it, both still manual.

UB Memory Allocation. MicroAPI introduces the register file as a planned address space. RegTensor<T> values occupy vector registers; LoadAlign / StoreAlign are the explicit move operations between UB and registers. UB layout itself is still managed by the basic_api TPipe / TQue machinery. Conclusion: UB allocation remains a basic_api concern; MicroAPI adds a second-tier register-allocation problem that the programmer (or a DSL compiler) must solve — evidenced by softmax’s end-to-end rewrite from membase/ to regbase/ on A5.

Overall: MicroAPI is a new level below basic_api, not a replacement. It exposes the register file and predication; it does not automate sync or memory planning.

SIMT-API — brief note (not a pyasc2 target)

SIMT-API is a CUDA-like per-thread scalar level (AbsImpl, AtomicCasImpl, warp-level primitives, a CPU-debug shim). It is A5-exclusive and represents a different programming model entirely: instead of tile-level intrinsics it exposes per-thread operations, with the compiler and hardware responsible for SIMT-style lane grouping and latency hiding.

pyasc2 does not target SIMT-API. SIMT-style per-thread code on Ascend is not expected to reach the performance ceiling needed by a pyasc2 kernel; the DSL’s goal is ≥90% of peak hardware potential and SIMT lowers that ceiling by giving up tile-level orchestration. SIMT-API is noted here for completeness and because its atomics are mirrored into basic_api.

pyasc2 implication 

pyasc2 targets basic_api (baseline, portable across A2 and A5) and MicroAPI (A5-only, for register-file throughput and the new dtype matrix). SIMT-API is out of scope. A pyasc2 program targeting A5 must lower to a c310 build that mixes basic_api for data movement and TPipe orchestration with MicroAPI for compute inside tiles; targeting A2 lowers to basic_api only (c220). This level split is the core portability decision §4 must make explicit.

2. Triton 

Sync Insertion 

No sync primitives are exposed to the user. The TritonGPU IR lowering pass inserts async copy wait groups and shared memory barriers automatically, based on data flow analysis. ^[10]

@triton.jit
def add_kernel(X, Y, Z, N, BLOCK: tl.constexpr):
    offs = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK)
    x = tl.load(X + offs)   # no barrier needed — compiler inserts it
    y = tl.load(Y + offs)
    tl.store(Z + offs, x + y)

Conclusion: Sync is fully automated. The user cannot and does not insert barriers manually.

Ping-Pong 

The user writes a plain loop. The compiler’s software pipelining pass automatically reorders instructions across iteration boundaries — issuing async loads for iteration K+N-1 while computing iteration K. Pipeline depth is controlled by a single num_stages hint. ^[10]

@triton.jit
def matmul_kernel(A, B, C, M, N, K,
                  BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr,
                  BLOCK_K: tl.constexpr, num_stages: tl.constexpr):
    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
    for k in tl.range(0, K, BLOCK_K, num_stages=num_stages):  # compiler pipelines this loop
        a = tl.load(...)
        b = tl.load(...)
        acc += tl.dot(a, b)
    tl.store(...)

Conclusion: Loop restructuring and barrier insertion are fully automated. num_stages is a user-facing tuning parameter — in practice set via autotuning, not manually.

UB Memory Allocation and Reuse 

Shared memory (GPU equivalent of Ascend UB) is managed entirely by the compiler. No user-visible allocation calls exist. If the total shared memory footprint exceeds hardware limits, Triton raises a compile-time error — unlike AscendC which silently corrupts memory at runtime.

Triton handles shared memory allocation through two separate mechanisms:

1. Encoding assignment — decides what goes to shared memory:

Pure elementwise ops (add, mul, relu) — encoding passes through unchanged (SameOperandsAndResultEncoding trait); tensors stay in registers. No shared memory allocated.
Reductions (softmax, sum, max) — ReduceDataDuplication pass assigns SwizzledSharedEncodingAttr to reduction intermediates, inserting LocalAllocOp for shared memory staging. ^[12]
Matrix ops (tl.dot) — BlockedToMMA pass converts operand encodings from Blocked to MMA shared memory encoding, inserting layout conversion ops automatically. ^[13]

2. Liveness analysis — decides where in shared memory each buffer lives: AllocationAnalysis class in Allocation.cpp runs three phases: ^[11]

getValuesAndSizes() — collects all shared memory values and their sizes
resolveLiveness() — computes live ranges via MLIR standard liveness analysis
computeOffsets() — assigns offsets using interference graph + first-fit graph coloring, reusing freed regions for non-overlapping buffers

Conclusion: The user writes the same code regardless of whether shared memory is needed. Encoding assignment and liveness-based reuse are fully automated by the compiler. Unlike AscendC, fit within hardware limits is validated at compile time.

3. cuTile 

Sync Insertion 

No sync primitives are exposed to the user. The TileIR compiler automatically injects all necessary barriers during code generation. Explicit synchronization or communication within a block is not permitted by design — this is a fundamental constraint of the cuTile model. ^[14]

This includes TMA (Tensor Memory Accelerator) async operations: the convert-tileaa-to-tileas pass converts tileaa.tiled_load into async loads, and the convert-pipeline-to-nvvm pass generates the corresponding nvvm.mbarrier.* intrinsics for memory-vs-compute synchronization. ^{[15] [16]}

TODO: Exmplanation above is about MMA. Better here to insert GEMM example (not vecadd)

@ct.kernel
def vector_add(a, b, c, tile_size: ct.Constant[int]):
    pid = ct.bid(0)
    a_tile = ct.load(a, index=(pid,), shape=(tile_size,))
    b_tile = ct.load(b, index=(pid,), shape=(tile_size,))
    result = a_tile + b_tile
    ct.store(c, index=(pid,), tile=result)

An 86-line Python cuTile kernel expands to ~1,900 lines of PTX with 20 barrier objects — none written by the developer. ^[15]

TODO: if possible to get references - insert perf numbers B200 GEMM (matmul kernels) Conclusion: Sync is fully automated and not even expressible by the user. cuTile is the most restrictive model — no escape hatch for manual barriers.

Ping-Pong 

The user writes a simple loop. The TileIR compiler transforms it into a pipelined loop through three passes: ^{[15] [16]}

convert-tileaa-to-tileas — converts tiled loads into async loads with pipeline ops

TODO: clarify if this “tiled loads are about MMA or vector ops

tileas-materialize-async — creates the async pipeline structure with multi-buffering
convert-pipeline-to-nvvm — lowers to NVVM barrier intrinsics (nvvm.mbarrier.*)

The result is a three-phase loop — prologue (pre-load N iterations), steady-state (overlap load K+N with compute K), and epilogue (drain remaining computes).

Unlike Triton, there is no user-facing num_stages parameter — the compiler determines pipeline depth automatically.

Performance evidence: GEMM achieves ~90% of cuBLAS on Blackwell with zero user-written pipelining. ^[17] Attention kernels are still a work in progress — TiledAttention research implementation reaches 0.63x vs fused PyTorch SDPA. ^[18]

Conclusion: Ping-pong is fully automated. No loop restructuring, no buffer management, no pipeline depth hint from the user. Strong results for GEMM; attention still maturing.

UB Memory Allocation and Reuse 

SIMT kernels (shared memory): The compiler decides what goes to shared memory and allocates it automatically. The user has no shared memory API — no __shared__ declarations, no size hints. In practice, the compiler allocated 180 KB of shared memory for an 86-line MOE kernel without any user input. ^{[15] [19]}

Tensor core kernels on Blackwell (TMEM): SM100 introduces TMEM — 256 KB per SM, dedicated to tensor cores, separate from shared memory. Operand A lives in TMEM or SMEM, operand B in SMEM, accumulator in TMEM exclusively. Allocation is dynamic via tcgen05.alloc (32-column minimum, power-of-2 granularity). ^[20]

cuTile handles TMEM automatically including contention handling — retry with 100ns backoff when allocation fails. On Hopper, matrix operands competed for register file space; on Blackwell, TMEM decouples tensor cores from CUDA cores entirely. ^[19]

TODO: about not clear. what is 100us? what is matrix operands “competed”?

Conclusion: All on-chip memory allocation — shared memory for SIMT, TMEM for tensor cores — is fully compiler-managed. No user-visible allocation API exists.

4. TileLang-Ascend 

Sync Insertion 

TileLang-Ascend offers two modes — manual and automatic.

Manual: user writes explicit primitives — T.set_flag(), T.wait_flag(), T.set_cross_flag() / T.wait_cross_flag() (for Cube↔Vector sync), T.pipe_barrier(), T.barrier_all(). ^[21]

Automatic: enabled via pass config: ^[22]

pass_configs = {
    tilelang.PassConfigKey.TL_ASCEND_AUTO_SYNC: True,
    tilelang.PassConfigKey.TL_ASCEND_AUTO_CV_SYNC: True,
}

The AscendSyncInsert pass analyzes producer-consumer data dependencies, separates Cube/Vector code regions (CombineCV pass), and inserts fine-grained flag operations at synchronization points. ^[23]

Known limitation: with double buffering + auto sync, the compiler currently generates conservative PipeBarrier<PIPE_ALL> instead of fine-grained flags, which may hurt performance. ^[24]

Conclusion: Hybrid model — auto sync available but not yet mature. Manual sync remains the practical default for performance-critical kernels.

Ping-Pong 

T.Pipelined(range, num_stages=N) is the primary API. The user sets pipeline depth and marks which loop to pipeline — necessary in nested loops where different levels produce different overlap behavior. The CrossCorePipeline pass extends buffer dimensions and restructures the loop. ^[21]

for m in T.range(M // block_M):
    for k in T.Pipelined(K // block_K, num_stages=2):  # only this loop is pipelined
        T.copy(A_global, A_L1)
        T.gemm_v0(A_L1, B_L1, C_L0C)

Similar to Triton’s num_stages — user-facing tuning parameter, not fully automated. Combined with the auto sync limitation (Issue #110), pipelined kernels currently get conservative PipeBarrier<PIPE_ALL> instead of fine-grained flags. ^[24]

Conclusion: Loop pipelining is compiler-supported via T.Pipelined, but pipeline depth and loop selection are user-chosen. Barrier insertion is still conservative.

UB Memory Allocation and Reuse 

User explicitly allocates buffers at each memory hierarchy level: ^[21]

q_l1     = T.alloc_L1([block_M, dim], dtype)        # L1 buffer
acc_l0c  = T.alloc_L0C([block_M, block_N], dtype)   # L0C (register-like SRAM)
softmax  = T.alloc_ub([block_M // 2], dtype)         # UB (256 KB limit)

Default mode: no liveness analysis. User must manually assign UB offsets via T.annotate_address(buffer, offset) — similar to AscendC’s manual planning. No compile-time validation; exceeding 256 KB causes silent corruption.

Opt-in mode: TL_ASCEND_MEMORY_PLANNING=True enables the AscendMemoryPlanning pass — liveness analysis, interference graph, first-fit coloring for buffer reuse. Eliminates manual T.annotate_address(). Validates UB fit at compile time. ^[25]

L1 and L0A/B/C allocation has no compile-time limit checking in either mode.

Conclusion: Memory hierarchy placement is always user-chosen (unlike Triton). Buffer reuse and validation are available but opt-in — not yet the default.

TODO: Clarify why Buffer reuse and validation is not yet default option.

5. Triton-Ascend 

Triton-Ascend is Huawei’s fork of Triton for Ascend NPU. From the user’s perspective — standard Triton Python API. Internally, a completely different compilation path. ^[28]

Compilation pipeline — does not use TTGIR (GPU-specific), stays at TTIR: ^{[28] [29]}

TTIR → triton-to-linalg → Linalg IR → HFusion → HIVM
  → HIVMToStandard (HIVM ops → CCE device library calls, e.g. vadd_2d_f16)
  → LLVM IR → kernel.o (linked with CCE libs, loaded via ascendcl)

HIVM (Hybrid ISA Virtual Machine) is the core dialect of AscendNPU-IR — abstracts Ascend computation, data movement, and synchronization at tile level. Built on standard MLIR dialects: Linalg, MemRef, SCF, Bufferization, Tensor. ^[29]

Confirmed working on Ascend (with torch_npu): matmul (tl.dot), fused softmax, layer norm, fused attention (Flash Attention v2), vector add. ^[27]

Sync Insertion 

Automated — same as standard Triton, no explicit barriers from the user. ^[26]

Internally handled by HIVM passes (bishengir/lib/Dialect/HIVM/Transforms/): ^[30]

InjectSync/ — intra-core sync: inserts set_flag/wait_flag between MTE and compute pipelines based on data dependency analysis
GraphSyncSolver/ — graph-based solver for optimal barrier placement, minimizing redundant barriers while preserving correctness
InjectBlockSync.cpp — inter-block sync for cross-core data dependencies
SplitMixKernel.cpp — splits CV-fused kernels into separate AIC (cube) and AIV (vector) functions, inserting cross-core sync at data exchange points

Confirmed by lit tests: inject-sync.mlir, sync-solver.mlir, sync-solver-cross-core.mlir, inject-block-sync.mlir. ^[30]

Key Ascend-specific adaptation at user level: grid is fixed to the number of physical cores (not thousands of blocks like GPU). For large data, two-level tiling with BLOCK_SIZE_SUB — an inner loop to fit 192 KB UB: ^[26]

@triton.jit
def kernel(inp, out, N, BLOCK_SIZE: tl.constexpr, BLOCK_SIZE_SUB: tl.constexpr):
    pid = tl.program_id(0)
    NUM_CORE = tl.num_programs(0)
    num_blocks = tl.cdiv(N, BLOCK_SIZE)
    for block_idx in range(pid, num_blocks, NUM_CORE):  # round-robin across cores
        base = block_idx * BLOCK_SIZE
        for sub_idx in range(BLOCK_SIZE // BLOCK_SIZE_SUB):  # fit UB
            offs = base + sub_idx * BLOCK_SIZE_SUB + tl.arange(0, BLOCK_SIZE_SUB)
            x = tl.load(inp + offs, mask=offs < N)
            tl.store(out + offs, x, mask=offs < N)

What works as-is: small-scale kernels where one block fits in UB. What requires adaptation: large data (add BLOCK_SIZE_SUB loop), grid sizing (fix to physical core count), i32/i64 comparisons (cast to float32), tail axis alignment (32-byte for vector, 512-byte for CV ops). ^[26]

Conclusion: Sync is automated. Standard Triton kernels can run on Ascend for small-scale ops; large-scale kernels need two-level tiling and grid adaptation.

Ping-Pong 

multiBuffer is enabled by default — no user annotation required. ^[26]

Implemented by HIVM passes (bishengir/lib/Dialect/HIVM/Transforms/): ^[30]

EnableMultiBuffer.cpp — enables double buffering for eligible allocations
MarkMultiBuffer.cpp — marks which buffers should be multi-buffered based on access patterns and loop structure
CVPipelining.cpp — reorders cube and vector code to enable CV core pipeline parallelism (load on one core while other computes)
OptMemPlanForPipeline.cpp — adjusts memory plan to accommodate pipelined execution

The BLOCK_SIZE_SUB inner loop creates the pipelining opportunity: each sub-block iteration is a natural stage for overlap. After doublebuffer is enabled, UB capacity is halved (192 KB → 96 KB effective). ^[26]

Confirmed by lit tests: enable-multi-buffer.mlir, mark-multi-buffer.mlir, cv-pipelining.mlir. ^[30]

Conclusion: Ping-pong is fully automated via multiBuffer (on by default). No num_stages hint like Triton GPU, no T.Pipelined like TileLang-Ascend.

UB Memory Allocation and Reuse 

UB allocation is handled by the hivm-plan-memory pass. Two-phase approach: ^[30]

MemLivenessAnalysis — uses standard MLIR Liveness analysis, extended with Ascend-specific buffer aliasing, multi-buffer tracking, and gen/kill maps
MemPlan — assigns UB offsets based on liveness, reuses non-overlapping buffers, handles multi-buffer (ping-pong) reuse scenarios

Hardware limits hardcoded in PlanMemory.cpp: UB = 192 KB, L1 = 512 KB, L0C = 128 KB. Overflow caught at compile time. ^[30]

Additional HIVM passes (bishengir/lib/Dialect/HIVM/Transforms/): ^[30]

AutoInferBufferSize.cpp — automatic buffer size inference
SetBufferSize.cpp — buffer size assignment after inference

Confirmed by lit tests: plan-memory.mlir, hivm-auto-infer-buffer-size.mlir, hivm-set-buffer-size.mlir. ^[30]

Conclusion: UB allocation is fully automated — standard MLIR liveness analysis extended with Ascend-specific buffer reuse. GPU Triton’s AllocationAnalysis is not used; the entire path is Ascend-specific via HIVM dialect.

6. PyPTO 

6.1 PyPTO-main 

PyPTO (Parallel Tensor/Tile Operation) is a tile-based programming framework for Ascend NPU. The user writes tensor-level code; the compiler handles tiling, memory, sync, and pipelining automatically. ^[31]

Pipeline (C++ framework, not MLIR): ^[31]

Python (@pypto.frontend.jit) → Tensor Graph → Tile Graph → Block Graph
  → Execution Graph → PTO Virtual Instructions → CANN Backend → NPU binary

User-facing code — no sync, no memory management, no barriers:

@pypto.frontend.jit
def matmul_kernel(a: pypto.Tensor([], pypto.DT_FP32),
                  b: pypto.Tensor([], pypto.DT_FP32),
                  out: pypto.Tensor([], pypto.DT_FP32)):
    pypto.set_cube_tile_shapes([32, 32], [64, 64], [64, 64])
    out[:] = pypto.matmul(a, b, pypto.DT_FP32)

Sync Insertion 

Fully automated. The InsertSync pass at Block Graph level performs: ^[32]

RAW/WAW/WAR data dependency analysis using interval trees (DataDependencySearcher)
Pipe-aware set_flag/wait_flag injection across all pipeline pairs (AIC_MTE2, AIC_M, AIV_MTE2, AIV_V, AIV_MTE3, etc.)
Event ID allocation with deadlock detection and recovery
Cross-core sync for CV-fused kernels

Post-optimization: tune_sync_for_vf relaxes unnecessary barriers for vector-fusion subgraphs. ^[33]

Conclusion: User writes zero sync primitives. The compiler performs full pipeline-aware dependency analysis and barrier injection.

Ping-Pong 

Automated at Tile Graph level. The n_buffer_merge pass doubles buffers for overlapping data transfer and computation. ^[34] Combined with schedule_ooo (out-of-order scheduling) and add_alloc at Block Graph level for pipeline parallelism. ^[35]

The user controls tile shapes via pypto.set_vec_tile_shapes() or pypto.set_cube_tile_shapes() — this indirectly determines buffer count and pipeline depth. No explicit num_stages or T.Pipelined.

Conclusion: Ping-pong is fully automated. User only sets tile shapes.

UB Memory Allocation and Reuse 

Multi-stage memory management across the compilation pipeline:

Tile Graph: assign_memory_type — decides which memory level each tile uses (DDR, L1, L0A/B/C, UB) based on operation type ^[36]
Block Graph: add_alloc / remove_alloc — inserts buffer allocations ^[35]; memory_reuse/ contains liveness-based reuse passes including global_memory_reuse.cpp, merge_src_dst_buffer.cpp, and connection matrix analysis for non-overlapping buffer sharing ^[37]

Conclusion: UB allocation is fully automated — memory space assignment, buffer allocation, and liveness-based reuse are all compiler-managed. No user-visible memory APIs.

6.2 PyPTOv3 (redesign)

PyPTOv3 is a ground-up rewrite of PyPTO with a new IR system, new DSL, and multi-level abstraction. Active development at github.com/hw-native-sys/pypto. Targets A2 and A5. ^[38]

Pipeline — custom C++ AST-based IR for passes, MLIR as output format: ^[38]

Python (@pl.program / @pl.function) → Custom C++ IR → SSA conversion
  → passes (tensor→tile, memory, sync)
  → PTO codegen → MLIR text (.pto file: func.func, arith.*, memref)
  → ptoas (PTO assembler) → NPU binary

Key design difference from PyPTO-main: explicit memory spaces in the user API. The user specifies data movement targets — compiler handles the rest:

@pl.program
class MatmulExample:
    @pl.function
    def main(self, a: pl.Tensor[[M, K], pl.BF16],
                   b: pl.Tensor[[K, N], pl.BF16]) -> pl.Tensor[[M, N], pl.FP32]:
        a_l1 = pl.load(a, [0, 0], [32, 32], target_memory=pl.Mem.Mat)   # DDR → L1
        b_l1 = pl.load(b, [0, 0], [32, 32], target_memory=pl.Mem.Mat)
        a_l0a = pl.move(a_l1, target_memory=pl.Mem.Left)                 # L1 → L0A
        b_l0b = pl.move(b_l1, target_memory=pl.Mem.Right)                # L1 → L0B
        c_acc = pl.matmul(a_l0a, b_l0b)                                  # → L0C (Acc)
        return pl.store(c_acc, [0, 0], out)

Sync Insertion 

Fully automated via a 4-phase SyncInserter algorithm on the SSA IR: ^[39]

CollectSyncPairs — walks the AST tracking last writers/readers per MemRef. Detects RAW, WAW, WAR hazards across statements. Handles if/else branches via state merging, for loops via fixed-point iteration. Removes transitive and linear redundant pairs.
AdjustScopeCrossings — moves sync points when producer and consumer are in different control flow scopes (e.g. one inside if, other outside).
AssignEventIds — allocates from 8 hardware event IDs per pipe pair (EventIdManager), with position-based free tracking.
ApplyInsertions — mutates AST to insert set_flag/wait_flag calls at computed positions.

Pipeline assignment per op is backend-specific (backend->InferPipe(call)) — supports A2 and A5 with different pipe configurations. ^[39]

Cross-core sync for mixed AIC/AIV kernels via expand_mixed_kernel_pass. ^[40]

Conclusion: User writes no sync. The compiler performs full MemRef-level dependency analysis with scope-aware insertion and hardware event ID management.

Ping-Pong 

Handled by the optimization pipeline (Default strategy): ^[38]

InitMemRef — assigns memory spaces and inserts buffer allocations ^[41]
MemoryReuse — shares buffers with non-overlapping lifetimes ^[42]
LegalizePTOBufferReuse — legalizes buffer reuse for PTO backend ^[43]

No explicit num_stages or T.Pipelined. The Qwen3 decode example (pypto-lib PR #25) shows manual tile sizing for TILELET (2 KB vector) and TILE (16 KB cube) budgets — the user controls chunk sizes, the compiler handles pipelining. ^[44]

Conclusion: Pipelining is compiler-managed. User controls tile shapes to fit hardware budgets.

UB Memory Allocation and Reuse 

Three-pass approach: ^{[45] [41] [46]}

InferTileMemorySpace — infers memory space (Vec/Mat/Left/Right/Acc) for each tile based on operation semantics ^[45]
InitMemRef — creates MemRef descriptors with sizes, assigns concrete memory spaces, inserts alloc/free points ^[41]
AllocateMemoryAddr — assigns concrete byte addresses within each memory space ^[46]

Between steps 2 and 3, MemoryReuse performs liveness analysis and merges non-overlapping buffers into shared memory regions. ^[42]

Conclusion: UB allocation is fully automated — infer space, create buffers, reuse by liveness, assign addresses. User specifies target_memory on loads but never manages addresses or buffer sizes.

7. Pallas (JAX)

Pallas is Google’s kernel DSL for TPU (and GPU via Triton backend). Part of JAX. No decorator, no custom parser — kernels are regular Python functions passed to pl.pallas_call(), which triggers JAX tracing. ^[47]

Pipeline: ^[48]

pl.pallas_call(kernel_fn, grid, BlockSpecs)
  → JAX tracing (symbolic execution) → Jaxpr (JAX IR)
  → Pallas IR (Jaxpr + BlockSpec + grid)
  → Mosaic compiler (MLIR: vector + arith + tpu dialects)
  → TPU IR → LLO → TPU machine code

User-facing code — kernel is plain Python, pallas_call defines tiling:

def matmul_kernel(a_ref, b_ref, o_ref):
    o_ref[:, :] = jnp.dot(a_ref[:, :], b_ref[:, :])

result = pl.pallas_call(
    matmul_kernel,
    out_shape=jax.ShapeDtypeStruct((64, 64), jnp.float32),
    grid=(1,),
    in_specs=[pl.BlockSpec((64, 128), lambda i: (0, 0)),
              pl.BlockSpec((128, 64), lambda i: (0, 0))],
    out_specs=pl.BlockSpec((64, 64), lambda i: (0, 0)),
)(a, b)

Sync Insertion 

TPU has two compute units (MXU for matmul, VPU for vector ops) but they share a single instruction stream — no parallel pipeline hazards between them. ^[49]

Pipelining on TPU is DMA vs compute overlap: the DMA engine loads the next tile while compute processes the current one. This is the only source of concurrency within a core. ^[49]

Automatic (default): user writes pallas_call with BlockSpec. Mosaic compiler generates all DMA operations, semaphores, and double-buffer swaps from the BlockSpec — user never writes sync code: ^[47]

def kernel(a_ref, b_ref, o_ref):
    o_ref[:, :] = a_ref[:, :] + b_ref[:, :]  # no DMA, no sync

pl.pallas_call(kernel, grid=(N,),
    in_specs=[pl.BlockSpec((128,), lambda i: (i,))],
    out_specs=pl.BlockSpec((128,), lambda i: (i,)),
    ...)(a, b)
# Mosaic: BlockSpec → DMA schedule → semaphore insertion → double buffering

Manual (expert): for advanced patterns (paged attention, custom prefetch), user writes pltpu.async_copy + semaphores directly: ^[50]

sem = pltpu.SemaphoreType.DMA((2,))
pltpu.async_copy(hbm_ref.at[slice], vmem_buffer, sem).wait()

Compare: Ascend has 3+ parallel units (MTE2, Cube, Vector, MTE3) requiring fine-grained set_flag/wait_flag between each pair. TPU has 2 (DMA + compute).

Conclusion: Sync is compiler-generated from BlockSpec by default. Manual semaphores available as escape hatch. Simpler than Ascend — only DMA↔compute overlap, no Cube↔Vector hazards.

Ping-Pong 

TPU on-chip memory: VMEM (Vector Memory) — 32 MB SRAM per core, equivalent to Ascend’s UB but ~170× larger (32 MB vs 192 KB). ^[49]

Default: 2-level double buffering for all inputs and outputs, generated automatically by Mosaic from BlockSpec. Sufficient for most kernels because VMEM is large enough that double buffering rarely causes pressure. ^[49]

Advanced control via pltpu.emit_pipeline — supports nested pipelines and lookahead prefetch (fetch next block as soon as a buffer slot is free, not just one iteration ahead). ^[51]

Conclusion: Ping-pong is automatic (2-stage default from BlockSpec). Advanced pipelining API available but rarely needed due to large VMEM.

UB Memory Allocation and Reuse 

TPU memory hierarchy: ^[49]

HBM — main memory (GBs), slow
VMEM — 32 MB vector SRAM per core, holds tiles during compute
VREG — vector registers (8×128 tiles for fp32), fastest
SMEM — ~4-8 KB scalar SRAM, for control/metadata only

User controls tiling via BlockSpec(shape, index_map) — this determines how much VMEM each tile consumes. The compiler handles VMEM allocation internally. ^[47]

Tiling constraints: last 2 dimensions must be divisible by 8 and 128 respectively (matching 8×128 vector register shape). ^[49]

No user-visible VMEM address management — unlike Ascend’s manual UB offset planning. If total VMEM usage exceeds 32 MB (including double-buffer overhead), compilation fails.

Conclusion: VMEM allocation is fully compiler-managed. User only controls tile shapes via BlockSpec. The 32 MB budget is ~170× Ascend’s UB — memory pressure is rarely an issue on TPU.

8. Mojo (Modular)

Mojo is a Python superset by Modular (Chris Lattner). Systems-level language with explicit GPU kernel programming. Uses MLIR internally via KGEN compiler. ^[52]

Pipeline: ^[52]

Mojo source → KGEN compiler (MLIR-based) → platform-specific backend
  → PTX (NVIDIA) / HIP (AMD) / Metal (Apple) → GPU binary

Abstraction level: between CUDA and Triton. Thread-centric like CUDA (explicit threads, blocks, sync), but with tile-level abstractions (TileTensor, TileIO, TilePipeline) for structured kernels. ^[53]

Supports NVIDIA, AMD, and Apple GPUs. No TPU or Ascend support.

Simple kernel — CUDA-like thread model, Python-like syntax: ^[55]

def scalar_add(vector: UnsafePointer[Float32], size: Int, scalar: Float32):
    idx = block_idx.x * block_dim.x + thread_idx.x
    if idx < size:
        vector[idx] += scalar

ctx.enqueue_function[scalar_add](device_buffer, 20, Float32(5.0),
                                  grid_dim=1, block_dim=20)

Sync Insertion 

Manual. User writes explicit barriers — closer to CUDA than to Triton: ^[54]

barrier() — block-level sync (like CUDA __syncthreads())
syncwarp() — warp-level sync
named_barrier(id, count) — multiple independent barriers

Platform-specific differences managed by user: ^[53]

NVIDIA Blackwell: hardware mbarrier with automatic byte counting for TMA
AMD MI300X: no mbarrier — explicit atomic counters with s_sleep 0 yield

Structured kernel pattern (“three pillars”) separates concerns: ^[53]

TileIO — data movement between memory levels
TilePipeline — producer-consumer with acquire/release semantics
TileOp — MMA/compute instructions

# Three warp roles execute in parallel:
# Load warp:    acquires producer stage → TMA load → signals completion
# MMA warp:     waits on input pipeline → computes → signals output pipeline
# Epilogue warp: reads from TMEM → writes to global memory
with producer.acquire() as tiles:
    tile_io.load(tiles, barrier, k_coord)  # explicit DMA + barrier

Conclusion: Sync is fully manual. The “three pillars” pattern provides structure but the user orchestrates all barriers. No compiler-automated sync insertion.

Ping-Pong 

Manual. Mojo intentionally rejects Triton-style automation — their position: automatic compilation improves accessibility but hits a performance ceiling for production inference, forcing developers back to low-level code. ^[56]

Instead, the TilePipeline abstraction provides structured manual control via acquire/release context managers: ^[53]

with producer.acquire() as tiles:
    tile_io.load(tiles, barrier, k_coord)   # load next tile
# release: signals consumer that data is ready

with consumer.acquire() as tiles:
    tile_op.mma(tiles)                       # compute on current tile
# release: signals producer that buffer is free

Three warp roles execute in parallel: load warp (DMA), MMA warp (compute), epilogue warp (store). User assigns warps to roles explicitly. ^[53]

Result: 48% less code than CUTLASS with identical performance, but still fully user-orchestrated. ^[56]

Conclusion: Pipelining is user-structured via TilePipeline, not compiler-generated. Design choice: structured manual control over automation with escape hatches.

UB Memory Allocation and Reuse 

Manual. User declares memory explicitly: ^[55]

Shared memory: Shared type for block-visible scratchpad
Registers: small compile-time-sized arrays auto-allocated by compiler
Global memory: DeviceBuffer types

TileTensor carries compile-time layout information (shape, stride, swizzle) and memory address space placement — enabling the compiler to verify layouts at compile time without managing allocation itself. ^[57]

Platform-specific memory handled by user: ^[53]

NVIDIA Blackwell: TMEM (256 KB) managed via explicit context managers — entering signals readiness, exiting deallocates
AMD MI300X: accumulators live in VGPRs, compiler manages register allocation

No liveness-based buffer reuse by the compiler. User controls all allocation and lifetime.

Conclusion: Memory is user-managed. TileTensor provides type-safe layout verification but not automatic allocation or reuse. Consistent with Mojo’s philosophy: structured control over automation.

9. Comparison 

	AscendC	Triton	cuTile	TileLang-Ascend	Triton-Ascend	PyPTO-main	PyPTOv3	Pallas	Mojo
Abstraction level	Instruction	Tile/block	Tile	Tile (hybrid)	Tile/block	Tensor	Tensor + Tile (expert)	Tile (BlockSpec)	Thread + Tile (structured)
Sync insertion	Manual	Auto	Auto (no escape)	Hybrid (auto opt-in)	Auto (HIVM)	Auto	Auto (4-phase)	Auto (Mosaic from BlockSpec)	Manual (three pillars)
Ping-pong	Manual	Auto (`num_stages`)	Auto (no hint)	`T.Pipelined(num_stages)`	Auto (multiBuffer)	Auto	Auto	Auto (2-stage default)	Manual (`TilePipeline`)
On-chip memory	UB 192 KB	Shared ~228 KB	Shared + TMEM 256 KB	UB 192 KB	UB 192 KB	UB 192 KB	UB 192 KB	VMEM 32 MB	Shared (GPU)
Memory mgmt	Manual	Auto (AllocationAnalysis)	Auto	Hybrid (opt-in planning)	Auto (PlanMemory/HIVM)	Auto	Auto (3-pass)	Auto (BlockSpec)	Manual (TileTensor)
Overflow detection	Silent corruption	Compile-time error	Compile-time error	Partial (opt-in)	Compile-time error	?	?	Compile-time error	Compile-time (layout)
IR type	C++ (AscendC)	MLIR (TTIR/TTGIR)	MLIR (TileIR)	TVM TIR	MLIR (TTIR→HIVM)	Custom C++ graphs	Custom C++ AST	Jaxpr → Mosaic (MLIR)	MLIR (KGEN)
User memory control	Full manual	None	None	Explicit hierarchy	Indirect (BLOCK_SIZE)	None	`target_memory` hint	BlockSpec (tile shape)	Full manual (Shared, TMEM)
Hardware	Ascend A2/A3	NVIDIA/AMD/Intel	NVIDIA Blackwell	Ascend A2/A3	Ascend A2/A3	Ascend A2	A2, A5	Google TPU	NVIDIA/AMD/Apple

Key Observations for PyAsc2 Design 

1. Full automation is achievable on Ascend — performance cost varies. Among Ascend-targeting frameworks, Triton-Ascend, TileLang-Ascend, PyPTO-main, and PyPTOv3 all automate sync, ping-pong, and UB allocation. But automation quality differs: TileLang-Ascend’s auto sync generates conservative PipeBarrier<PIPE_ALL> instead of fine-grained flags. No published benchmarks compare automated vs hand-optimized AscendC for the same kernels. The gap between “correct” and “optimal” automatic code is the main engineering challenge for PyAsc2.

2. The right abstraction level is tensor — with tile escape hatch. PyPTOv3’s two-level design covers both use cases. Triton’s tile-only model forces all users into low-level thinking. cuTile’s no-escape model limits advanced optimization. PyAsc2 should follow PyPTOv3’s approach: tensor-level default, tile-level available.

3. Memory hierarchy hints > full manual control. cuTile (fully implicit) and AscendC (fully explicit) are the two extremes. PyPTOv3’s target_memory=pl.Mem.Mat is the sweet spot — the user declares intent, the compiler handles mechanics.

4. Liveness-based UB reuse is critical — more so than on GPU. On GPU, shared memory reuse pressure is low — up to 228 KB available, rarely exhausted. On Ascend, UB is 192 KB (96 KB with double buffering) and almost always the bottleneck. Suboptimal reuse forces smaller tiles or spills to global memory. PyAsc2 should provide automatic reuse by default with an option for expert override of buffer placement.

5. Compile-time UB overflow detection is non-negotiable. AscendC’s silent corruption is the #1 developer pain point. Every modern framework catches overflow at compile time. PyAsc2 must validate UB/L1/L0 fit at compile time with clear error messages.

6. num_stages vs fully automatic pipelining. Triton and TileLang-Ascend expose pipeline depth to users. cuTile, PyPTO-main, and PyPTOv3 hide it entirely. For PyAsc2: default to automatic, provide num_stages as optional expert hint.

7. Ascend-native IR outperforms adapted GPU IR. Triton-Ascend’s TTIR→HIVM path works but requires user-level rewrites (BLOCK_SIZE_SUB, fixed grid, alignment). Native Ascend frameworks handle these in the compiler. PyAsc2 should be Ascend-native from the start.

References 

[10] Triton matmul tutorial — pipelined kernel with no user-written barriers.
https://triton-lang.org/main/getting-started/tutorials/03-matrix-multiplication.html

[11] Triton shared memory allocation — AllocationAnalysis.
https://github.com/triton-lang/triton/blob/main/lib/Analysis/Allocation.cpp

[12] Triton encoding assignment for reductions — ReduceDataDuplication.
https://github.com/triton-lang/triton/blob/main/lib/Dialect/TritonGPU/Transforms/ReduceDataDuplication.cpp

[13] Triton encoding assignment for dot/matmul — BlockedToMMA.
https://github.com/triton-lang/triton/blob/main/lib/Dialect/TritonGPU/Transforms/AccelerateMatmul.cpp

[14] cuTile execution model — no intra-block sync.
https://docs.nvidia.com/cuda/cutile-python/execution.html

[15] TileIR internals — from cuTile to MLIR/LLVM to SASS.
https://maknee.github.io/blog/2026/NVIDIA-TileIR-Internals-from-CuTile-to-MLIR-LLVM-to-SASS/

[16] TileIR source code (NVIDIA, open-source).
https://github.com/NVIDIA/cuda-tile

[17] NVIDIA blog — high-performance matrix multiply in cuTile.
https://developer.nvidia.com/blog/how-to-write-high-performance-matrix-multiply-in-nvidia-cuda-tile

[18] TiledAttention — SDPA in cuTile (arXiv:2603.01960).
https://arxiv.org/abs/2603.01960

[19] CuTile on Blackwell — compiler moat analysis (TMEM, mbarrier details).
https://patricktoulme.substack.com/p/cutile-on-blackwell-nvidias-compiler

[20] tcgen05 for dummies — TMEM allocation, operand placement.
https://gau-nernst.github.io/tcgen05/

[21] TileLang-Ascend repo — examples and sync primitives.
https://github.com/tile-ai/tilelang-ascend

[22] TileLang-Ascend auto sync config (Issue #98).
https://github.com/tile-ai/tilelang-ascend/issues/98

[23] TileLang-Ascend CombineCV + AscendSyncInsert passes.
https://github.com/tile-ai/tilelang-ascend/blob/ascendc_pto/src/transform/ascend_combinecv.cc

[24] TileLang-Ascend double buffer sync issue (#110).
https://github.com/tile-ai/tilelang-ascend/issues/110

[25] TileLang-Ascend roadmap — memory planning, autotuner (Issue #3).
https://github.com/tile-ai/tilelang-ascend/issues/3

[26] Triton-Ascend programming guide (gitcode.com primary repo).
https://gitcode.com/Ascend/triton-ascend/blob/main/docs/en/programming_guide.md

[27] Triton-Ascend examples — confirmed on Ascend with torch_npu.
https://gitcode.com/Ascend/triton-ascend/tree/main/docs/en/examples

[28] Triton-Ascend architecture design and core features.
https://gitcode.com/Ascend/triton-ascend/blob/main/docs/en/architecture_design_and_core_features.md

[29] AscendNPU-IR architecture — HIVM/HFusion/HACC dialect definitions.
https://gitcode.com/Ascend/AscendNPU-IR/blob/main/docs/source/en/introduction/architecture.md

[30] AscendNPU-IR PlanMemory pass — liveness-based UB allocation.
https://gitcode.com/Ascend/AscendNPU-IR/blob/main/bishengir/lib/Dialect/HIVM/Transforms/PlanMemory.cpp

[31] PyPTO-main matmul example.
https://gitcode.com/cann/pypto/blob/master/examples/01_beginner/compute/matmul_ops.py

[32] PyPTO InsertSync pass — RAW/WAW/WAR dependency analysis.
https://gitcode.com/cann/pypto/blob/master/framework/src/passes/block_graph_pass/insert_sync.cpp

[33] PyPTO tune_sync_for_vf — barrier relaxation for vector fusion.
https://gitcode.com/cann/pypto/blob/master/framework/src/passes/block_graph_pass/tune_sync_for_vf.cpp

[34] PyPTO n_buffer_merge — double buffering at Tile Graph level.
https://gitcode.com/cann/pypto/blob/master/framework/src/passes/tile_graph_pass/graph_partition/n_buffer_merge.cpp

[35] PyPTO add_alloc / schedule_ooo — Block Graph allocation and scheduling.
https://gitcode.com/cann/pypto/blob/master/framework/src/passes/block_graph_pass/schedule_ooo/add_alloc.cpp

[36] PyPTO assign_memory_type — memory space assignment at Tile Graph.
https://gitcode.com/cann/pypto/blob/master/framework/src/passes/tile_graph_pass/data_path/assign_memory_type.cpp

[37] PyPTO memory_reuse — liveness-based buffer reuse.
https://gitcode.com/cann/pypto/blob/master/framework/src/passes/block_graph_pass/memory_reuse/global_memory_reuse.cpp

[38] PyPTOv3 language guide — DSL, memory hierarchy, optimization pipeline.
https://github.com/hw-native-sys/pypto/blob/main/docs/en/user/01-language_guide.md

[39] PyPTOv3 insert_sync_pass — 4-phase sync insertion algorithm.
https://github.com/hw-native-sys/pypto/blob/main/src/ir/transforms/insert_sync_pass.cpp

[40] PyPTOv3 expand_mixed_kernel_pass — AIC/AIV split + cross-core sync.
https://github.com/hw-native-sys/pypto/blob/main/src/ir/transforms/expand_mixed_kernel_pass.cpp

[41] PyPTOv3 init_memref — buffer allocation and memory space assignment.
https://github.com/hw-native-sys/pypto/blob/main/src/ir/transforms/init_memref.cpp

[42] PyPTOv3 memory_reuse_pass — liveness-based buffer sharing.
https://github.com/hw-native-sys/pypto/blob/main/src/ir/transforms/memory_reuse_pass.cpp

[43] PyPTOv3 legalize_pto_buffer_reuse — PTO backend buffer legalization.
https://github.com/hw-native-sys/pypto/blob/main/src/ir/transforms/legalize_pto_buffer_reuse_pass.cpp

[44] PyPTOv3 Qwen3 decode example — tilelet-aware tiling.
https://github.com/hw-native-sys/pypto-lib/pull/25

[45] PyPTOv3 infer_tile_memory_space — memory space inference.
https://github.com/hw-native-sys/pypto/blob/main/src/ir/transforms/infer_tile_memory_space_pass.cpp

[46] PyPTOv3 allocate_memory_addr — concrete address assignment.
https://github.com/hw-native-sys/pypto/blob/main/src/ir/transforms/allocate_memory_addr_pass.cpp

[47] Pallas — JAX kernel language overview.
https://jax.readthedocs.io/en/latest/pallas/index.html

[48] Pallas design notes — compilation pipeline.
https://jax.readthedocs.io/en/latest/pallas/design/design.html

[49] Pallas TPU details — memory, sync, tiling constraints.
https://docs.jax.dev/en/latest/pallas/tpu/details.html

[50] Pallas paged attention kernel — async_copy + semaphore usage.
https://github.com/jax-ml/jax/blob/main/jax/experimental/pallas/ops/tpu/paged_attention/paged_attention_kernel.py

[51] Pallas TPU pipelining — emit_pipeline, lookahead prefetch.
https://docs.jax.dev/en/latest/pallas/tpu/pipelining.html

[52] Mojo MLIR-based compilation — KGEN compiler (arXiv:2509.21039).
https://arxiv.org/abs/2509.21039

[53] Structured Mojo Kernels Part 2 — three pillars (TileIO, TilePipeline, TileOp).
https://www.modular.com/blog/structured-mojo-kernels-part-2-the-three-pillars

[54] Mojo GPU sync primitives — barrier, syncwarp, named_barrier.
https://docs.modular.com/mojo/manual/gpu/block-and-warp/

[55] Mojo GPU fundamentals — kernel model, thread indexing.
https://docs.modular.com/mojo/manual/gpu/fundamentals/

[56] Structured Mojo Kernels Part 1 — design philosophy, performance.
https://www.modular.com/blog/structured-mojo-kernels-part-1-peak-performance-half-the-code

[57] TileTensor — parametric tile-level tensors in Mojo.
https://www.modular.com/blog/tiletensor-part-1-safer-more-efficient-gpu-kernels

[58] CANN 9 SDK preview — 950/A5 programming model findings (internal notes from SDK source inspection).