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

    
    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");
    

    int segment_count, 
    uint8_t* d_data_compressed,
    const uint64_t* const d_block_list,
    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;
