     out6 = even3 - even_diff * 0.707106781f;
     out7 = odd_diff1 - odd_diff4;
 }
+
+/** Constant memory copy of transposed quantization table pre-divided with DCT output weights. */
+__constant__ float gpujpeg_dct_gpu_quantization_table_const[64];
 
 /**
  * Performs 8x8 block-wise Forward Discrete Cosine Transform of the given

            d_quantization_table
        );
        cudaThreadSynchronize();
        gpujpeg_cuda_check_error("Forward Integer DCT failed");
    }
}


        // Get quantization table
        enum gpujpeg_component_type type = (comp == 0) ? GPUJPEG_COMPONENT_LUMINANCE : GPUJPEG_COMPONENT_CHROMINANCE;
        const float* const d_quantization_table = encoder->table_quantization[type].d_table_forward;
        
        // copy the quantization table into constant memory for devices of CC < 2.0
        if( encoder->coder.cuda_cc_major < 2 ) {
            cudaMemcpyToSymbol(
                gpujpeg_dct_gpu_quantization_table_const,
                d_quantization_table,
                sizeof(gpujpeg_dct_gpu_quantization_table_const),
                0,
                cudaMemcpyDeviceToDevice
            );
            cudaThreadSynchronize();
            gpujpeg_cuda_check_error("Quantization table memcpy failed");
        }

        int roi_width = component->data_width;
        int roi_height = component->data_height;
