     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

    in += source_stride;
    dct_t src7 = *in;
    
    
    
    
    
        // destination pointer into shared transpose buffer (each thread saves one column)
    volatile dct_t * const s_dest = s_transposition + dct_idx;
    
    dct(src0, src1, src2, src3, src4, src5, src6, src7,

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

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