@triton.jit
def _sparse_semi_structured_to_dense_kernel(sparse_ptr, meta_reordered_ptr,
    dense_ptr, m, k, BLOCK_SIZE: tl.constexpr):
    row_idx = tl.program_id(0)
    group, interweave = 32, 4
    dest_row = row_idx // 32 * 32 + row_idx % 8 * 4 + row_idx % group // 8
    if dest_row % 2 == 0:
        dest_row_ = (row_idx // 32 * 32 + row_idx % 8 * 4 + row_idx % group //
            8 + tl.arange(0, BLOCK_SIZE // 16) % 2)
        dest_col_ = tl.arange(0, BLOCK_SIZE // 16) // 2 * 2
        index = dest_col_ // 2 * m * 2 + dest_row_ * 2 + dest_col_ % 2
        meta = tl.load(meta_reordered_ptr + index, mask=tl.arange(0, 
            BLOCK_SIZE // 16) < k // 16, other=-float('inf'))
    else:
        dest_row_ = (row_idx // 32 * 32 + row_idx % 8 * 4 + row_idx % group //
            8 - (tl.arange(0, BLOCK_SIZE // 16) + 1) % 2)
        dest_col_ = tl.arange(0, BLOCK_SIZE // 16) // 2 * 2 + 1
        index = dest_col_ // 2 * m * 2 + dest_row_ * 2 + dest_col_ % 2
        meta = tl.load(meta_reordered_ptr + index, mask=tl.arange(0, 
            BLOCK_SIZE // 16) < k // 16, other=-float('inf'))
    meta_20 = (meta & 3) + (row_idx * k + 16 * tl.arange(0, BLOCK_SIZE // 16))
    meta_21 = (meta >> 2 & 3) + (row_idx * k + 16 * tl.arange(0, BLOCK_SIZE //
        16))
    meta_22 = (meta >> 4 & 3) + (row_idx * k + 16 * tl.arange(0, BLOCK_SIZE //
        16) + 4)
    meta_23 = (meta >> 6 & 3) + (row_idx * k + 16 * tl.arange(0, BLOCK_SIZE //
        16) + 4)
    meta_24 = (meta >> 8 & 3) + (row_idx * k + 16 * tl.arange(0, BLOCK_SIZE //
        16) + 8)
    meta_25 = (meta >> 10 & 3) + (row_idx * k + 16 * tl.arange(0, 
        BLOCK_SIZE // 16) + 8)
    meta_26 = (meta >> 12 & 3) + (row_idx * k + 16 * tl.arange(0, 
        BLOCK_SIZE // 16) + 12)
    meta_27 = (meta >> 14 & 3) + (row_idx * k + 16 * tl.arange(0, 
        BLOCK_SIZE // 16) + 12)
    row0 = tl.load(sparse_ptr + row_idx * k // 2 + 8 * tl.arange(0, 
        BLOCK_SIZE // 16), mask=tl.arange(0, BLOCK_SIZE // 16) < k // 16,
        other=-float('inf'))
    row1 = tl.load(sparse_ptr + row_idx * k // 2 + 8 * tl.arange(0, 
        BLOCK_SIZE // 16) + 1, mask=tl.arange(0, BLOCK_SIZE // 16) < k // 
        16, other=-float('inf'))
    row2 = tl.load(sparse_ptr + row_idx * k // 2 + 8 * tl.arange(0, 
        BLOCK_SIZE // 16) + 2, mask=tl.arange(0, BLOCK_SIZE // 16) < k // 
        16, other=-float('inf'))
    row3 = tl.load(sparse_ptr + row_idx * k // 2 + 8 * tl.arange(0, 
        BLOCK_SIZE // 16) + 3, mask=tl.arange(0, BLOCK_SIZE // 16) < k // 
        16, other=-float('inf'))
    row4 = tl.load(sparse_ptr + row_idx * k // 2 + 8 * tl.arange(0, 
        BLOCK_SIZE // 16) + 4, mask=tl.arange(0, BLOCK_SIZE // 16) < k // 
        16, other=-float('inf'))
    row5 = tl.load(sparse_ptr + row_idx * k // 2 + 8 * tl.arange(0, 
        BLOCK_SIZE // 16) + 5, mask=tl.arange(0, BLOCK_SIZE // 16) < k // 
        16, other=-float('inf'))
    row6 = tl.load(sparse_ptr + row_idx * k // 2 + 8 * tl.arange(0, 
        BLOCK_SIZE // 16) + 6, mask=tl.arange(0, BLOCK_SIZE // 16) < k // 
        16, other=-float('inf'))
    row7 = tl.load(sparse_ptr + row_idx * k // 2 + 8 * tl.arange(0, 
        BLOCK_SIZE // 16) + 7, mask=tl.arange(0, BLOCK_SIZE // 16) < k // 
        16, other=-float('inf'))
    tl.store(dense_ptr + meta_20, row0, mask=tl.arange(0, BLOCK_SIZE // 16) <
        k // 16)
    tl.store(dense_ptr + meta_21, row1, mask=tl.arange(0, BLOCK_SIZE // 16) <
        k // 16)
    tl.store(dense_ptr + meta_22, row2, mask=tl.arange(0, BLOCK_SIZE // 16) <
        k // 16)
    tl.store(dense_ptr + meta_23, row3, mask=tl.arange(0, BLOCK_SIZE // 16) <
        k // 16)
    tl.store(dense_ptr + meta_24, row4, mask=tl.arange(0, BLOCK_SIZE // 16) <
        k // 16)
    tl.store(dense_ptr + meta_25, row5, mask=tl.arange(0, BLOCK_SIZE // 16) <
        k // 16)
    tl.store(dense_ptr + meta_26, row6, mask=tl.arange(0, BLOCK_SIZE // 16) <
        k // 16)
    tl.store(dense_ptr + meta_27, row7, mask=tl.arange(0, BLOCK_SIZE // 16) <
        k // 16)
