     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;
+            float coefficient = (int16_t)(source[i * source_stride]);
+            coefficient -= 128.0f;
             block_ptr[i * GPUJPEG_DCT_THREAD_BLOCK_STRIDE] = coefficient;
         }
     }


    // 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();


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