     if ( block_x < block_count_x && block_y < block_count_y ) {
         #pragma unroll
         for(int i = 0; i < GPUJPEG_BLOCK_SIZE; i++)
-            output[i * GPUJPEG_BLOCK_SIZE] = block_ptr[i * GPUJPEG_DCT_THREAD_BLOCK_STRIDE];
+            output[i * GPUJPEG_BLOCK_SIZE] = round(block_ptr[i * GPUJPEG_DCT_THREAD_BLOCK_STRIDE]);
     }
 }
 


    // Quantization
    for(int i = 0; i < GPUJPEG_BLOCK_SIZE; i++) {
        uint16_t quantization = quantization_table[i * GPUJPEG_BLOCK_SIZE + threadIdx.x];
        int coefficient = block_ptr[i * GPUJPEG_DCT_THREAD_BLOCK_STRIDE];

        if ( coefficient < 0 ) {
            coefficient = -coefficient;
            coefficient = (coefficient * quantization + 16384) / 32767;
            coefficient = -coefficient;
        } else {
            coefficient = (coefficient * quantization + 16384) / 32767;
        }

        block_ptr[i * GPUJPEG_DCT_THREAD_BLOCK_STRIDE] = coefficient;
    }
    __syncthreads();


 * @return None
 */
__device__ void
gpujpeg_dct_gpu_kernel_inplace(float* SrcDst, int Stride)
{
    dct(SrcDst[Stride * 0], SrcDst[Stride * 1], SrcDst[Stride * 2], SrcDst[Stride * 3],
        SrcDst[Stride * 4], SrcDst[Stride * 5], SrcDst[Stride * 6], SrcDst[Stride * 7],
