     const int rle_zero_count = category_id ? min(category_id >> 4, 63) : 63;
     
     // save all the info into the right place in the destination table
-    d_table_dest->codes[bits].code_nbits = code_nbits;
-    d_table_dest->codes[bits].value_nbits = value_nbits;
-    d_table_dest->codes[bits].rle_zero_count = rle_zero_count;
+    gpujpeg_huffman_gpu_decoder_tables[dest_offset + bits].code_nbits = code_nbits;
+    gpujpeg_huffman_gpu_decoder_tables[dest_offset + bits].value_nbits = value_nbits;
+    gpujpeg_huffman_gpu_decoder_tables[dest_offset + bits].rle_zero_count = rle_zero_count;
 }
 
 /**

gpujpeg_huffman_gpu_decoder_table_setup(
    const int bits, 
    const struct gpujpeg_table_huffman_decoder* const d_table_src,
    struct gpujpeg_table_huffman_decoder_fast* const d_table_dest
) {
    // 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)

) {
    // 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, 0x00000);
    gpujpeg_huffman_gpu_decoder_table_setup(idx, d_table_y_ac, 0x10000);
    gpujpeg_huffman_gpu_decoder_table_setup(idx, d_table_cbcr_dc, 0x20000);
    gpujpeg_huffman_gpu_decoder_table_setup(idx, d_table_cbcr_ac, 0x30000);
}

/** Documented at declaration */
