  * @return 0 if succeeds, otherwise nonzero
  */
 __device__ int
-gpujpeg_huffman_gpu_encoder_encode_block(int & out_bits, int & dc, short2* data, uint8_t* & data_compressed, 
+gpujpeg_huffman_gpu_encoder_encode_block(const int leftover_value, int & out_bits, int & dc, short2* data, uint8_t* & data_compressed, 
     struct gpujpeg_table_huffman_encoder* d_table_dc, struct gpujpeg_table_huffman_encoder* d_table_ac)
 {
     // Encode the DC coefficient difference per section F.1.2.1
     int temp = data[0].x - dc;
     dc = data[0].x;
     
+    __threadfence_block();
+    // put leftover bits back
+    *(int*)data = leftover_value;
+    
     int cword_size, cword_value;
     gpujpeg_huffman_gpu_encode_value(cword_size, cword_value, 0, temp, d_table_dc);
-    gpujpeg_huffman_gpu_encoder_emit_bits(out_bits, data_compressed, cword_size, cword_value);
+    gpujpeg_huffman_gpu_encoder_emit_bits(out_bits, (uint8_t*)data, cword_size, cword_value);
     
     // Encode the AC coefficients per section F.1.2.2 (r = run length of zeros)
     for ( int k = 1; k < 64; k++ ) 
     {
+        __threadfence_block();
         const int value = data[k].x;
         const int r = data[k].y;
-
+        __threadfence_block();
         // if last coefficient is zero
         if ( r == -1 ) {
-            gpujpeg_huffman_gpu_encoder_emit_bits(out_bits, data_compressed, d_table_ac->size[256], d_table_ac->code[256]);
+            gpujpeg_huffman_gpu_encoder_emit_bits(out_bits, (uint8_t*)data, d_table_ac->size[256], d_table_ac->code[256]);
             break;
         }
         
         gpujpeg_huffman_gpu_encode_value(cword_size, cword_value, r, value, d_table_ac);
-        gpujpeg_huffman_gpu_encoder_emit_bits(out_bits, data_compressed, cword_size, cword_value);
+        gpujpeg_huffman_gpu_encoder_emit_bits(out_bits, (uint8_t*)data, cword_size, cword_value);
+        __threadfence_block();
     }
 
     return 0;

        // 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, s_in, s_out, out_size, &last_dc, tid, d_table_dc, d_table_ac) != 0)
                break;
            block += component->mcu_size;
        }

                        }
                        
                        // Encode 8x8 block
                        gpujpeg_huffman_gpu_encoder_encode_block(block, data_compressed, s_temp, s_out, remaining_bits, remaining_bytes, &last_dc, tid, d_table_dc, d_table_ac);
                    }
                }
            }
