  * 
  * Each entry consists of:
  *   - Number of bits of code corresponding to this entry (0 - 16, both inclusive) - bits 4 to 8
- *   - Number of run-length coded zeros before currently decoded coefficient (0 - 64, both inclusive) - bits 9 to 15
+ *   - Number of run-length coded zeros before currently decoded coefficient + 1 (1 - 64, both inclusive) - bits 9 to 15
  *   - Number of bits representing the value of currently decoded coefficient (0 - 15, both inclusive) - bits 0 to 3
  * bit #:    15                      9   8               4   3           0
  *         +---+---+---+---+---+---+---+---+---+---+---+---+---+---+---+---+
  * value:  |      RLE zero count       |   code bit size   | value bit size|
  *         +---+---+---+---+---+---+---+---+---+---+---+---+---+---+---+---+
  */
-__device__ uint16_t gpujpeg_huffman_gpu_decoder_tables[4 * (1 << 16)];
-
+__device__ uint16_t gpujpeg_huffman_gpu_decoder_tables_full[4 * (1 << 16)];
+
+/** Number of code bits to be checked first (with high chance for the code to fit into this number of bits). */
+#define QUICK_CHECK_BITS 8
+
+/** Table with same format as the full table, except that all-zero-entry means that the full table should be consulted. */
+__device__ uint16_t gpujpeg_huffman_gpu_decoder_tables_quick[4 * (1 << QUICK_CHECK_BITS)];
 
 
 

    const int rle_zero_count = category_id ? min(1 + (category_id >> 4), 64) : 64;
    
    // save all the info into the right place in the destination table
    gpujpeg_huffman_gpu_decoder_tables[dest_offset + bits] = (rle_zero_count << 9) + (code_nbits << 4) + value_nbits;
}

/**

) {
    // Each thread uses all 4 Huffman tables to "decode" one symbol from its unique 16bits.
    const int idx = threadIdx.x + blockIdx.x * blockDim.x;
    gpujpeg_huffman_gpu_decoder_table_setup(idx, d_table_y_dc, 0);
    gpujpeg_huffman_gpu_decoder_table_setup(idx, d_table_y_ac, 1);
    gpujpeg_huffman_gpu_decoder_table_setup(idx, d_table_cbcr_dc, 2);
    gpujpeg_huffman_gpu_decoder_table_setup(idx, d_table_cbcr_ac, 3);
}

/** Documented at declaration */
