     // index of row/column processed by this thread within its 8x8 block
     const int dct_idx = threadIdx.x & 7;
     
-    
-    
     // Data type of transformed coefficients
     typedef float dct_t;
     

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

    V8[3] = sh3.hInt;
}



template <typename T>
__device__ static inline void
dct(const T in0, const T in1, const T in2, const T in3, const T in4, const T in5, const T in6, const T in7,
    volatile T & out0, volatile T & out1, volatile T & out2, volatile T & out3, volatile T & out4, volatile T & out5, volatile T & out6, volatile T & out7,
    const float level_shift = 0.0f)
{
    /* Load data into workspace */
    const float tmp0 = in0 + in7;
    const float tmp7 = in0 - in7;
    const float tmp1 = in1 + in6;
    const float tmp6 = in1 - in6;
    const float tmp2 = in2 + in5;
    const float tmp5 = in2 - in5;
    const float tmp3 = in3 + in4;
    const float tmp4 = in3 - in4;

    {
        /* Even part */

        const float tmp10 = tmp0 + tmp3;        /* phase 2 */
        const float tmp13 = tmp0 - tmp3;
        const float tmp11 = tmp1 + tmp2;
        const float tmp12 = tmp1 - tmp2;

        /* Apply unsigned->signed conversion */
        out0 = tmp10 + tmp11 + level_shift; /* phase 3 */
        out4 = tmp10 - tmp11;

        const float z1 = (tmp12 + tmp13) * 0.707106781f; /* c4 */
        out2 = tmp13 + z1;    /* phase 5 */
        out6 = tmp13 - z1;
    }
    
    /* Odd part */

    const float tmp10 = tmp4 + tmp5;        /* phase 2 */
    const float tmp11 = tmp5 + tmp6;
    const float tmp12 = tmp6 + tmp7;

    /* The rotator is modified from fig 4-8 to avoid extra negations. */
    const float z5 = (tmp10 - tmp12) * 0.382683433f; /* c6 */
    const float z2 = 0.541196100f * tmp10 + z5; /* c2-c6 */
    const float z4 = 1.306562965f * tmp12 + z5; /* c2+c6 */
    const float z3 = tmp11 * 0.707106781f; /* c4 */

    const float z11 = tmp7 + z3;            /* phase 5 */
    const float z13 = tmp7 - z3;

    out5 = z13 + z2;      /* phase 6 */
    out3 = z13 - z2;
    out1 = z11 + z4;
    out7 = z11 - z4;
}

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