@triton.jit
def dropconnect_dx_kernel(dy_ptr, w_ptr, dx_ptr, seed, M, K, N, stride_dym,
    stride_dyn, stride_wk, stride_wn, stride_dm, stride_dk, stride_dn,
    stride_xm, stride_xk, BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.
    constexpr, BLOCK_SIZE_K: tl.constexpr, ALLOWTF32: tl.constexpr):
    """ 
    dY_m = Y.grad
    dO_m = dY_m.view(M,1,N).broadcast_to(M,K,N)
    dx_m_cast = dO_m * WD_m
    dx_m = dx_m_cast.sum(dim=2) """
    pid_m = tl.program_id(0)
    pid_k = tl.program_id(1)
    offset_m = pid_m * BLOCK_SIZE_M
    offset_n = 0
    offset_k = pid_k * BLOCK_SIZE_K
    dy_offsets = block_offsets_2d(M, N, stride_dym, stride_dyn, offset_m,
        offset_n, BLOCK_SIZE_M, BLOCK_SIZE_N)
    w_offsets = block_offsets_2d(K, N, stride_wk, stride_wn, offset_k,
        offset_n, BLOCK_SIZE_K, BLOCK_SIZE_N)
    d_offsets = block_offsets_3d(M, K, N, stride_dm, stride_dk, stride_dn,
        offset_m, offset_k, offset_n, BLOCK_SIZE_M, BLOCK_SIZE_K, BLOCK_SIZE_N)
    dy_offsets = dy_offsets.reshape(BLOCK_SIZE_M, 1, BLOCK_SIZE_N)
    w_offsets = w_offsets.reshape(1, BLOCK_SIZE_K, BLOCK_SIZE_N)
    offs_n = tl.arange(0, BLOCK_SIZE_N)
    dy_tile = dy_ptr + dy_offsets
    w_tile = w_ptr + w_offsets
    ASM: tl.constexpr = 'cvt.rna.tf32.f32 $0, $1;'
    accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_K), dtype=tl.float32)
    for n in range(0, tl.cdiv(N, BLOCK_SIZE_N)):
        random_masks = tl.random.rand(seed, d_offsets) > 0.5
        n_mask = offs_n[None, None, :] < N - n * BLOCK_SIZE_N
        dy_load = tl.load(dy_tile, mask=n_mask, other=0.0)
        w_load = tl.load(w_tile, mask=n_mask, other=0.0)
        dy = tl.where(random_masks, dy_load, 0.0)
        wd = tl.where(random_masks, w_load, 0.0)
        mul = dy * wd
        accumulator += tl.sum(mul, axis=2)
        dy_tile += BLOCK_SIZE_N * stride_dyn
        w_tile += BLOCK_SIZE_N * stride_wn
        d_offsets += BLOCK_SIZE_N * stride_dn
    dx_offset, dx_mask = block_offsets_2d(M, K, stride_xm, stride_xk,
        offset_m, offset_k, BLOCK_SIZE_M, BLOCK_SIZE_K, True)
    dx_tile = dx_ptr + dx_offset
    dx = accumulator.to(dx_tile.dtype.element_ty)
    tl.store(dx_tile, dx, mask=dx_mask)
