                         }
                         
                         // Encode 8x8 block
-                        gpujpeg_huffman_gpu_encoder_encode_block(block, data_compressed, &last_dc, tid, d_table_dc, d_table_ac);
+                        gpujpeg_huffman_gpu_encoder_encode_block(block, data_compressed, s_out, remaining_codewords, &last_dc, tid, d_table_dc, d_table_ac);
                     }
                 }
             }
         }
     }
 
+    // flush remaining codewords
+    gpujpeg_huffman_gpu_encoder_flush_codewords(s_out, data_compressed, remaining_codewords, tid);
+    
     // Set number of codewords.
     if (tid == 0 ) {
-        segment->data_compressed_size = (data_compressed - data_compressed_start) * 4;
+        segment->data_compressed_size = data_compressed - data_compressed_start;
     }
     __syncthreads();
 }

    int warpidx = threadIdx.x >> 5;
    int tid = threadIdx.x & 31;

//     __shared__ int s_out_all[256 * WARPS_NUM];
// 
//     uint4 * s_out  =  s_out_all + warpidx * 256;
    
    // Select Segment
    int segment_index = blockIdx.x * WARPS_NUM + warpidx;

        dc[comp] = 0;
    
    // Prepare data pointers
    unsigned int * data_compressed = (unsigned int*)(d_data_compressed + segment->data_compressed_index);
    unsigned int * data_compressed_start = data_compressed;
    
    // Non-interleaving mode
    if ( comp_count == 1 ) {
