     int segment_count, 
     uint8_t* d_data_compressed,
     const uint64_t* const d_block_list,
-    int16_t* const d_data_quantized
+    int16_t* const d_data_quantized,
+    struct gpujpeg_component* const d_component,
+    const int comp_count
 ) {    
     int warpidx = threadIdx.x >> 5;
     int tid = threadIdx.x & 31;

    
    assert(coder->param.restart_interval > 0);
    
    // Configure more shared memory
    cudaFuncSetCacheConfig(gpujpeg_huffman_encoder_encode_kernel, cudaFuncCachePreferShared);
    cudaFuncSetCacheConfig(gpujpeg_huffman_encoder_serialization_kernel, cudaFuncCachePreferShared);
    
    // Run encoder kernel
    dim3 thread(32 * WARPS_NUM);
    dim3 grid = gpujpeg_huffman_gpu_encoder_grid_size(gpujpeg_div_and_round_up(coder->segment_count, (thread.x / 32)));
    gpujpeg_huffman_encoder_encode_kernel<<<grid, thread>>>(
        coder->d_segment, 
        coder->segment_count, 
        coder->d_temp_huffman,
        coder->d_block_list,
        coder->d_data_quantized
    );
    cudaThreadSynchronize();
    gpujpeg_cuda_check_error("Huffman encoding failed");
    

    // Pre-add thread ID to output pointer (it's allways used only with it)
    data_compressed += (tid * 4);
    
    // Encode all block in segment
    if(CONTINUOUS_BLOCK_LIST) {
        // Get component for current scan
        const struct gpujpeg_component* component = &d_component[segment->scan_index];
        
        // mcu size of the component
        const int comp_mcu_size = component->mcu_size;
        
        // Get component data for MCU (first block)
        const int16_t* block = component->d_data_quantized + (segment->scan_segment_index * component->segment_mcu_count) * comp_mcu_size;
        
        // Get huffman table offset
        const int huffman_table_offset = component->type == GPUJPEG_COMPONENT_LUMINANCE ? 0 : (256 + 1) * 2; // possibly skips luminance tables
        
        // Encode MCUs in segment
        for (int block_count = segment->mcu_count; block_count--;) {
            // Encode 8x8 block
            gpujpeg_huffman_gpu_encoder_encode_block(block, data_compressed, s_out, remaining_codewords, 256, tid, huffman_table_offset);
            
            // Advance to next block
            block += comp_mcu_size;
        }
    } else {
        // Pointer to segment's list of 8x8 blocks and their count
        const uint64_t* packed_block_info_ptr = d_block_list + segment->block_index_list_begin;
        
        // Encode all blocks
        for(int block_count = segment->block_count; block_count--;) {
            // Get pointer to next block input data and info about its color type
            const uint64_t packed_block_info = *(packed_block_info_ptr++);
            
            // Get coder parameters
            const int last_dc_idx = 256 + (packed_block_info & 0x7f);
            
            // Get offset to right part of huffman table
            const int huffman_table_offset = packed_block_info & 0x80 ? (256 + 1) * 2 : 0; // possibly skips luminance tables
            
            // Source data pointer
            int16_t* block = &d_data_quantized[packed_block_info >> 8];
                        
            // Encode 8x8 block
            gpujpeg_huffman_gpu_encoder_encode_block(block, data_compressed, s_out, remaining_codewords, last_dc_idx, tid, huffman_table_offset);
        }
    }

    // flush remaining codewords
