     for ( int k = 1; k < 64; k++ ) 
     {
         __threadfence_block();
-        const int value = data[k].x;
-        const int r = data[k].y;
+        const int size = data[k] & 31;
+        const int value = data[k] >> 5;
         __threadfence_block();
-        // if last coefficient is zero
-        if ( r == -1 ) {
-            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, (uint8_t*)data, cword_size, cword_value);
-        __threadfence_block();
+        gpujpeg_huffman_gpu_encoder_emit_bits(out_bits, (uint8_t*)data, size, value);
     }
 
     return 0;

 * @return 0 if succeeds, otherwise nonzero
 */
__device__ int
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);

    // even if even is zero or none if even value itself is nonzero)
    const int zeros_before_odd = in_even || !tid ? 0 : zeros_before_even + 1;
    
    // (TODO: remove later) save leftover bits from previous iteration
    uint8_t leftover_value = *(uint8_t*)s_temp;
    
    // clear the buffer in parallel
    ((uint64_t*)s_temp)[tid] = 0;
    
    // put leftover value back
    if(0 == tid) {
        *(uint8_t*)s_temp = leftover_value;
    }
    
    
    
    
    // pointer to LUT for encoding thread's even value 
    // (only thread #0 uses DC table, others use AC table)
    const struct gpujpeg_table_huffman_encoder * d_table_even = d_table_ac;
    
    // first thread handles special DC coefficient
    if(0 == tid) {
        // first thread uses DC table for its even value
        d_table_even = d_table_dc;
        
        // update last DC coefficient
        const int original_in_even = in_even;
        in_even -= *last_dc;
        *last_dc = original_in_even;
    }
    
    // each thread encodes its two pixels
    int even_code_size = 0, even_code_value = 0, odd_code_size = 0, odd_code_value = 0;
    if(nonzero_follows || !tid) {
        gpujpeg_huffman_gpu_encode_value(even_code_size, even_code_value, zeros_before_even & 0xf, in_even, d_table_even);
        gpujpeg_huffman_gpu_encode_value(odd_code_size, odd_code_value, zeros_before_odd & 0xf, in_odd, d_table_ac);
    }
    if(tid == 31 && !in_odd) {
        odd_code_size = d_table_ac->size[256];
        odd_code_value = d_table_ac->code[256];
    }
    
    // replace values in shared memory with tuples (value, preceding zero count)
    s_temp[load_idx] = even_code_size + 32 * even_code_value;
    s_temp[load_idx + 1] = odd_code_size + 32 * odd_code_value;
    __threadfence_block();
    
    
    int result = 0;
    if(0 == tid) {
        for ( int k = 0; k < 64; k++ ) 
        {
            const int size = s_temp[k] & 31;
            const int value = s_temp[k] >> 5;
            if(k == 0) {
                *s_temp = leftover_value;   // put leftover bits back
            }
            gpujpeg_huffman_gpu_encoder_emit_bits(remaining_bits, (uint8_t*)s_temp, size, value);
        }
        
        
        
        
        
        
        __threadfence_block();
        
        // apply byte stuffing to encoded bytes
        gpujpeg_huffman_gpu_encoder_byte_stuff((uint8_t*)s_temp, remaining_bits, (uint8_t*)s_out, remaining_bytes);
