 __launch_bounds__(WARPS_NUM * 32, 1024 / (WARPS_NUM * 32))
 #endif 
 __global__ static void
-gpujpeg_huffman_encoder_encode_kernel(
+gpujpeg_huffman_encoder_encode_kernel_warp(
     struct gpujpeg_segment* d_segment,
     int segment_count, 
     uint8_t* d_data_compressed,

    
    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");
    
    // Run codeword serialization kernel
    const int num_serialization_tblocks = gpujpeg_div_and_round_up(coder->segment_count, SERIALIZATION_THREADS_PER_TBLOCK);
    const dim3 serialization_grid = gpujpeg_huffman_gpu_encoder_grid_size(num_serialization_tblocks);
    gpujpeg_huffman_encoder_serialization_kernel<<<num_serialization_tblocks, SERIALIZATION_THREADS_PER_TBLOCK>>>(
        coder->d_segment, 
        coder->segment_count, 
        coder->d_temp_huffman
    );
    cudaThreadSynchronize();
    gpujpeg_cuda_check_error("Codeword serialization failed");
    
    // Run output compaction kernel (one warp per segment)
    const dim3 compaction_thread(32, WARPS_NUM);
    const dim3 compaction_grid = gpujpeg_huffman_gpu_encoder_grid_size(gpujpeg_div_and_round_up(coder->segment_count, WARPS_NUM));
    gpujpeg_huffman_encoder_compaction_kernel<<<compaction_grid, compaction_thread>>>(
        coder->d_segment,
        coder->segment_count,
        coder->d_temp_huffman,
        coder->d_data_compressed
    );
    cudaThreadSynchronize();
    gpujpeg_cuda_check_error("Huffman output compaction failed");
    
    // Read and return number of occupied bytes
    cudaMemcpyFromSymbol(output_byte_count, gpujpeg_huffman_output_byte_count, sizeof(unsigned int), 0, cudaMemcpyDeviceToHost);
    gpujpeg_cuda_check_error("Huffman output size getting failed");
    
    return 0;
}

    const uint8_t* const d_src,
    uint8_t* const d_dest
) {
#if __CUDA_ARCH__ >= 200
    
    // get some segment (size of threadblocks is 32 x N, so threadIdx.y is warp index)
    const int block_idx = blockIdx.x + blockIdx.y * gridDim.x;
