@triton.jit
def dropconnect_fwd_kernel(x_ptr, w_ptr, y_ptr, seed, M, K, N, stride_xm,
    stride_xk, stride_wk, stride_wn, stride_ym, stride_yn, stride_dm,
    stride_dk, stride_dn, BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.
    constexpr, BLOCK_SIZE_K: tl.constexpr, ALLOWTF32: tl.constexpr):
    pid_m = tl.program_id(0)
    pid_n = tl.program_id(1)
    offset_m = pid_m * BLOCK_SIZE_M
    offset_n = pid_n * BLOCK_SIZE_N
    offset_k = 0
    x_offsets = block_offsets_2d(M, K, stride_xm, stride_xk, offset_m,
        offset_k, BLOCK_SIZE_M, BLOCK_SIZE_K)
    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)
    x_offsets = x_offsets.reshape(BLOCK_SIZE_M, BLOCK_SIZE_K, 1)
    w_offsets = w_offsets.reshape(1, BLOCK_SIZE_K, BLOCK_SIZE_N)
    offs_k = tl.arange(0, BLOCK_SIZE_K)
    x_tile = x_ptr + x_offsets
    w_tile = w_ptr + w_offsets
    ASM: tl.constexpr = 'cvt.rna.tf32.f32 $0, $1;'
    accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
    for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):
        random_masks = tl.random.rand(seed, d_offsets) > 0.5
        k_mask = offs_k[None, :, None] < K - k * BLOCK_SIZE_K
        x_load = tl.load(x_tile, mask=k_mask, other=0.0)
        w_load = tl.load(w_tile, mask=k_mask, other=0.0)
        a = tl.where(random_masks, x_load, 0.0)
        b = tl.where(random_masks, w_load, 0.0)
        mul = a * b
        accumulator += tl.sum(mul, axis=1)
        x_tile += BLOCK_SIZE_K * stride_xk
        w_tile += BLOCK_SIZE_K * stride_wk
        d_offsets += BLOCK_SIZE_K * stride_dk
    y_offset, y_mask = block_offsets_2d(M, N, stride_ym, stride_yn,
        offset_m, offset_n, BLOCK_SIZE_M, BLOCK_SIZE_N, True)
    y_tile = y_ptr + y_offset
    y = accumulator.to(y_tile.dtype.element_ty)
    tl.store(y_tile, y, mask=y_mask)
