     // Since zeroes are skipped, output area must be cleared beforehand
     for ( int k = 1; k < 64; k++ ) {
         // s: (run, category)
-//         int s = gpujpeg_huffman_gpu_decoder_get_category(get_bits, get_buff, data, data_size, table_ac);
-        int s = gpujpeg_huffman_gpu_decoder_decode_special_decode(table_ac, 1, r_bit, r_bit_count, s_byte, s_byte_idx, d_byte, d_byte_chunk_count);
+        int s = gpujpeg_huffman_gpu_decoder_get_category(r_bit, r_bit_count, s_byte, s_byte_idx, d_byte, d_byte_chunk_count, table_ac);
         // r: run length for ac zero, 0 <= r < 16
         int r = s >> 4;
         // s: category for this non-zero ac

    return (r_bit >> r_bit_count) & ((1 << nbits) - 1);
}


/**
 * Special Huffman decode:

    return (offset < half[category]) ? (offset + start[category]) : offset;    
}

/**
 * Get category number for dc, or (0 run length, ac category) for ac.
 * The max length for Huffman codes is 15 bits; so we use 32 bits buffer    
 * m_nGetBuff, with the validated length is m_nGetBits.
 * Usually, more than 95% of the Huffman codes will be 8 or fewer bits long
 * To speed up, we should pay more attention on the codes whose length <= 8
 * 
 * @param table
 * @param get_bits
 * @param get_buff
 * @param data
 * @param data_size
 * @return int
 */
__device__ inline int
gpujpeg_huffman_gpu_decoder_get_category(
                unsigned int & r_bit, unsigned int & r_bit_count, uint4* const s_byte,
                unsigned int & s_byte_idx, const uint4* & d_byte, unsigned int & d_byte_chunk_count,
                const struct gpujpeg_table_huffman_decoder* const table)
{
    // Peek the first valid byte
    const unsigned int look = gpujpeg_huffman_gpu_decoder_peek_bits(8, r_bit, r_bit_count, s_byte, s_byte_idx, d_byte, d_byte_chunk_count);
    const int nb = table->look_nbits[look];

    if ( nb ) { 
        gpujpeg_huffman_gpu_decoder_discard_bits(nb, r_bit, r_bit_count);
        return table->look_sym[look]; 
    } else {
        //Decode long codes with length >= 9
        return gpujpeg_huffman_gpu_decoder_decode_special_decode(table, 9, r_bit, r_bit_count, s_byte, s_byte_idx, d_byte, d_byte_chunk_count);
    }
}

/**
 * Decode one 8x8 block
