 /** Natural order in constant memory */
 __constant__ int gpujpeg_huffman_gpu_encoder_order_natural[GPUJPEG_ORDER_NATURAL_SIZE];
 
+/** Value decomposition in constant memory (input range from -4096 to 4095  ... both inclusive) */
+__device__ unsigned int gpujpeg_huffman_value_decomposition[8 * 1024];
+
 /**
  * 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) {
    unsigned int value_code = value;
    int absolute = value;
    if ( value < 0 ) {
        // valu eis now absolute value of input
        absolute = -absolute;
        // For a negative input, want temp2 = bitwise complement of abs(input)
        // This code assumes we are on a two's complement machine
        value_code--;
    }

    // Find the number of bits needed for the magnitude of the coefficient
    unsigned int value_nbits = 0;
    while ( absolute ) {
        value_nbits++;
        absolute >>= 1;
    }
    
    // upshift the value, trimming remaining bits
    value_code <<= 32 - value_nbits;
    
    // find prefix of the codeword and size of the prefix
    const int prefix_idx = preceding_zero_count * 16 + value_nbits;
    unsigned int codeword = d_table->code[prefix_idx];
    
    // compose packed codeword with its size
    return (codeword + value_nbits) | (value_code >> (codeword & 31));
}



    // 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, 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
