@triton.jit
def _kernel_dequantize_mx4(A, mx4_lookup_table, out, M, GROUPS_PER_THREAD,
    GROUP_SIZE: tl.constexpr, GROUP_LOAD: tl.constexpr, USE_INT64: tl.constexpr
    ) ->None:
    """Dequantize a packed MX4 tensor and apply scaling.

    Args:
        A (Tensor): [M] MX4 tensor packed into int8.
        shared_exp (Tensor): Int8 tensor representing group exponent.
        mx4_lookup_table (Tensor): Map from mx4 integer value to floating point.
        M (int): Total number of elements in input.
        GROUPS_PER_THREAD (int): Number of groups each thread is responsible for.
        GROUP_SIZE (int): Size of chunks that use the same shared exponent.
        GROUP_LOAD (int): Number of groups to process simultaneously.
        USE_INT64 (bool): Whether to use int64 for indexing.
    """
    MX4_BIT_MASK: tl.constexpr = 15
    FP32_EXP_BIAS: tl.constexpr = 127
    PACKED_GROUP_SIZE: tl.constexpr = GROUP_SIZE // 2 + 1
    pid = tl.program_id(0)
    if USE_INT64:
        pid = pid.to(tl.int64)
        M = tl.cast(M, tl.int64)
        GROUPS_PER_THREAD = tl.cast(GROUPS_PER_THREAD, tl.int64)
    INPUT_CHUNK_SIZE = GROUPS_PER_THREAD * PACKED_GROUP_SIZE
    OUTPUT_CHUNK_SIZE = GROUPS_PER_THREAD * GROUP_SIZE
    OUTPUT_SIZE = M // PACKED_GROUP_SIZE * GROUP_SIZE
    input_start = pid * (GROUPS_PER_THREAD * PACKED_GROUP_SIZE)
    exp_start = input_start + GROUP_SIZE // 2
    output_start = pid * OUTPUT_CHUNK_SIZE
    input_offset = tl.arange(0, GROUP_LOAD * GROUP_SIZE // 2)
    exp_indices = input_offset // (GROUP_SIZE // 2)
    input_offset = input_offset + exp_indices + input_start
    output_offset = tl.arange(0, GROUP_LOAD * GROUP_SIZE) + output_start
    exp_offset = exp_indices * PACKED_GROUP_SIZE + exp_start
    for _k in range(0, tl.cdiv(GROUPS_PER_THREAD, GROUP_LOAD)):
        a = tl.load(A + input_offset, mask=(input_offset < M) & (
            input_offset < INPUT_CHUNK_SIZE * (pid + 1)), other=0.0)
        low_mx4 = a & MX4_BIT_MASK
        high_mx4 = a >> 4 & MX4_BIT_MASK
        low_fp32 = tl.load(mx4_lookup_table + low_mx4)
        high_fp32 = tl.load(mx4_lookup_table + high_mx4)
        exp = tl.load(A + exp_offset, mask=(exp_offset < M) & (exp_offset <
            INPUT_CHUNK_SIZE * (pid + 1)), other=0.0)
        exp = exp.to(tl.int16) - FP32_EXP_BIAS
        scale = tl.exp2(exp.to(tl.float64)).to(tl.float32)
        scaled_low_fp32 = scale * low_fp32
        scaled_high_fp32 = scale * high_fp32
        scaled_fp32 = tl.interleave(scaled_low_fp32, scaled_high_fp32)
        tl.store(out + output_offset, scaled_fp32, mask=(output_offset <
            OUTPUT_SIZE) & (output_offset < OUTPUT_CHUNK_SIZE * (pid + 1)))
        input_offset += GROUP_LOAD * PACKED_GROUP_SIZE
        exp_offset += GROUP_LOAD * PACKED_GROUP_SIZE
        output_offset += GROUP_LOAD * GROUP_SIZE
