 gpujpeg_huffman_gpu_decoder_table_setup(
     const int bits, 
     const struct gpujpeg_table_huffman_decoder* const d_table_src,
-    const int dest_offset
+    const int table_idx
 ) {
     // Decode one codeword from given bits to get following:
     //  - minimal number of bits actually needed to decode the codeword (up to 16 bits, 0 for invalid ones)

    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 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 + 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_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)];



