     struct gpujpeg_table_huffman_decoder* d_table_cbcr_dc,
     struct gpujpeg_table_huffman_decoder* d_table_cbcr_ac
 ) {
-    int segment_index = blockIdx.x * blockDim.x + threadIdx.x;
+    int segment_index = blockIdx.x * THREADS_PER_TBLOCK + threadIdx.x;
     if ( segment_index >= segment_count )
         return;
     

 * 
 * @return void
 */
__global__ void
gpujpeg_huffman_decoder_decode_kernel(
    struct gpujpeg_component* d_component,

    cudaThreadSynchronize();
    gpujpeg_cuda_check_error("Huffman decoder table setup failed");
    
    // Run decoding kernel
    dim3 thread(THREADS_PER_TBLOCK);
    dim3 grid(gpujpeg_div_and_round_up(decoder->segment_count, THREADS_PER_TBLOCK));
    gpujpeg_huffman_decoder_decode_kernel<THREADS_PER_TBLOCK><<<grid, thread>>>(
        coder->d_component, 
        coder->d_segment, 
        comp_count,
