Repository navigation
[Bug] auto-sync should track aliasing through pto.subview #536
Description
Activity
p_fp16_half = pto.subset(
p_fp16,
offsets=[row_off_sb, 0],
sizes=[S0_HALF, S1_TILE],
)pto.subsetwill be deprecated. Trypto.SubviewOpinstead here, it's supposed to work.Thanks for reporting this.
I re-validated on current
main(commit123b664) with the same dependency shape but usingpto.subview(as suggested in-thread), and I can no longer reproduce the “missing auto-sync” behavior from the originalpto.subsetreport.What I tested:
- Static half-split subviews (
[0:16, :]and[16:32, :]) writing subviews then reading parent tile. - Dynamic row-offset subview (
row_off = subblock_idx * 16) writing subview then reading parent tile.
In both cases, auto-sync now emits ordering ops around the hazard path, e.g.:
pipe_barrier(PIPE_V)set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0)wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0)
So for the subview-based path, this issue appears resolved on latest main.
Notes:
pto.subsetis currently not recognized in this path (parse error), consistent with deprecation/migration topto.subview.- If needed, I can upstream the exact reproducer as a lit regression so this behavior stays guarded.
- Static half-split subviews (
Thanks for reporting this.
I re-validated on current
main(commit123b664) with the same dependency shape but usingpto.subview(as suggested in-thread), and I can no longer reproduce the “missing auto-sync” behavior from the originalpto.subsetreport.What I tested:
* Static half-split subviews (`[0:16, :]` and `[16:32, :]`) writing subviews then reading parent tile. * Dynamic row-offset subview (`row_off = subblock_idx * 16`) writing subview then reading parent tile.In both cases, auto-sync now emits ordering ops around the hazard path, e.g.:
* `pipe_barrier(PIPE_V)` * `set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0)` * `wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0)`So for the subview-based path, this issue appears resolved on latest main.
Notes:
* `pto.subset` is currently not recognized in this path (parse error), consistent with deprecation/migration to `pto.subview`. * If needed, I can upstream the exact reproducer as a lit regression so this behavior stays guarded.I forgot to add that I'm running the nightly build 0.28, your fixes are probably not included in that version right? I wrote also a minimal reproducer where the runtime result is: top half = expected, bottom half = uninitialized garbage.
module { func.func @subview_shared_repro(%arg0: !pto.ptr<f32>) { %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index %c32 = arith.constant 32 : index %c256 = arith.constant 256 : index %0 = pto.get_subblock_idx %1 = arith.index_cast %0 : i64 to index %c16 = arith.constant 16 : index %2 = arith.muli %1, %c16 : index pto.section.vector { %3 = pto.alloc_tile : !pto.tile_buf<vec, 32x256xf32> %4 = pto.subview %3[%2, %c0] sizes [16, 256] : !pto.tile_buf<vec, 32x256xf32> -> !pto.tile_buf<vec, 16x256xf32> %5 = arith.cmpi eq, %1, %c0 : index %cst = arith.constant 1.000000e+00 : f32 %cst_0 = arith.constant 2.000000e+00 : f32 %6 = arith.select %5, %cst, %cst_0 : f32 %cst_1 = arith.constant 0.000000e+00 : f32 pto.tmuls ins(%4, %cst_1 : !pto.tile_buf<vec, 16x256xf32>, f32) outs(%4 : !pto.tile_buf<vec, 16x256xf32>) pto.tadds ins(%4, %6 : !pto.tile_buf<vec, 16x256xf32>, f32) outs(%4 : !pto.tile_buf<vec, 16x256xf32>) pto.barrier <PIPE_V> %7 = arith.cmpi eq, %1, %c0 : index scf.if %7 { %8 = pto.make_tensor_view %arg0, shape = [%c32, %c256], strides = [%c256, %c1] : !pto.tensor_view<?x?xf32> %9 = pto.partition_view %8, offsets = [%c0, %c0], sizes = [%c32, %c256] : !pto.tensor_view<?x?xf32> -> !pto.partition_tensor_view<32x256xf32> pto.tstore ins(%3 : !pto.tile_buf<vec, 32x256xf32>) outs(%9 : !pto.partition_tensor_view<32x256xf32>) } } return } }
Generated cpp using:
ptoas --enable-insert-sync ./kernel.pto -o ./kernel.cpp
and compiled using:
bisheng \ -I${PTO_LIB_PATH}/include \ -fPIC -shared -D_FORTIFY_SOURCE=2 -O2 -std=c++17 \ -Wno-macro-redefined -Wno-ignored-attributes -fstack-protector-strong \ -xcce -Xhost-start -Xhost-end \ -mllvm -cce-aicore-stack-size=0x8000 \ -mllvm -cce-aicore-function-stack-size=0x8000 \ -mllvm -cce-aicore-record-overflow=true \ -mllvm -cce-aicore-addr-transform \ -mllvm -cce-aicore-dcci-insert-for-scalar=false \ --npu-arch=dav-2201 -DMEMORY_BASE \ -std=gnu++17 \ -DKERNEL_CPP="\"${SCRIPT_DIR}/kernel.cpp\"" \ ./caller.cpp \ -o ./lib.so#include "pto/pto-inst.hpp" using namespace pto; enum class PTOAutoSyncTailMode : int { kBarrierAll = 0, kSetWaitMte3ToSEvent0 = 1, }; static AICORE inline void ptoas_auto_sync_tail( PTOAutoSyncTailMode mode = PTOAutoSyncTailMode::kBarrierAll) { switch (mode) { case PTOAutoSyncTailMode::kSetWaitMte3ToSEvent0: set_flag(PIPE_MTE3, PIPE_S, EVENT_ID0); wait_flag(PIPE_MTE3, PIPE_S, EVENT_ID0); break; case PTOAutoSyncTailMode::kBarrierAll: default: pipe_barrier(PIPE_ALL); break; } } __global__ AICORE void subview_shared_repro(__gm__ float* v1) { unsigned v2 = 1; unsigned v3 = 256; unsigned v4 = 0; const int32_t v5 = 256; const int32_t v6 = 1; const int32_t v7 = 0; const int32_t v8 = 16; const float v9 = 1.0f; const float v10 = 2.0f; const float v11 = 0.0f; const int64_t v12 = 0; using T = float; int64_t v13 = get_subblockid(); int32_t v14 = (int32_t) ((int64_t) v13); #if defined(__DAV_VEC__) set_mask_norm(); set_vector_mask(-1, -1); Tile<TileType::Vec, float, 32, 256, BLayout::RowMajor, 32, 256, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> v15; TASSIGN(v15, v12); __ubuf__ float* v16 = v15.data(); __ubuf__ float* v17 = v16 + (v4 + (unsigned) ((int32_t) (uint32_t) v14 * (uint32_t) v8) * v3 + v4 * v2); __ubuf__ float* v18 = (__ubuf__ float*) v17; Tile<TileType::Vec, float, 32, 256, BLayout::RowMajor, 16, 256, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> v19; uint64_t v20 = reinterpret_cast<uint64_t>(v17); TASSIGN(v19, v20); bool v21 = v14 == v7; float v22 = v21 ? v9 : v10; TMULS(v19, v19, v11); pipe_barrier(PIPE_V); TADDS(v19, v19, v22); set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); pipe_barrier(PIPE_V); wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); if (v21) { pto::Shape<1, 1, 1, 32, 256> v23 = pto::Shape<1, 1, 1, 32, 256>(); pto::Stride<8192, 8192, 8192, 256, 1> v24 = pto::Stride<8192, 8192, 8192, 256, 1>(); GlobalTensor<float, pto::Shape<1, 1, 1, 32, 256>, pto::Stride<8192, 8192, 8192, 256, 1>, pto::Layout::ND> v25 = GlobalTensor<float, pto::Shape<1, 1, 1, 32, 256>, pto::Stride<8192, 8192, 8192, 256, 1>, pto::Layout::ND>(v1 + (v4 + v4 * (unsigned) v5 + v4 * (unsigned) v6), v23, v24); TSTORE(v25, v15); } #endif // __DAV_VEC__ ptoas_auto_sync_tail(PTOAutoSyncTailMode::kBarrierAll); return; }
UPDATE: this works if I only let AIV write and store. The MLIR is generated by the following DSL:
# Single-AIV control test for subview_shared_repro. # # Same parent allocation, same two subviews, same store — but only one AIV # does everything (the other does nothing). This isolates the per-AIV- # private-alloc_tile issue (Bug #2) from any subview mechanics: # # * If THIS test passes → subview works fine inside one AIV; the only # failure mode is cross-AIV sharing. # * If THIS test also fails → subview itself is broken even within a # single AIV. # # Test: # On AIV0 only: # - alloc parent [32, 256] # - subview_top = parent[ 0..16, :] ← write 1.0 # - subview_bottom = parent[16..32, :] ← write 2.0 # - store full parent to GM # AIV1 does nothing. # # Expected: GM rows 0..16 = 1.0, rows 16..32 = 2.0. from ptodsl import pto, tile, to_ir_module from ptodsl import scalar as s const = s.const ROWS = 32 HALF = ROWS // 2 COLS = 256 def meta_data(): fp32 = pto.float32 ptr_fp32 = pto.PtrType(fp32) o_tensor_ty = pto.TensorType(rank=2, dtype=fp32) o_sub_ty = pto.SubTensorType(shape=[ROWS, COLS], dtype=fp32) parent_ty = pto.TileBufType( shape=[ROWS, COLS], dtype=fp32, memory_space="VEC" ) return locals() @to_ir_module(meta_data=meta_data) def subview_single_aiv( gm_o: "ptr_fp32", ) -> None: c0 = const(0) c1 = const(1) cROWS = const(ROWS) cCOLS = const(COLS) cHALF = const(HALF) sb_idx = s.index_cast(pto.get_subblock_idx()) is_aiv0 = s.eq(sb_idx, c0) with pto.vector_section(): # Both AIVs allocate (alloc_tile is per-AIV-private), but only AIV0 # writes and stores. AIV1's parent stays unused. parent = pto.alloc_tile(parent_ty) with pto.if_context(is_aiv0): # Two subviews of the SAME AIV0-private parent. top = pto.subview( parent, offsets=[c0, c0], sizes=[HALF, COLS] ) bot = pto.subview( parent, offsets=[cHALF, c0], sizes=[HALF, COLS] ) # Fill top half with 1.0, bottom half with 2.0. tile.muls(top, const(0.0, s.float32), top) tile.adds(top, const(1.0, s.float32), top) tile.muls(bot, const(0.0, s.float32), bot) tile.adds(bot, const(2.0, s.float32), bot) pto.pipe_barrier("V") # Store full parent. tv_o = pto.as_tensor( o_tensor_ty, ptr=gm_o, shape=[cROWS, cCOLS], strides=[cCOLS, c1], ) o_view = pto.slice_view( o_sub_ty, source=tv_o, offsets=[c0, c0], sizes=[cROWS, cCOLS], ) pto.store(parent, o_view) if __name__ == "__main__": print(subview_single_aiv)
module { func.func @subview_single_aiv(%arg0: !pto.ptr<f32>) { %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index %c32 = arith.constant 32 : index %c256 = arith.constant 256 : index %c16 = arith.constant 16 : index %0 = pto.get_subblock_idx %1 = arith.index_cast %0 : i64 to index %2 = arith.cmpi eq, %1, %c0 : index pto.section.vector { %3 = pto.alloc_tile : !pto.tile_buf<vec, 32x256xf32> scf.if %2 { %4 = pto.subview %3[%c0, %c0] sizes [16, 256] : !pto.tile_buf<vec, 32x256xf32> -> !pto.tile_buf<vec, 16x256xf32> %5 = pto.subview %3[%c16, %c0] sizes [16, 256] : !pto.tile_buf<vec, 32x256xf32> -> !pto.tile_buf<vec, 16x256xf32> %cst = arith.constant 0.000000e+00 : f32 pto.tmuls ins(%4, %cst : !pto.tile_buf<vec, 16x256xf32>, f32) outs(%4 : !pto.tile_buf<vec, 16x256xf32>) %cst_0 = arith.constant 1.000000e+00 : f32 pto.tadds ins(%4, %cst_0 : !pto.tile_buf<vec, 16x256xf32>, f32) outs(%4 : !pto.tile_buf<vec, 16x256xf32>) %cst_1 = arith.constant 0.000000e+00 : f32 pto.tmuls ins(%5, %cst_1 : !pto.tile_buf<vec, 16x256xf32>, f32) outs(%5 : !pto.tile_buf<vec, 16x256xf32>) %cst_2 = arith.constant 2.000000e+00 : f32 pto.tadds ins(%5, %cst_2 : !pto.tile_buf<vec, 16x256xf32>, f32) outs(%5 : !pto.tile_buf<vec, 16x256xf32>) pto.barrier <PIPE_V> %6 = pto.make_tensor_view %arg0, shape = [%c32, %c256], strides = [%c256, %c1] : !pto.tensor_view<?x?xf32> %7 = pto.partition_view %6, offsets = [%c0, %c0], sizes = [%c32, %c256] : !pto.tensor_view<?x?xf32> -> !pto.partition_tensor_view<32x256xf32> pto.tstore ins(%3 : !pto.tile_buf<vec, 32x256xf32>) outs(%7 : !pto.partition_tensor_view<32x256xf32>) } } return } }
with generated cpp:
#include "pto/pto-inst.hpp" using namespace pto; enum class PTOAutoSyncTailMode : int { kBarrierAll = 0, kSetWaitMte3ToSEvent0 = 1, }; static AICORE inline void ptoas_auto_sync_tail( PTOAutoSyncTailMode mode = PTOAutoSyncTailMode::kBarrierAll) { switch (mode) { case PTOAutoSyncTailMode::kSetWaitMte3ToSEvent0: set_flag(PIPE_MTE3, PIPE_S, EVENT_ID0); wait_flag(PIPE_MTE3, PIPE_S, EVENT_ID0); break; case PTOAutoSyncTailMode::kBarrierAll: default: pipe_barrier(PIPE_ALL); break; } } __global__ AICORE void subview_single_aiv(__gm__ float* v1) { unsigned v2 = 16; unsigned v3 = 1; unsigned v4 = 256; unsigned v5 = 0; const int32_t v6 = 256; const int32_t v7 = 1; const int32_t v8 = 0; const float v9 = 0.0f; const float v10 = 1.0f; const float v11 = 2.0f; const int64_t v12 = 0; using T = float; int64_t v13 = get_subblockid(); #if defined(__DAV_VEC__) set_mask_norm(); set_vector_mask(-1, -1); Tile<TileType::Vec, float, 32, 256, BLayout::RowMajor, 32, 256, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> v14; TASSIGN(v14, v12); if ((int32_t) ((int64_t) v13) == v8) { __ubuf__ float* v15 = v14.data(); __ubuf__ float* v16 = v15 + (v5 + v5 * v4 + v5 * v3); __ubuf__ float* v17 = (__ubuf__ float*) v16; Tile<TileType::Vec, float, 32, 256, BLayout::RowMajor, 16, 256, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> v18; uint64_t v19 = reinterpret_cast<uint64_t>(v16); TASSIGN(v18, v19); __ubuf__ float* v20 = v14.data(); __ubuf__ float* v21 = v20 + (v5 + v2 * v4 + v5 * v3); __ubuf__ float* v22 = (__ubuf__ float*) v21; Tile<TileType::Vec, float, 32, 256, BLayout::RowMajor, 16, 256, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> v23; uint64_t v24 = reinterpret_cast<uint64_t>(v21); TASSIGN(v23, v24); TMULS(v18, v18, v9); pipe_barrier(PIPE_V); TADDS(v18, v18, v10); TMULS(v23, v23, v9); pipe_barrier(PIPE_V); TADDS(v23, v23, v11); set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); pipe_barrier(PIPE_V); pto::Shape<1, 1, 1, 32, 256> v25 = pto::Shape<1, 1, 1, 32, 256>(); pto::Stride<8192, 8192, 8192, 256, 1> v26 = pto::Stride<8192, 8192, 8192, 256, 1>(); GlobalTensor<float, pto::Shape<1, 1, 1, 32, 256>, pto::Stride<8192, 8192, 8192, 256, 1>, pto::Layout::ND> v27 = GlobalTensor<float, pto::Shape<1, 1, 1, 32, 256>, pto::Stride<8192, 8192, 8192, 256, 1>, pto::Layout::ND>(v1 + (v5 + v5 * (unsigned) v6 + v5 * (unsigned) v7), v25, v26); wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); TSTORE(v27, v14); } #endif // __DAV_VEC__ ptoas_auto_sync_tail(PTOAutoSyncTailMode::kBarrierAll); return; }
- changed the title
[-][Bug] auto-sync should track aliasing through `pto.subset`[/-][+][Bug] auto-sync should track aliasing through `pto.subview`[/+]on Apr 21, 2026 In Split mode,you does not need subview to get the half part of parent tile. Please refer to https://github.com/hw-native-sys/PTOAS/blob/main/test/samples/TPushTPop/test4/kernel.pto
Metadata
Metadata
Assignees
Labels
Type
Projects
- Status
Component
PTO Dialect / ODS (include/PTO/IR)
Description
The auto-sync pass treats the result of
pto.subsetas an independentbuffer rather than an alias of its source tile. As a consequence, no
sync flag (
set_flag/wait_flagpair) is emitted between a writeto a sub-view and a later read/write of either the parent tile or an
overlapping sub-view. Kernels that use
pto.subsetto expresssub-block parallelism (each AIV sub-block owning a row-half of a
shared tile) produce wrong data.
Reproduction (minimal)
Vector kernel pattern (simplified):
The RAW dependency from
tile.cvt(... p_fp16_half)topto.tpush_to_aic(p_fp16, 0)is invisible to the auto-sync passbecause
p_fp16_halfandp_fp16are distinct SSA values. No flag isemitted; the push reads stale data.
Manual
pto.pipe_barrier("V")calls inserted at every cross-sub-blockseam reduce but do not eliminate the mismatch (because they only order
within a sub-block, not across).
Diagnostic
To isolate the alias-tracking issue from the orthogonal "pipe ops
under sub-block parallelism" question (separate issue), I ran the
following control:
(
row_off_sb = 0for both).(≈ 50% of elements baseline-wrong).
Result: 47.8% mismatch — i.e. less than the 50% baseline, meaning
the upper 16 rows are computed correctly. So:
pto.subsetop itself are correct.halves of the same parent (the realistic case), mismatch jumps to
~64% — the cross-sub-block writes are not getting the WAR/RAW
protection auto-sync would normally insert for the parent.
Expected behavior
Two ops
A,Baccessing memory through any of:t,subset(t, ...)or any chain of subsets rooted att,should be treated by the auto-sync dependency tracker as touching the
same memory. Standard 'set_flag / wait_flag' analysis on the parent tile. If the
subsets can be proven non-overlapping at compile time (constant
offsets + sizes), the pass may skip emitting a flag for that pair as
an optimisation.
Actual behavior / error logs
Result: 47.8% mismatch against torch npu reference — i.e. less than the 50% baseline, meaning the upper 16 rows are computed correctly. So:
pto.subsetop itself are correct.halves of the same parent (the realistic case), mismatch jumps to
~64% — the cross-sub-block writes are not getting the WAR/RAW
protection auto-sync would normally insert for the parent.
Git commit
07d7fc3
Host platform
Linux (aarch64)
Target Ascend arch (if relevant)
a3
PTOAS build level (if relevant)
None