  * 
  * @return void
  */
+template <bool CONTINUOUS_BLOCK_LIST>
 __launch_bounds__(WARPS_NUM * 32, 1024 / (WARPS_NUM * 32))
 __global__ void
 gpujpeg_huffman_encoder_encode_kernel(

    // Pre-add thread ID to output pointer (it's allways used only with it)
    data_compressed += (tid * 4);
    
    // 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;
    int block_count = segment->block_count;
    
    // Encode MCUs in segment
    while(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

    
    assert(coder->param.restart_interval > 0);
    
    // Configure more shared memory for all kernels
    cudaFuncSetCacheConfig(gpujpeg_huffman_encoder_encode_kernel<true>, cudaFuncCachePreferShared);
    cudaFuncSetCacheConfig(gpujpeg_huffman_encoder_encode_kernel<false>, cudaFuncCachePreferShared);
    cudaFuncSetCacheConfig(gpujpeg_huffman_encoder_serialization_kernel, cudaFuncCachePreferShared);
    
    // Select encoder kernel which either expects continuos segments of blocks or uses block lists
    int comp_count = 1;
    if ( coder->param.interleaved == 1 )
        comp_count = coder->param_image.comp_count;
    assert(comp_count >= 1 && comp_count <= GPUJPEG_MAX_COMPONENT_COUNT);
    
    // 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)));
    if(comp_count == 1) {
        gpujpeg_huffman_encoder_encode_kernel<true><<<grid, thread>>>(
            coder->d_segment, 
            coder->segment_count, 
            coder->d_temp_huffman,
            coder->d_block_list,
            coder->d_data_quantized,
            coder->d_component,
            comp_count
        );
    } else { 
        gpujpeg_huffman_encoder_encode_kernel<false><<<grid, thread>>>(
            coder->d_segment, 
            coder->segment_count, 
            coder->d_temp_huffman,
            coder->d_block_list,
            coder->d_data_quantized,
            coder->d_component,
            comp_count
        );
    }
    cudaThreadSynchronize();
    gpujpeg_cuda_check_error("Huffman encoding failed");
    
