  * @return None
  */
 __device__ void
-gpujpeg_dct_gpu_kernel_inplace(int16_t* SrcDst, int Stride)
+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],

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

    // Determine position into shared buffer
    int16_t* 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);


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

