         // Encode MCUs in segment
         for ( int mcu_index = 0; mcu_index < segment->mcu_count; mcu_index++ ) {
             // Encode 8x8 block
-            if (gpujpeg_huffman_gpu_encoder_encode_block(block, data_compressed, &last_dc, tid, d_table_dc, d_table_ac) != 0)
+            if (gpujpeg_huffman_gpu_encoder_encode_block(block, data_compressed, s_out, remaining_codewords, &last_dc, tid, d_table_dc, d_table_ac) != 0)
                 break;
             block += component->mcu_size;
         }

    ((uint2*)data_compressed)[tid] = codewords;
    
    // advance the pointer
    data_compressed += (32 * sizeof(uint2)) / sizeof(*data_compressed);
    
    // nothing to fail here
    return 0;

}


__device__ void
gpujpeg_huffman_gpu_encoder_flush_codewords(unsigned int * const s_out, unsigned int * &data_compressed, int & remaining_codewords, const int tid) {
    // this works for up to 4 * 32 remaining codewords
    if(remaining_codewords) {
        // pad remianing codewords with extra zero-sized codewords, not to have to use special case in serialization kernel, which saves 4 codewords at once
        s_out[remaining_codewords + tid] = 0;
        
        // save all remaining codewords at once (together with some zero sized padding codewords)
        ((uint4*)data_compressed)[tid] = ((uint4*)s_out)[tid];
        
        // update codeword counter
        data_compressed += remaining_codewords;
        remaining_codewords = 0;
    }
}


/**
 * Encode one 8x8 block
 *
 * @return 0 if succeeds, otherwise nonzero
 */
__device__ int
gpujpeg_huffman_gpu_encoder_encode_block(int16_t * block, unsigned int * &data_compressed, unsigned int * const s_out,
                int & remaining_codewords, int *last_dc, int tid,
                struct gpujpeg_table_huffman_encoder* d_table_dc, struct gpujpeg_table_huffman_encoder* d_table_ac)
{
    // each thread loads a pair of values (pair after zigzag reordering)
