     int result = 0;
     if(0 == tid) {
         int & dc = *last_dc;
-        int & put_value = s_out[0];
-        int & put_bits = s_out[1];
-        if(out_size == 0) {
-            put_value = 0;
-            put_bits = 0;
-            out_size = 1;
-        }
-        result = gpujpeg_huffman_gpu_encoder_encode_block(put_value, put_bits, dc, (short2*)s_in, data_compressed, d_table_dc, d_table_ac);
+        result = gpujpeg_huffman_gpu_encoder_encode_block(out_size, dc, (short2*)s_in, data_compressed, d_table_dc, d_table_ac);
     }
     return __ballot(result);
     

//     int even_bit_size
//     if()
//     
//     
//     typedef uint64_t loading_t;
//     const int loading_iteration_count = 64 * 2 / sizeof(loading_t);
//     
//     // Load block to shared memory
//     __shared__ int16_t s_data[64 * THREAD_BLOCK_SIZE];
//     for ( int i = 0; i < loading_iteration_count; i++ ) {
//         ((loading_t*)s_data)[loading_iteration_count * threadIdx.x + i] = ((loading_t*)data)[i];
//     }
//     int data_start = 64 * threadIdx.x;
// 
//     // Encode the DC coefficient difference per section F.1.2.1
//     int temp = s_data[data_start + 0] - dc;
//     dc = s_data[data_start + 0];
//     
//     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;
//     }
// 
//     // 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;
//     }
// 
//     // Write category offset (EmitBits rejects calls with size 0)
//     if ( nbits ) {
//         if ( gpujpeg_huffman_gpu_encoder_emit_bits((unsigned int) temp2, nbits, put_value, put_bits, data_compressed) != 0 )
//             return -1;
//     }
//     
//     // Encode the AC coefficients per section F.1.2.2 (r = run length of zeros)
//     int r = 0;
//     for ( int k = 1; k < 64; k++ ) 
//     {
//         temp = s_data[data_start + gpujpeg_huffman_gpu_encoder_order_natural[k]];
//         if ( temp == 0 ) {
//             r++;
//         }
//         else {
//             // If run length > 15, must emit special run-length-16 codes (0xF0)
//             while ( r > 15 ) {
//                 if ( gpujpeg_huffman_gpu_encoder_emit_bits(d_table_ac->code[0xF0], d_table_ac->size[0xF0], put_value, put_bits, data_compressed) != 0 )
//                     return -1;
//                 r -= 16;
//             }
// 
//             temp2 = temp;
//             if ( temp < 0 ) {
//                 // temp is abs value of input
//                 temp = -temp;        
//                 // This code assumes we are on a two's complement machine
//                 temp2--;
//             }
// 
//             // Find the number of bits needed for the magnitude of the coefficient
//             // there must be at least one 1 bit
//             nbits = 1;
//             while ( (temp >>= 1) )
//                 nbits++;
// 
//             // Emit Huffman symbol for run length / number of bits
//             int i = (r << 4) + nbits;
//             if ( gpujpeg_huffman_gpu_encoder_emit_bits(d_table_ac->code[i], d_table_ac->size[i], put_value, put_bits, data_compressed) != 0 )
//                 return -1;
// 
//             // Write Category offset
//             if ( gpujpeg_huffman_gpu_encoder_emit_bits((unsigned int) temp2, nbits, put_value, put_bits, data_compressed) != 0 )
//                 return -1;
// 
//             r = 0;
//         }
//     }
// 
//     // If all the left coefs were zero, emit an end-of-block code
//     if ( r > 0 ) {
//         if ( gpujpeg_huffman_gpu_encoder_emit_bits(d_table_ac->code[0], d_table_ac->size[0], put_value, put_bits, data_compressed) != 0 )
//             return -1;
//     }
// 
//     return 0;
}

/**


__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);
    }
}

