@triton.jit
def quantize_kernel(X, Y, N, K, K1, scale, zero_point, BITS: tl.constexpr,
    QFACT: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.
    constexpr, BLOCK_SIZE_K1: tl.constexpr):
    dtype = X.dtype.element_ty
    scale_ty = tl.full((), scale, dtype=dtype)
    zero_point_ty = tl.full((), zero_point, dtype=dtype)
    QMASK = (1 << BITS) - 1
    QMASK_FLT = tl.full((), QMASK, dtype=dtype)
    pid_n = tl.program_id(0)
    pid_k = tl.program_id(1)
    bk = tl.arange(0, BLOCK_SIZE_K)
    bk1 = tl.arange(0, BLOCK_SIZE_K1)
    x_shift = BITS * (bk % QFACT)
    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] * K + rk[None, :])
    Y1 = Y + (rn[:, None] * K1 + rk1[None, :])
    x = tl.load(X1, mask=mask_x)
    xf = clamp(x / scale_ty + zero_point_ty, 0.0, QMASK_FLT)
    xi = tl.math.rint(xf).to(tl.uint8)
    xq = xi << x_shift
    mat = tl.reshape(xq, (BLOCK_SIZE_N, BLOCK_SIZE_K1, QFACT))
    out = tl.sum(mat, axis=2)
    tl.store(Y1, out, mask=mask_y)
