     // bits loaded into the register and their count
     unsigned int r_bit_count = 0;
     unsigned int r_bit = 0; // LSB-aligned
+    
+    // TODO: try using block lists
     
     // Non-interleaving mode
     if ( comp_count == 1 ) {

    __shared__ uint4 s_byte_all[2 * THREADS_PER_TBLOCK]; // 32 bytes per thread
    uint4 * const s_byte = s_byte_all + 2 * threadIdx.x;
    
    // Last DC coefficient values
    int dc[GPUJPEG_MAX_COMPONENT_COUNT];
    for ( int comp = 0; comp < GPUJPEG_MAX_COMPONENT_COUNT; comp++ )
        dc[comp] = 0;

    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)
{
    // TODO: try unified decoding of DC/AC coefficients
    
    // Index of next coefficient to be decoded (in ZIG-ZAG order)
    unsigned int coefficient_idx = 0;
    
