@triton.jit
def dropconnect_dx_kernel(dy_ptr, x_ptr, dw_ptr, seed, M, K, N, stride_dym,
    stride_dyn, stride_xm, stride_xk, stride_dm, stride_dk, stride_dn,
    stride_dwk, stride_dwn, 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)
    dWD_m = dO_m * x_m_cast
    dw_m_cast = dWD_m * D
    dw_m = dw_m_cast.sum(dim=0) """
    pid_k = tl.program_id(0)
    pid_n = tl.program_id(1)
    offset_m = 0
    offset_k = pid_k * BLOCK_SIZE_K
    offset_n = pid_n * BLOCK_SIZE_N
    dy_offsets = block_offsets_2d(M, N, stride_dym, stride_dyn, offset_m,
        offset_n, BLOCK_SIZE_M, BLOCK_SIZE_N)
    x_offsets = block_offsets_2d(M, K, stride_xm, stride_xk, offset_m,
        offset_k, BLOCK_SIZE_M, BLOCK_SIZE_K)
    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)
    x_offsets = x_offsets.reshape(BLOCK_SIZE_M, BLOCK_SIZE_K, 1)
    offs_m = tl.arange(0, BLOCK_SIZE_M)
    dy_tile = dy_ptr + dy_offsets
    x_tile = x_ptr + x_offsets
    ASM: tl.constexpr = 'cvt.rna.tf32.f32 $0, $1;'
    accumulator = tl.zeros((BLOCK_SIZE_K, BLOCK_SIZE_N), dtype=tl.float32)
    for m in range(0, tl.cdiv(M, BLOCK_SIZE_M)):
        random_masks = tl.random.rand(seed, d_offsets) > 0.5
        m_mask = offs_m[:, None, None] < M - m * BLOCK_SIZE_M
        dy_load = tl.load(dy_tile, mask=m_mask, other=0.0)
        x_load = tl.load(x_tile, mask=m_mask, other=0.0)
        dy = tl.where(random_masks, dy_load, 0.0)
        wd = tl.where(random_masks, x_load, 0.0)
        mul = dy * wd
        accumulator += tl.sum(mul, axis=0)
        dy_tile += BLOCK_SIZE_M * stride_dym
        x_tile += BLOCK_SIZE_M * stride_xm
        d_offsets += BLOCK_SIZE_M * stride_dm
    dw_offset, dw_mask = block_offsets_2d(K, N, stride_dwk, stride_dwn,
        offset_k, offset_n, BLOCK_SIZE_K, BLOCK_SIZE_N, True)
    dw_tile = dw_ptr + dw_offset
    dw = accumulator.to(dw_tile.dtype.element_ty)
    tl.store(dw_tile, dw, mask=dw_mask)
