             coder->d_segment, 
             comp_count,
             decoder->segment_count,
-            coder->d_data_compressed
+            coder->d_data_compressed,
+            coder->d_block_list,
+            coder->d_data_quantized
         );
     } else {
         gpujpeg_huffman_decoder_decode_kernel<false, THREADS_PER_TBLOCK><<<grid, thread>>>(

            coder->d_segment, 
            comp_count,
            decoder->segment_count,
            coder->d_data_compressed
        );
    }
    cudaError cuerr = cudaThreadSynchronize();

    struct gpujpeg_segment* d_segment,
    int comp_count,
    int segment_count, 
    uint8_t* d_data_compressed,
    const uint64_t* d_block_list,
    int16_t* d_data_quantized
) {
    int segment_index = blockIdx.x * THREADS_PER_TBLOCK + threadIdx.x;
    if ( segment_index >= segment_count )
