     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]);
     }
 }
 

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

    int thread_x_permutated = (thread_x & 0xFFFFFFE0) | (((thread_x << 1) | ((thread_x >> 4) & 0x1)) & 0x1F);

    // Determine position into shared buffer
    float* block_ptr = block + IMAD(thread_y, GPUJPEG_DCT_THREAD_BLOCK_STRIDE, thread_x);

    // Determine position in source buffer and apply it
    int source_x = IMAD(block_x, GPUJPEG_BLOCK_SIZE, threadIdx.x);
