     
     int warpidx  = threadIdx.x >> 5;
     int tid    = threadIdx.x & 31;
-    int out_size = 0;
-
-    __shared__ int s_in_all[64 * WARPS_NUM];
+    int remaining_bytes = 0;
+    int remaining_bits = 0;
+
+    __shared__ int s_temp_all[64 * WARPS_NUM];
     __shared__ int s_out_all[192 * WARPS_NUM];
 
-    int * s_in  =  s_in_all + warpidx * 64;
-    int * s_out = s_out_all + warpidx * 192;
+    int * s_temp = s_temp_all + warpidx * 64;
+    int * s_out  =  s_out_all + warpidx * 192;
     
     // Select Segment
     int segment_index = blockIdx.x * WARPS_NUM + warpidx;

 * @return 0 if succeeds, otherwise nonzero
 */
__device__ int
gpujpeg_huffman_gpu_encoder_encode_block(int16_t * block, uint8_t * &data_compressed, int * s_in, int * s_out, int &out_size, 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)




// apply byte stuffing to encoded bytes
__device__ void
gpujpeg_huffman_gpu_encoder_byte_stuff(uint8_t* src, int & remaining_bits, uint8_t * & dest) {
    const int byte_count = remaining_bits / 8;
    for(int i = 0; i < byte_count; i++) {
        const uint8_t value = src[i];
        *(dest++) = value;
        if(0xff == value) {
            *(dest++) = 0;
        }
    }
    remaining_bits &= 7;
    src[0] = src[byte_count];
}


__device__ void
gpujpeg_huffman_gpu_encoder_emit_left_bits(uint8_t * &data_compressed, int * s_temp, int * s_out, int & remaining_bits, int & remaining_bytes, int tid) {
    if(tid == 0) {
        gpujpeg_huffman_gpu_encoder_emit_bits(remaining_bits, (uint8_t*)s_temp, 7, 0x7f);
        gpujpeg_huffman_gpu_encoder_byte_stuff((uint8_t*)s_temp, remaining_bits, data_compressed);
    }
}



/**
 * Encode one 8x8 block
