  * @return 0 if succeeds, otherwise nonzero
  */
 __device__ int
-gpujpeg_huffman_gpu_encoder_encode_block(int & put_value, int & put_bits, int & dc, short2* data, uint8_t* & data_compressed, 
+gpujpeg_huffman_gpu_encoder_encode_block(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;
     
-    int temp2 = temp;
-    if ( temp < 0 ) {
-        // Temp is abs value of input
-        temp = -temp;
-        // For a negative input, want temp2 = bitwise complement of abs(input)
-        // This code assumes we are on a two's complement machine
-        temp2--;
-    }
-
-    // Find the number of bits needed for the magnitude of the coefficient
-    int nbits = 0;
-    while ( temp ) {
-        nbits++;
-        temp >>= 1;
-    }
-
+    int temp2, nbits;
+    gpujpeg_huffman_gpu_encoder_decompose(temp, nbits, temp2);
+    
     // Write category number
-    if ( gpujpeg_huffman_gpu_encoder_emit_bits(d_table_dc->code[nbits], d_table_dc->size[nbits], put_value, put_bits, data_compressed) != 0 ) {
-        return -1;
-    }
+    gpujpeg_huffman_gpu_encoder_emit_bits(out_bits, data_compressed, d_table_dc->size[nbits], d_table_dc->code[nbits]);
 
     // Write category offset (EmitBits rejects calls with size 0)
-    if ( gpujpeg_huffman_gpu_encoder_emit_bits((unsigned int) temp2, nbits, put_value, put_bits, data_compressed) != 0 )
-        return -1;
-    
-//     // find last nonzero value
-//     int last_nonzero_idx = 64;
-//     while( last_nonzero_idx-- ) {
-//         if( data[last_nonzero_idx].x ) {
-//             break;
-//         }
-//     }
+    gpujpeg_huffman_gpu_encoder_emit_bits(out_bits, data_compressed, nbits, temp2);
     
     // Encode the AC coefficients per section F.1.2.2 (r = run length of zeros)
     for ( int k = 1; k < 64; k++ ) 

    *data_compressed = (uint8_t)(marker); \
    data_compressed++; }

/**
 * Output bits to the file. Only the right 24 bits of put_buffer are used; 
 * the valid bits are left-justified in this part.  At most 16 bits can be 
 * passed to EmitBits in one call, and we never retain more than 7 bits 
 * in put_buffer between calls, so 24 bits are sufficient.
 * 
 * @param coder  Huffman coder structure
 * @param code  Huffman code
 * @param size  Size in bits of the Huffman code
 * @return void
 */
__device__ inline int
gpujpeg_huffman_gpu_encoder_emit_bits(unsigned int code, int size, int & put_value, int & put_bits, uint8_t* & data_compressed)
{
    // This routine is heavily used, so it's worth coding tightly
    int _put_buffer = (int)code;
    int _put_bits = put_bits;
    // Mask off any extra bits in code
    _put_buffer &= (((int)1) << size) - 1; 
    // New number of bits in buffer
    _put_bits += size;                    
    // Align incoming bits
    _put_buffer <<= 24 - _put_bits;        
    // And merge with old buffer contents
    _put_buffer |= put_value;    
    // If there are more than 8 bits, write it out
    unsigned char uc;
    while ( _put_bits >= 8 ) {
        // Write one byte out
        uc = (unsigned char) ((_put_buffer >> 16) & 0xFF);
        gpujpeg_huffman_gpu_encoder_emit_byte(data_compressed, uc);
        // If need to stuff a zero byte
        if ( uc == 0xFF ) {  
            // Write zero byte out
            gpujpeg_huffman_gpu_encoder_emit_byte(data_compressed, 0);
        }
        _put_buffer <<= 8;
        _put_bits -= 8;
    }
    // update state variables
    put_value = _put_buffer; 
    put_bits = _put_bits;
    return 0;
}

/**
 * Emit left bits
 * 
 * @param coder  Huffman coder structure
 * @return void
 */
__device__ inline void
gpujpeg_huffman_gpu_encoder_emit_left_bits(int & put_value, int & put_bits, uint8_t* & data_compressed)
{
    // Fill 7 bits with ones
    if ( gpujpeg_huffman_gpu_encoder_emit_bits(0x7F, 7, put_value, put_bits, data_compressed) != 0 )
        return;
    
    //unsigned char uc = (unsigned char) ((put_value >> 16) & 0xFF);
    // Write one byte out
    //gpujpeg_huffman_gpu_encoder_emit_byte(data_compressed, uc);
    
    put_value = 0; 
    put_bits = 0;
}

/**
 * Decomposes given value into number of bits and one's complement value.


__device__ void
gpujpeg_huffman_gpu_encoder_emit_left_bits(uint8_t * &data_compressed, int * s_out, int &out_size, int tid) {
    if(tid == 0) {
        gpujpeg_huffman_gpu_encoder_emit_bits(out_size, data_compressed, 7, 0x7f);
        data_compressed += (out_size >> 3);
    }
}

