     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;
     

        comp_count = coder->param_image.comp_count;
    assert(comp_count >= 1 && comp_count <= GPUJPEG_MAX_COMPONENT_COUNT);
    
    // Configure more Shared memory for both kernels
    cudaFuncSetCacheConfig(gpujpeg_huffman_decoder_table_kernel, cudaFuncCachePreferShared);
    cudaFuncSetCacheConfig(gpujpeg_huffman_decoder_decode_kernel, cudaFuncCachePreferShared);
    
    // Setup GPU tables (one thread for each of 65536 entries)
    gpujpeg_huffman_decoder_table_kernel<<<256, 256>>>(

    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,
