@triton.jit
def dequantize_kernel(X, Y, N, K, K1, scale, zero_point, BITS: tl.constexpr,
    BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr, BLOCK_SIZE_K1:
    tl.constexpr):
    dtype = Y.dtype.element_ty
    scale_ty = tl.full((), scale, dtype=dtype)
    zero_point_ty = tl.full((), zero_point, dtype=dtype)
    QFACT = 8 // BITS
    QMASK = (1 << BITS) - 1
    QMASK_INT = tl.full((), QMASK, dtype=tl.uint8)
    pid_n = tl.program_id(0)
    pid_k = tl.program_id(1)
    bk = tl.arange(0, BLOCK_SIZE_K)
    bk1, bq1 = bk // QFACT, bk % QFACT
    x_shift = BITS * bq1
    rn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
    rk = pid_k * BLOCK_SIZE_K + bk
    rk1 = pid_k * BLOCK_SIZE_K1 + bk1
    mask_x = rn[:, None] < N
    mask_y = rn[:, None] < N
    X1 = X + (rn[:, None] * K1 + rk1[None, :])
    Y1 = Y + (rn[:, None] * K + rk[None, :])
    x = tl.load(X1, mask=mask_x)
    xi = x >> x_shift & QMASK_INT
    xf = scale_ty * (xi.to(dtype) - zero_point_ty)
    tl.store(Y1, xf, mask=mask_y)
