             // Get coder parameters
             int & component_dc = dc[segment->scan_index];
             
-            // Get huffman tables
-            struct gpujpeg_table_huffman_decoder* d_table_dc = NULL;
-            struct gpujpeg_table_huffman_decoder* d_table_ac = NULL;
-            if ( component->type == GPUJPEG_COMPONENT_LUMINANCE ) {
-                d_table_dc = d_table_y_dc;
-                d_table_ac = d_table_y_ac;
-            } else {
-                d_table_dc = d_table_cbcr_dc;
-                d_table_ac = d_table_cbcr_ac;
-            }
+            // Get huffman tables offset
+            const unsigned int table_offset = component->type == GPUJPEG_COMPONENT_LUMINANCE ? 0x00000 : 0x20000;
             
             // Encode 8x8 block
-            if ( gpujpeg_huffman_gpu_decoder_decode_block(component_dc, block, d_table_dc, d_table_ac, r_bit, r_bit_count, s_byte, s_byte_idx, d_byte, d_byte_chunk_count) != 0 )
+            if ( gpujpeg_huffman_gpu_decoder_decode_block(component_dc, block, table_offset, r_bit, r_bit_count, s_byte, s_byte_idx, d_byte, d_byte_chunk_count) != 0 )
                 break;
         } 
     }

 * @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);
    }
}

/**

 */
__device__ inline int
gpujpeg_huffman_gpu_decoder_decode_block(
    int & dc, int16_t* const data_output, const unsigned int table_offset,
    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)
{
    // Index of next coefficient to be decoded (in ZIG-ZAG order)
    unsigned int coefficient_idx = 0;
    
    // Section F.2.2.1: decode the DC coefficient difference
    // Get the coefficient value (using DC coding table)
    int dc_coefficient_value = gpujpeg_huffman_gpu_decoder_get_coefficient(r_bit, r_bit_count, s_byte, s_byte_idx, d_byte, d_byte_chunk_count, table_offset, coefficient_idx);

    // Convert DC difference to actual value, update last_dc_val
    dc = dc_coefficient_value += dc;

    // Output the DC coefficient (assumes gpujpeg_natural_order[0] = 0)
    // TODO: try to skip saving of zero coefficients
    data_output[0] = dc_coefficient_value;
    
    // TODO: error check: coefficient_idx must still be 0 in valid codestreams
    coefficient_idx = 1;
    
    // Section F.2.2.2: decode the AC coefficients
    // Since zeroes are skipped, output area must be cleared beforehand
    do {
        // decode next coefficient, updating its destination index
        const int coefficient_value = gpujpeg_huffman_gpu_decoder_get_coefficient(r_bit, r_bit_count, s_byte, s_byte_idx, d_byte, d_byte_chunk_count, table_offset + 0x10000, coefficient_idx);
        
        // stop with this block if have all coefficients
        if(coefficient_idx >= 64) {
            break;
        }
        
        // save the coefficient   TODO: try to ommit saving 0 coefficients
        data_output[gpujpeg_huffman_gpu_decoder_order_natural[coefficient_idx++]] = coefficient_value;
    } while(coefficient_idx < 64);
    
    return 0;
}
