     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_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("Codeword serialization failed");
+    gpujpeg_cuda_check_error("Output compaction failed");
+    
+    
+    // TODO: read and return number of occupied bytes
     
     
     return 0;

        coder->d_segment, 
        comp_count,
        coder->segment_count, 
        coder->d_data_compressed
    #ifndef GPUJPEG_HUFFMAN_CODER_TABLES_IN_CONSTANT
        ,encoder->d_table_huffman[GPUJPEG_COMPONENT_LUMINANCE][GPUJPEG_HUFFMAN_DC]
        ,encoder->d_table_huffman[GPUJPEG_COMPONENT_LUMINANCE][GPUJPEG_HUFFMAN_AC]

    // save result packed into unsigned int (value bits are left aligned in MSBs and size is right aligned in LSBs)
    gpujpeg_huffman_value_decomposition[tid] = value_nbits | (value_code << (32 - value_nbits));
//     printf("%+04d: %08x\n", value, gpujpeg_huffman_value_decomposition[tid]);
    
    // first thread also initializes size of final output, not to have to do it in separate memcpy/kernel
    if(0 == tid) {
        gpujpeg_huffman_output_byte_count = 0;
    }
}


