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

    // Configure more shared memory
    cudaFuncSetCacheConfig(gpujpeg_huffman_encoder_encode_kernel, cudaFuncCachePreferShared);
    cudaFuncSetCacheConfig(gpujpeg_huffman_encoder_serialization_kernel, cudaFuncCachePreferShared);
            
    // Run encoder kernel
    dim3 thread(32 * WARPS_NUM);
    dim3 grid(gpujpeg_div_and_round_up(coder->segment_count, (thread.x / 32)));


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


