 
 //Almost the same as naive scan1Inclusive, but doesn't need __syncthreads()
 //assuming size <= WARP_SIZE and size is power of 2
-template <class T>
-inline __device__ T warpScanInclusive(T idata, volatile T *s_Data)
-{
-    Ncv32u pos = 2 * threadIdx.x - (threadIdx.x & (K_WARP_SIZE - 1));
+//template <class T>
+//inline __device__ T warpScanInclusive(T idata, volatile T *s_Data)
+//{
+//    Ncv32u pos = 2 * threadIdx.x - (threadIdx.x & (K_WARP_SIZE - 1));
+//    s_Data[pos] = 0;
+//    pos += K_WARP_SIZE;
+//    s_Data[pos] = idata;
+//
+//    s_Data[pos] += s_Data[pos - 1];
+//    s_Data[pos] += s_Data[pos - 2];
+//    s_Data[pos] += s_Data[pos - 4];
+//    s_Data[pos] += s_Data[pos - 8];
+//    s_Data[pos] += s_Data[pos - 16];
+//
+//    return s_Data[pos];
+//}
+
+
+//template <class T>
+//inline __device__ T warpScanExclusive(T idata, volatile T *s_Data)
+//{
+//    return warpScanInclusive(idata, s_Data) - idata;
+//}
+//
+//
+//template <class T, Ncv32u tiNumScanThreads>
+//inline __device__ T blockScanInclusive(T idata, volatile T *s_Data)
+//{
+//    if (tiNumScanThreads > K_WARP_SIZE)
+//    {
+//        //Bottom-level inclusive warp scan
+//        T warpResult = warpScanInclusive(idata, s_Data);
+//
+//        //Save top elements of each warp for exclusive warp scan
+//        //sync to wait for warp scans to complete (because s_Data is being overwritten)
+//        __syncthreads();
+//        if( (threadIdx.x & (K_WARP_SIZE - 1)) == (K_WARP_SIZE - 1) )
+//        {
+//            s_Data[threadIdx.x >> K_LOG2_WARP_SIZE] = warpResult;
+//        }
+//
+//        //wait for warp scans to complete
+//        __syncthreads();
+//
+//        if( threadIdx.x < (tiNumScanThreads / K_WARP_SIZE) )
+//        {
+//            //grab top warp elements
+//            T val = s_Data[threadIdx.x];
+//            //calculate exclusive scan and write back to shared memory
+//            s_Data[threadIdx.x] = warpScanExclusive(val, s_Data);
+//        }
+//
+//        //return updated warp scans with exclusive scan results
+//        __syncthreads();
+//        return warpResult + s_Data[threadIdx.x >> K_LOG2_WARP_SIZE];
+//    }
+//    else
+//    {
+//        return warpScanInclusive(idata, s_Data);
+//    }
+//}
+
+template <Ncv32u size>
+__device__ Ncv32u warpScanInclusive(Ncv32u idata, volatile Ncv32u* s_Data)
+{
+    Ncv32u pos = 2 * threadIdx.x - (threadIdx.x & (size - 1));
     s_Data[pos] = 0;
-    pos += K_WARP_SIZE;
+    pos += size;
     s_Data[pos] = idata;
 
-    s_Data[pos] += s_Data[pos - 1];
-    s_Data[pos] += s_Data[pos - 2];
-    s_Data[pos] += s_Data[pos - 4];
-    s_Data[pos] += s_Data[pos - 8];
-    s_Data[pos] += s_Data[pos - 16];
+    for(Ncv32u offset = 1; offset < size; offset <<= 1)
+        s_Data[pos] += s_Data[pos - offset];
 
     return s_Data[pos];
 }
 
-
-template <class T>
-inline __device__ T warpScanExclusive(T idata, volatile T *s_Data)
-{
-    return warpScanInclusive(idata, s_Data) - idata;
-}
-
-
-template <class T, Ncv32u tiNumScanThreads>
-inline __device__ T blockScanInclusive(T idata, volatile T *s_Data)
-{
-    if (tiNumScanThreads > K_WARP_SIZE)
+template <Ncv32u size>
+__forceinline__ __device__ Ncv32u warpScanExclusive(Ncv32u idata, volatile Ncv32u *s_Data)
+{
+    return warpScanInclusive<size>(idata, s_Data) - idata;
+}
+
+template <Ncv32u size, Ncv32u tiNumScanThreads>
+__device__ Ncv32u scan1Inclusive(Ncv32u idata, volatile Ncv32u *s_Data)
+{
+    if(size > K_WARP_SIZE)
     {
         //Bottom-level inclusive warp scan
-        T warpResult = warpScanInclusive(idata, s_Data);
+        Ncv32u warpResult = warpScanInclusive<K_WARP_SIZE>(idata, s_Data);
 
         //Save top elements of each warp for exclusive warp scan
         //sync to wait for warp scans to complete (because s_Data is being overwritten)
         __syncthreads();
         if( (threadIdx.x & (K_WARP_SIZE - 1)) == (K_WARP_SIZE - 1) )
-        {
             s_Data[threadIdx.x >> K_LOG2_WARP_SIZE] = warpResult;
-        }
 
         //wait for warp scans to complete
         __syncthreads();
-
         if( threadIdx.x < (tiNumScanThreads / K_WARP_SIZE) )
         {
             //grab top warp elements
-            T val = s_Data[threadIdx.x];
-            //calculate exclusive scan and write back to shared memory
-            s_Data[threadIdx.x] = warpScanExclusive(val, s_Data);
+            Ncv32u val = s_Data[threadIdx.x];
+            //calculate exclsive scan and write back to shared memory
+            s_Data[threadIdx.x] = warpScanExclusive<(size >> K_LOG2_WARP_SIZE)>(val, s_Data);
         }
 
         //return updated warp scans with exclusive scan results

    fdata.resize(fsize);
    Ncv32u dataOffset = 0;
    fseek(fp, 0, SEEK_SET);
    ncvAssertReturn(1 == fread(&fdata[0], fsize, 1, fp), NCV_FILE_ERROR);
    fclose(fp);

    //data


    NcvBool bTexCacheCascade = devProp.major < 2;
    NcvBool bTexCacheIImg = true; //this works better even on Fermi so far
    NcvBool bDoAtomicCompaction = false;// devProp.major >= 2 || (devProp.major == 1 && devProp.minor >= 3);

    NCVVector<Ncv32u> *d_ptrNowData = &d_vecPixelMask;
    NCVVector<Ncv32u> *d_ptrNowTmp = &d_vecPixelMaskTmp;
