`ext.sync_block_all` causes device deadlock on Ascend 910B2

[h1074112368]

Describe the bug

ext.sync_block_all causes device deadlock on Ascend 910B2

Describe the bug

ext.sync_block_all("all", event_id) causes a device deadlock (30s timeout) when launched with more than 1 program on Ascend 910B2. The kernel compiles successfully but hangs at the sync point during execution, resulting in npuSynchronizeDevice timeout error.

The C++ builder binding exists (ascendnpu_ir_builder.sync_block_all is True), and the Python API (triton.language.extra.cann.extension.sync_block_all) is importable, but the generated hardware barrier instruction does not release.

Minimal Reproducible Example

import torch
import torch_npu
torch_npu.npu.set_device(0)
import triton
import triton.language as tl
from triton.language.extra.cann import extension as ext

N = 4800  # 48 cores × 100 elements each

@triton.jit
def k_sync_sum(data_ptr, ws_ptr, out_ptr, N: tl.constexpr, RB: tl.constexpr):
    pid = tl.program_id(0)
    ncores = 48

    # Each core computes partial sum over its [lo, hi) slice
    chunk = (N + ncores - 1) // ncores
    lo = pid * chunk
    hi = tl.minimum(lo + chunk, N)
    acc = tl.zeros((RB,), tl.float32)
    for i in range(lo, hi, RB):
        offs = i + tl.arange(0, RB)
        mask = offs < hi
        acc += tl.load(data_ptr + offs, mask=mask, other=0.0)
    partial = tl.sum(acc, 0)

    # Write partial to workspace slot
    tl.store(ws_ptr + pid, partial)

    # Cross-core synchronization — DEADLOCKS HERE
    ext.sync_block_all("all", 0)

    # Core 0 reads all 48 partials and produces final sum
    if pid == 0:
        total = tl.zeros((64,), tl.float32)
        for j in range(0, ncores, 64):
            idx = j + tl.arange(0, 64)
            m = idx < ncores
            total += tl.load(ws_ptr + idx, mask=m, other=0.0)
        tl.store(out_ptr, tl.sum(total, 0))


data = torch.randn(N, device="npu")
ws = torch.zeros(48, device="npu")
out = torch.zeros(1, device="npu")
ref = data.sum()

k_sync_sum[(48,)](data, ws, out, N=N, RB=64)
torch.npu.synchronize()
print(f"got={out.item():.4f} ref={ref.item():.4f}")

Observed Result

RuntimeError: npuSynchronizeDevice:torch_npu/csrc/core/npu/NPUStream.cpp:748
NPU function error: AclrtSynchronizeDeviceWithTimeout, error code: 507033

The kernel compiles without errors. At runtime, all 48 programs reach the sync_block_all call and deadlock — the device watchdog kills them after the 30-second timeout.

With grid=(1,) (single core), the same kernel also hangs — even a single-core barrier against itself does not release.

Expected Result

got=15.7032 ref=15.7032

All 48 cores should synchronize at the barrier, then core 0 reads the fully-populated workspace and produces the correct total sum.

Environment details

Component	Version
Hardware	Ascend 910B2 (aarch64)
CANN	9.1.0-beta.1
triton-ascend	3.2.2+git0d69f740 (source build)
triton	3.2.0
torch_npu	2.13.0+gita4623d3
PyTorch	2.13.0+cpu
bishengir-compile	1.2.0 (llvm 19.1.7)

[hongziqi]

Thanks for your feedback. We will arrange relevant engineers to investigate this issue later.