  */
 template <int WARP_COUNT>
 __global__ void
-gpujpeg_dct_gpu_kernel(int block_count_x, int block_count_y, uint8_t* source, int source_stride,
+gpujpeg_dct_gpu_kernel(int block_count_x, int block_count_y, uint8_t* source, const unsigned int source_stride,
                        int16_t* output, int output_stride)
 {
     // each warp processes 4 8x8 blocks (horizontally neighboring)

    V8[3] = sh3.hInt;
}

/** Quantization table */
#if __CUDA_ARCH__ < 200
__constant__ // quantization table in constant mempory is faster on devices without L2 cache

{
    return (x + 0x1000) >> 13;
}

/**
 * Performs in-place IDCT of vector of 8 elements (used to access columns in shared memory).
