     // each thread gets codeword for its two pixels
     unsigned int even_code = 0, odd_code = 0;
     if(nonzero_follows || !tid) {
-        even_code = gpujpeg_huffman_gpu_encode_value(zeros_before_even & 0xf, in_even, d_table_even);
-        odd_code = gpujpeg_huffman_gpu_encode_value(zeros_before_odd & 0xf, in_odd, d_table_ac);
+        even_code = gpujpeg_huffman_gpu_encode_value(zeros_before_even, in_even, d_table_even);
+        odd_code = gpujpeg_huffman_gpu_encode_value(zeros_before_odd, in_odd, d_table_ac);
     }
     
     // last thread writes "end of block" value if last coefficient is zero

/** Natural order in constant memory */
__constant__ int gpujpeg_huffman_gpu_encoder_order_natural[GPUJPEG_ORDER_NATURAL_SIZE];

/**
 * Write marker to compressed data
 * 


__device__ static unsigned int
gpujpeg_huffman_gpu_encode_value(const int preceding_zero_count, const int value,
                                 const struct gpujpeg_table_huffman_encoder * const d_table)
{
    // value bits are in MSBs (left aligned) and bit size of the value is in LSBs (right aligned)
    const unsigned int packed_value = gpujpeg_huffman_value_decomposition[4096 + value];
    
    // decompose value info into upshifted value and value's bit size
    const int value_nbits = packed_value & 0xf;
    const unsigned int value_code = packed_value & ~0xf;
    
    // find prefix of the codeword and size of the prefix
    const int prefix_idx = (preceding_zero_count & 0xf) * 16 + value_nbits;
    const unsigned int codeword = d_table->code[prefix_idx];
    const unsigned int prefix_nbits = codeword & 31;
    
    // compose packed codeword with its size
    return (codeword + value_nbits) | (value_code >> prefix_nbits);
}


