     
     struct gpujpeg_segment* segment = &d_segment[segment_index];
     
-    // Start coder
-    int get_buff = 0;
-    int get_bits = 0;
+    // Byte buffers in shared memory
+    __shared__ uint4 s_byte_all[2 * THREADS_PER_TBLOCK]; // 32 bytes per thread
+    uint4 * const s_byte = s_byte_all + 2 * threadIdx.x;
+    
+    // Last DC coefficient values
     int dc[GPUJPEG_MAX_COMPONENT_COUNT];
     for ( int comp = 0; comp < GPUJPEG_MAX_COMPONENT_COUNT; comp++ )
         dc[comp] = 0;
         
-    // Get compressed data
-    uint8_t* data_compressed = &d_data_compressed[segment->data_compressed_index];
-    int data_compressed_size = segment->data_compressed_size;
+    // Get aligned compressed data chunk pointer and load first 2 chunks of the data
+    const unsigned int d_byte_begin_idx = segment->data_compressed_index;
+    const unsigned int d_byte_begin_idx_aligned = d_byte_begin_idx & ~15; // loading 16byte chunks
+    const uint4* d_byte = (uint4*)(d_data_compressed + d_byte_begin_idx_aligned);
+    
+    // Get number of remaining global memory byte chunks (not to read bytes out of buffer)
+    const unsigned int d_byte_end_idx_aligned = (d_byte_begin_idx + segment->data_compressed_size + 15) & ~15;
+    unsigned int d_byte_chunk_count = (d_byte_end_idx_aligned - d_byte_begin_idx_aligned) / 16;
+    
+    // Load first 2 chunks of compressed data into the shared memory buffer and remember index of first code byte (skipping bytes read due to alignment)
+    s_byte[0] = d_byte[0];
+    s_byte[1] = d_byte[1];
+    d_byte += 2;
+    d_byte_chunk_count = max(d_byte_chunk_count, 2) - 2;
+    unsigned int s_byte_idx = d_byte_begin_idx - d_byte_begin_idx_aligned;
+    
+    // bits loaded into the register and their count
+    unsigned int r_bit_count = 0;
+    unsigned int r_bit = 0; // LSB-aligned
     
     // Non-interleaving mode
     if ( comp_count == 1 ) {

            }
            
            // Encode 8x8 block
            if ( gpujpeg_huffman_gpu_decoder_decode_block(get_buff, get_bits, component_dc, data_compressed, data_compressed_size, block, d_table_dc, d_table_ac) != 0 )
                break;
        } 
    }

 * @return int
 */
__device__ inline int
gpujpeg_huffman_gpu_decoder_decode_special_decode(
                struct gpujpeg_table_huffman_decoder* table, int min_bits, unsigned int & r_bit,
                unsigned int & r_bit_count, uint4 * const s_byte, 
                unsigned int & s_byte_idx, const uint4 * & d_byte, unsigned int & d_byte_chunk_count)
{
    // HUFF_DECODE has determined that the code is at least min_bits
    // bits long, so fetch that many bits in one swoop.
    int code = gpujpeg_huffman_gpu_decoder_get_bits(min_bits, r_bit, r_bit_count, s_byte, s_byte_idx, d_byte, d_byte_chunk_count);

    // Collect the rest of the Huffman code one bit at a time.
    // This is per Figure F.16 in the JPEG spec.
    int l = min_bits;
    while ( code > table->maxcode[l] ) {
        code <<= 1;
        code |= gpujpeg_huffman_gpu_decoder_get_bits(1, r_bit, r_bit_count, s_byte, s_byte_idx, d_byte, d_byte_chunk_count);
        l++;
    }

