@triton.autotune(configs=get_configs(), key=['m', 'k'])
@triton.jit
def _sparse_semi_structured_from_dense_triton_8(dense_ptr, sparse_ptr,
    meta_reordered_ptr, mask_ptr, dense_row_stride, sparse_row_stride,
    mask_row_stride, dense_col_stride, sparse_col_stride, mask_col_stride,
    m, k, seed, BLOCK_SIZE: tl.constexpr, PRUNE: tl.constexpr, ARRAY_LAYOUT:
    tl.constexpr):
    if ARRAY_LAYOUT == 'row':
        row_idx = tl.program_id(0)
        col_idx = tl.program_id(1) * 32 * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE
            ) * 32
        mask = col_idx < k
    elif ARRAY_LAYOUT == 'col':
        row_idx = tl.arange(0, BLOCK_SIZE) + tl.program_id(0) * BLOCK_SIZE
        col_idx = tl.program_id(1) * 32
        mask = row_idx < m
    dense_40 = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        0) * dense_col_stride, mask=mask)
    dense_41 = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        1) * dense_col_stride, mask=mask)
    dense_42 = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        2) * dense_col_stride, mask=mask)
    dense_43 = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        3) * dense_col_stride, mask=mask)
    dense_44 = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        4) * dense_col_stride, mask=mask)
    dense_45 = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        5) * dense_col_stride, mask=mask)
    dense_46 = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        6) * dense_col_stride, mask=mask)
    dense_47 = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        7) * dense_col_stride, mask=mask)
    dense_48 = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        8) * dense_col_stride, mask=mask)
    dense_49 = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        9) * dense_col_stride, mask=mask)
    dense_4A = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        10) * dense_col_stride, mask=mask)
    dense_4B = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        11) * dense_col_stride, mask=mask)
    dense_4C = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        12) * dense_col_stride, mask=mask)
    dense_4D = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        13) * dense_col_stride, mask=mask)
    dense_4E = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        14) * dense_col_stride, mask=mask)
    dense_4F = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        15) * dense_col_stride, mask=mask)
    dense_4G = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        16) * dense_col_stride, mask=mask)
    dense_4H = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        17) * dense_col_stride, mask=mask)
    dense_4I = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        18) * dense_col_stride, mask=mask)
    dense_4J = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        19) * dense_col_stride, mask=mask)
    dense_4K = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        20) * dense_col_stride, mask=mask)
    dense_4L = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        21) * dense_col_stride, mask=mask)
    dense_4M = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        22) * dense_col_stride, mask=mask)
    dense_4N = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        23) * dense_col_stride, mask=mask)
    dense_4O = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        24) * dense_col_stride, mask=mask)
    dense_4P = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        25) * dense_col_stride, mask=mask)
    dense_4Q = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        26) * dense_col_stride, mask=mask)
    dense_4R = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        27) * dense_col_stride, mask=mask)
    dense_4S = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        28) * dense_col_stride, mask=mask)
    dense_4T = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        29) * dense_col_stride, mask=mask)
    dense_4U = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        30) * dense_col_stride, mask=mask)
    dense_4V = tl.load(dense_ptr + row_idx * dense_row_stride + (col_idx + 
        31) * dense_col_stride, mask=mask)
    if PRUNE == 'mse':
        (_dense_40, _dense_41, _dense_42, _dense_43, _dense_44, _dense_45,
            _dense_46, _dense_47, _dense_48, _dense_49, _dense_4A,
            _dense_4B, _dense_4C, _dense_4D, _dense_4E, _dense_4F) = (
            dense_40, dense_41, dense_42, dense_43, dense_44, dense_45,
            dense_46, dense_47, dense_48, dense_49, dense_4A, dense_4B,
            dense_4C, dense_4D, dense_4E, dense_4F)
        (_dense_4G, _dense_4H, _dense_4I, _dense_4J, _dense_4K, _dense_4L,
            _dense_4M, _dense_4N, _dense_4O, _dense_4P, _dense_4Q,
            _dense_4R, _dense_4S, _dense_4T, _dense_4U, _dense_4V) = (
            dense_4G, dense_4H, dense_4I, dense_4J, dense_4K, dense_4L,
            dense_4M, dense_4N, dense_4O, dense_4P, dense_4Q, dense_4R,
            dense_4S, dense_4T, dense_4U, dense_4V)
        x1, x2, x3, x4, x5, x6 = tl.abs(_dense_40) > tl.abs(_dense_41), tl.abs(
            _dense_40) > tl.abs(_dense_42), tl.abs(_dense_40) > tl.abs(
            _dense_43), tl.abs(_dense_41) > tl.abs(_dense_42), tl.abs(_dense_41
            ) > tl.abs(_dense_43), tl.abs(_dense_42) > tl.abs(_dense_43)
        m0, m1, m2, m3 = (x2 & x3 | x1 & x2 | x1 & x3, ~x1 & x5 | x4 & x5 |
            ~x1 & x4, ~x2 & ~x4 | ~x2 & x6 | ~x4 & x6, ~x3 & ~x5 | ~x3 & ~
            x6 | ~x5 & ~x6)
        x1, x2, x3, x4, x5, x6 = tl.abs(_dense_44) > tl.abs(_dense_45), tl.abs(
            _dense_44) > tl.abs(_dense_46), tl.abs(_dense_44) > tl.abs(
            _dense_47), tl.abs(_dense_45) > tl.abs(_dense_46), tl.abs(_dense_45
            ) > tl.abs(_dense_47), tl.abs(_dense_46) > tl.abs(_dense_47)
        m4, m5, m6, m7 = (x2 & x3 | x1 & x2 | x1 & x3, ~x1 & x5 | x4 & x5 |
            ~x1 & x4, ~x2 & ~x4 | ~x2 & x6 | ~x4 & x6, ~x3 & ~x5 | ~x3 & ~
            x6 | ~x5 & ~x6)
        x1, x2, x3, x4, x5, x6 = tl.abs(_dense_48) > tl.abs(_dense_49), tl.abs(
            _dense_48) > tl.abs(_dense_4A), tl.abs(_dense_48) > tl.abs(
            _dense_4B), tl.abs(_dense_49) > tl.abs(_dense_4A), tl.abs(_dense_49
            ) > tl.abs(_dense_4B), tl.abs(_dense_4A) > tl.abs(_dense_4B)
        m8, m9, mA, mB = (x2 & x3 | x1 & x2 | x1 & x3, ~x1 & x5 | x4 & x5 |
            ~x1 & x4, ~x2 & ~x4 | ~x2 & x6 | ~x4 & x6, ~x3 & ~x5 | ~x3 & ~
            x6 | ~x5 & ~x6)
        x1, x2, x3, x4, x5, x6 = tl.abs(_dense_4C) > tl.abs(_dense_4D), tl.abs(
            _dense_4C) > tl.abs(_dense_4E), tl.abs(_dense_4C) > tl.abs(
            _dense_4F), tl.abs(_dense_4D) > tl.abs(_dense_4E), tl.abs(_dense_4D
            ) > tl.abs(_dense_4F), tl.abs(_dense_4E) > tl.abs(_dense_4F)
        mC, mD, mE, mF = (x2 & x3 | x1 & x2 | x1 & x3, ~x1 & x5 | x4 & x5 |
            ~x1 & x4, ~x2 & ~x4 | ~x2 & x6 | ~x4 & x6, ~x3 & ~x5 | ~x3 & ~
            x6 | ~x5 & ~x6)
        x1, x2, x3, x4, x5, x6 = tl.abs(_dense_4G) > tl.abs(_dense_4H), tl.abs(
            _dense_4G) > tl.abs(_dense_4I), tl.abs(_dense_4G) > tl.abs(
            _dense_4J), tl.abs(_dense_4H) > tl.abs(_dense_4I), tl.abs(_dense_4H
            ) > tl.abs(_dense_4J), tl.abs(_dense_4I) > tl.abs(_dense_4J)
        mG, mH, mI, mJ = (x2 & x3 | x1 & x2 | x1 & x3, ~x1 & x5 | x4 & x5 |
            ~x1 & x4, ~x2 & ~x4 | ~x2 & x6 | ~x4 & x6, ~x3 & ~x5 | ~x3 & ~
            x6 | ~x5 & ~x6)
        x1, x2, x3, x4, x5, x6 = tl.abs(_dense_4K) > tl.abs(_dense_4L), tl.abs(
            _dense_4K) > tl.abs(_dense_4M), tl.abs(_dense_4K) > tl.abs(
            _dense_4N), tl.abs(_dense_4L) > tl.abs(_dense_4M), tl.abs(_dense_4L
            ) > tl.abs(_dense_4N), tl.abs(_dense_4M) > tl.abs(_dense_4N)
        mK, mL, mM, mN = (x2 & x3 | x1 & x2 | x1 & x3, ~x1 & x5 | x4 & x5 |
            ~x1 & x4, ~x2 & ~x4 | ~x2 & x6 | ~x4 & x6, ~x3 & ~x5 | ~x3 & ~
            x6 | ~x5 & ~x6)
        x1, x2, x3, x4, x5, x6 = tl.abs(_dense_4O) > tl.abs(_dense_4P), tl.abs(
            _dense_4O) > tl.abs(_dense_4Q), tl.abs(_dense_4O) > tl.abs(
            _dense_4R), tl.abs(_dense_4P) > tl.abs(_dense_4Q), tl.abs(_dense_4P
            ) > tl.abs(_dense_4R), tl.abs(_dense_4Q) > tl.abs(_dense_4R)
        mO, mP, mQ, mR = (x2 & x3 | x1 & x2 | x1 & x3, ~x1 & x5 | x4 & x5 |
            ~x1 & x4, ~x2 & ~x4 | ~x2 & x6 | ~x4 & x6, ~x3 & ~x5 | ~x3 & ~
            x6 | ~x5 & ~x6)
        x1, x2, x3, x4, x5, x6 = tl.abs(_dense_4S) > tl.abs(_dense_4T), tl.abs(
            _dense_4S) > tl.abs(_dense_4U), tl.abs(_dense_4S) > tl.abs(
            _dense_4V), tl.abs(_dense_4T) > tl.abs(_dense_4U), tl.abs(_dense_4T
            ) > tl.abs(_dense_4V), tl.abs(_dense_4U) > tl.abs(_dense_4V)
        mS, mT, mU, mV = (x2 & x3 | x1 & x2 | x1 & x3, ~x1 & x5 | x4 & x5 |
            ~x1 & x4, ~x2 & ~x4 | ~x2 & x6 | ~x4 & x6, ~x3 & ~x5 | ~x3 & ~
            x6 | ~x5 & ~x6)
    elif PRUNE == 'mask':
        m0 = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 0) *
            mask_col_stride, mask=mask).to(tl.int1)
        m1 = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 1) *
            mask_col_stride, mask=mask).to(tl.int1)
        m2 = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 2) *
            mask_col_stride, mask=mask).to(tl.int1)
        m3 = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 3) *
            mask_col_stride, mask=mask).to(tl.int1)
        m4 = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 4) *
            mask_col_stride, mask=mask).to(tl.int1)
        m5 = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 5) *
            mask_col_stride, mask=mask).to(tl.int1)
        m6 = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 6) *
            mask_col_stride, mask=mask).to(tl.int1)
        m7 = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 7) *
            mask_col_stride, mask=mask).to(tl.int1)
        m8 = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 8) *
            mask_col_stride, mask=mask).to(tl.int1)
        m9 = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 9) *
            mask_col_stride, mask=mask).to(tl.int1)
        mA = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 10) *
            mask_col_stride, mask=mask).to(tl.int1)
        mB = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 11) *
            mask_col_stride, mask=mask).to(tl.int1)
        mC = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 12) *
            mask_col_stride, mask=mask).to(tl.int1)
        mD = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 13) *
            mask_col_stride, mask=mask).to(tl.int1)
        mE = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 14) *
            mask_col_stride, mask=mask).to(tl.int1)
        mF = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 15) *
            mask_col_stride, mask=mask).to(tl.int1)
        mG = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 16) *
            mask_col_stride, mask=mask).to(tl.int1)
        mH = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 17) *
            mask_col_stride, mask=mask).to(tl.int1)
        mI = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 18) *
            mask_col_stride, mask=mask).to(tl.int1)
        mJ = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 19) *
            mask_col_stride, mask=mask).to(tl.int1)
        mK = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 20) *
            mask_col_stride, mask=mask).to(tl.int1)
        mL = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 21) *
            mask_col_stride, mask=mask).to(tl.int1)
        mM = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 22) *
            mask_col_stride, mask=mask).to(tl.int1)
        mN = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 23) *
            mask_col_stride, mask=mask).to(tl.int1)
        mO = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 24) *
            mask_col_stride, mask=mask).to(tl.int1)
        mP = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 25) *
            mask_col_stride, mask=mask).to(tl.int1)
        mQ = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 26) *
            mask_col_stride, mask=mask).to(tl.int1)
        mR = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 27) *
            mask_col_stride, mask=mask).to(tl.int1)
        mS = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 28) *
            mask_col_stride, mask=mask).to(tl.int1)
        mT = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 29) *
            mask_col_stride, mask=mask).to(tl.int1)
        mU = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 30) *
            mask_col_stride, mask=mask).to(tl.int1)
        mV = tl.load(mask_ptr + row_idx * mask_row_stride + (col_idx + 31) *
            mask_col_stride, mask=mask).to(tl.int1)
    elif PRUNE == 'mvue':
        if ARRAY_LAYOUT == 'row':
            seed0 = seed + (tl.program_id(0) + tl.program_id(1) * m) * 2
            seed1 = seed + (tl.program_id(0) + tl.program_id(1) * m) * 2 + 1
        else:
            seed0 = seed + (tl.program_id(0) * k // 32 + tl.program_id(1)) * 2
            seed1 = seed + (tl.program_id(0) * k // 32 + tl.program_id(1)
                ) * 2 + 1
        random0, random1, random2, random3 = tl.rand4x(seed0, tl.arange(0,
            BLOCK_SIZE), n_rounds=5)
        random4, random5, random6, random7 = tl.rand4x(seed1, tl.arange(0,
            BLOCK_SIZE), n_rounds=5)
        dense_40, dense_41, dense_42, dense_43, m0, m1, m2, m3 = (
            _MVUE24_approx(dense_40, dense_41, dense_42, dense_43, random0,
            random1))
        dense_44, dense_45, dense_46, dense_47, m4, m5, m6, m7 = (
            _MVUE24_approx(dense_44, dense_45, dense_46, dense_47, random2,
            random3))
        dense_48, dense_49, dense_4A, dense_4B, m8, m9, mA, mB = (
            _MVUE24_approx(dense_48, dense_49, dense_4A, dense_4B, random4,
            random5))
        dense_4C, dense_4D, dense_4E, dense_4F, mC, mD, mE, mF = (
            _MVUE24_approx(dense_4C, dense_4D, dense_4E, dense_4F, random6,
            random7))
    else:
        m0 = dense_40 != 0
        m1 = dense_41 != 0
        m2 = dense_42 != 0
        m3 = dense_43 != 0
        m4 = dense_44 != 0
        m5 = dense_45 != 0
        m6 = dense_46 != 0
        m7 = dense_47 != 0
        m8 = dense_48 != 0
        m9 = dense_49 != 0
        mA = dense_4A != 0
        mB = dense_4B != 0
        mC = dense_4C != 0
        mD = dense_4D != 0
        mE = dense_4E != 0
        mF = dense_4F != 0
        mG = dense_4G != 0
        mH = dense_4H != 0
        mI = dense_4I != 0
        mJ = dense_4J != 0
        mK = dense_4K != 0
        mL = dense_4L != 0
        mM = dense_4M != 0
        mN = dense_4N != 0
        mO = dense_4O != 0
        mP = dense_4P != 0
        mQ = dense_4Q != 0
        mR = dense_4R != 0
        mS = dense_4S != 0
        mT = dense_4T != 0
        mU = dense_4U != 0
        mV = dense_4V != 0
    bit0 = ~m0 & m1
    bit1 = ~m0 & ~m1
    bit2 = bit1 | ~m2
    bit3 = bit0 | ~m1 | m2
    idxs0 = bit0 | bit1.to(tl.int64) << 1
    idxs1 = bit2 | bit3.to(tl.int64) << 1
    sparse0 = tl.where(bit1, tl.where(bit0, dense_43, dense_42), tl.where(
        bit0, dense_41, dense_40))
    sparse1 = tl.where(bit3, tl.where(bit2, dense_43, dense_42), tl.where(
        bit2, dense_41, dense_40))
    bit4 = ~m4 & m5
    bit5 = ~m4 & ~m5
    bit6 = bit5 | ~m6
    bit7 = bit4 | ~m5 | m6
    idxs2 = bit4 | bit5.to(tl.int64) << 1
    idxs3 = bit6 | bit7.to(tl.int64) << 1
    sparse2 = tl.where(bit5, tl.where(bit4, dense_47, dense_46), tl.where(
        bit4, dense_45, dense_44))
    sparse3 = tl.where(bit7, tl.where(bit6, dense_47, dense_46), tl.where(
        bit6, dense_45, dense_44))
    bit8 = ~m8 & m9
    bit9 = ~m8 & ~m9
    bitA = bit9 | ~mA
    bitB = bit8 | ~m9 | mA
    idxs4 = bit8 | bit9.to(tl.int64) << 1
    idxs5 = bitA | bitB.to(tl.int64) << 1
    sparse4 = tl.where(bit9, tl.where(bit8, dense_4B, dense_4A), tl.where(
        bit8, dense_49, dense_48))
    sparse5 = tl.where(bitB, tl.where(bitA, dense_4B, dense_4A), tl.where(
        bitA, dense_49, dense_48))
    bitC = ~mC & mD
    bitD = ~mC & ~mD
    bitE = bitD | ~mE
    bitF = bitC | ~mD | mE
    idxs6 = bitC | bitD.to(tl.int64) << 1
    idxs7 = bitE | bitF.to(tl.int64) << 1
    sparse6 = tl.where(bitD, tl.where(bitC, dense_4F, dense_4E), tl.where(
        bitC, dense_4D, dense_4C))
    sparse7 = tl.where(bitF, tl.where(bitE, dense_4F, dense_4E), tl.where(
        bitE, dense_4D, dense_4C))
    bitG = ~mG & mH
    bitH = ~mG & ~mH
    bitI = bitH | ~mI
    bitJ = bitG | ~mH | mI
    idxs8 = bitG | bitH.to(tl.int64) << 1
    idxs9 = bitI | bitJ.to(tl.int64) << 1
    sparse8 = tl.where(bitH, tl.where(bitG, dense_4J, dense_4I), tl.where(
        bitG, dense_4H, dense_4G))
    sparse9 = tl.where(bitJ, tl.where(bitI, dense_4J, dense_4I), tl.where(
        bitI, dense_4H, dense_4G))
    bitK = ~mK & mL
    bitL = ~mK & ~mL
    bitM = bitL | ~mM
    bitN = bitK | ~mL | mM
    idxsA = bitK | bitL.to(tl.int64) << 1
    idxsB = bitM | bitN.to(tl.int64) << 1
    sparseA = tl.where(bitL, tl.where(bitK, dense_4N, dense_4M), tl.where(
        bitK, dense_4L, dense_4K))
    sparseB = tl.where(bitN, tl.where(bitM, dense_4N, dense_4M), tl.where(
        bitM, dense_4L, dense_4K))
    bitO = ~mO & mP
    bitP = ~mO & ~mP
    bitQ = bitP | ~mQ
    bitR = bitO | ~mP | mQ
    idxsC = bitO | bitP.to(tl.int64) << 1
    idxsD = bitQ | bitR.to(tl.int64) << 1
    sparseC = tl.where(bitP, tl.where(bitO, dense_4R, dense_4Q), tl.where(
        bitO, dense_4P, dense_4O))
    sparseD = tl.where(bitR, tl.where(bitQ, dense_4R, dense_4Q), tl.where(
        bitQ, dense_4P, dense_4O))
    bitS = ~mS & mT
    bitT = ~mS & ~mT
    bitU = bitT | ~mU
    bitV = bitS | ~mT | mU
    idxsE = bitS | bitT.to(tl.int64) << 1
    idxsF = bitU | bitV.to(tl.int64) << 1
    sparseE = tl.where(bitT, tl.where(bitS, dense_4V, dense_4U), tl.where(
        bitS, dense_4T, dense_4S))
    sparseF = tl.where(bitV, tl.where(bitU, dense_4V, dense_4U), tl.where(
        bitU, dense_4T, dense_4S))
    col_idx = tl.program_id(1) * 16
    if ARRAY_LAYOUT == 'row':
        col_idx = tl.program_id(1) * 16 * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE
            ) * 16
        mask = col_idx < k // 2
    else:
        col_idx = tl.program_id(1) * 16
        mask = row_idx < m
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 0) *
        sparse_col_stride, sparse0, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 1) *
        sparse_col_stride, sparse1, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 2) *
        sparse_col_stride, sparse2, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 3) *
        sparse_col_stride, sparse3, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 4) *
        sparse_col_stride, sparse4, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 5) *
        sparse_col_stride, sparse5, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 6) *
        sparse_col_stride, sparse6, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 7) *
        sparse_col_stride, sparse7, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 8) *
        sparse_col_stride, sparse8, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 9) *
        sparse_col_stride, sparse9, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 10) *
        sparse_col_stride, sparseA, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 11) *
        sparse_col_stride, sparseB, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 12) *
        sparse_col_stride, sparseC, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 13) *
        sparse_col_stride, sparseD, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 14) *
        sparse_col_stride, sparseE, mask=mask)
    tl.store(sparse_ptr + row_idx * sparse_row_stride + (col_idx + 15) *
        sparse_col_stride, sparseF, mask=mask)
    meta_40 = idxs0 | idxs1 << 2
    meta_41 = idxs2 | idxs3 << 2
    meta_42 = idxs4 | idxs5 << 2
    meta_43 = idxs6 | idxs7 << 2
    meta_44 = idxs8 | idxs9 << 2
    meta_45 = idxsA | idxsB << 2
    meta_46 = idxsC | idxsD << 2
    meta_47 = idxsE | idxsF << 2
    meta = (meta_40 | meta_41 << 4 | meta_42 << 8 | meta_43 << 12 | meta_44 <<
        16 | meta_45 << 20 | meta_46 << 24 | meta_47 << 28)
    if ARRAY_LAYOUT == 'row':
        col_idx = tl.program_id(1) * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    elif ARRAY_LAYOUT == 'col':
        col_idx = tl.program_id(1)
    group, interweave = 16, 2
    dest_row = (row_idx // group * group + row_idx % 8 * interweave + 
        row_idx % group // 8)
    dest_col = col_idx
    topright = ((dest_row % 2 == 0) & (dest_col % 2 == 1)).to(tl.int8)
    bottomleft = ((dest_row % 2 == 1) & (dest_col % 2 == 0)).to(tl.int8)
    dest_row = dest_row + topright - bottomleft
    dest_col = dest_col - topright + bottomleft
    interleave = 2
    cols_maj = dest_col // interleave
    cols_min = dest_col % interleave
    meta_reordered_offsets = (cols_maj * m * interleave + dest_row *
        interleave + cols_min)
    if ARRAY_LAYOUT == 'row':
        mask = col_idx < k // 32
    elif ARRAY_LAYOUT == 'col':
        mask = row_idx < m
    tl.store(meta_reordered_ptr + meta_reordered_offsets, meta, mask=mask)
