@triton.jit
def _layer_norm_bwd_dx_fused(DX, DY, DW, DB, Y, W, B, V, Lock, stride, N,
    **META):
    GROUP_SIZE_M = META['GROUP_SIZE_M']
    BLOCK_SIZE_N = META['BLOCK_SIZE_N']
    row = tl.program_id(0)
    cols = tl.arange(0, BLOCK_SIZE_N)
    y_ptrs = Y + row * stride + cols
    dy_ptrs = DY + row * stride + cols
    w_ptrs = W + cols
    b_ptrs = B + cols
    lock_id = row % GROUP_SIZE_M
    Lock += lock_id
    Count = Lock + GROUP_SIZE_M
    y = tl.load(y_ptrs, mask=cols < N, other=0).to(tl.float32)
    dy = tl.load(dy_ptrs, mask=cols < N, other=0).to(tl.float32)
    w = tl.load(w_ptrs, mask=cols < N, other=0).to(tl.float32)
    b = tl.load(b_ptrs, mask=cols < N, other=0).to(tl.float32)
    rstd = tl.load(V + row)
    xhat = (y - b) / w
    wdy = w * dy
    xhat = tl.where(cols < N, xhat, 0.0)
    wdy = tl.where(cols < N, wdy, 0.0)
    mean1 = tl.sum(xhat * wdy, axis=0) / N
    mean2 = tl.sum(wdy, axis=0) / N
    dx = (wdy - (xhat * mean1 + mean2)) * rstd
    _store(dx, DX, stride, N, META)
    partial_dw = (dy * xhat).to(w.dtype)
    partial_db = dy.to(w.dtype)
    while tl.atomic_cas(Lock, 0, 1) == 1:
        pass
    count = tl.load(Count)
    dw_ptrs = DW + lock_id * N + cols
    db_ptrs = DB + lock_id * N + cols
    if count == 0:
        tl.atomic_xchg(Count, 1)
    else:
        partial_dw += tl.load(dw_ptrs, mask=cols < N, other=0.0)
        partial_db += tl.load(db_ptrs, mask=cols < N, other=0.0)
    tl.store(dw_ptrs, partial_dw, mask=cols < N)
    tl.store(db_ptrs, partial_db, mask=cols < N)
    tl.atomic_xchg(Lock, 0)
