     cudaThreadSynchronize();
     gpujpeg_cuda_check_error("Huffman decoder table setup failed");
     
-    // Run kernel
-    dim3 thread(32);
-    dim3 grid(gpujpeg_div_and_round_up(decoder->segment_count, thread.x));
-    gpujpeg_huffman_decoder_decode_kernel<<<grid, thread>>>(
+    // 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,

        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>>>(

 * 
 * @return void
 */
template <int THREADS_PER_TBLOCK>
__global__ void
gpujpeg_huffman_decoder_decode_kernel(
    struct gpujpeg_component* d_component,
