  *  @param[out] g_ybar All \f$P_{m,n}(\bar{Y})\f$
  *  @param[out] g_vhat All \f$P^T_{m,n}(\hat{V})\f$
  */
-__global__ void algSAT_stage1(const int c_width, const int c_height, const float *g_in, float *g_ybar,
-    float *g_vhat )
-{
-
-  const int tx = threadIdx.x, ty = threadIdx.y,
-        bx = blockIdx.x, by = blockIdx.y, col = bx*WS+tx, row0 = by*WS;
-
-  __shared__ float s_block[ WS ][ WS+1 ];
-
-  float (*bdata)[WS+1] = (float (*)[WS+1]) &s_block[ty][tx];
-
-  g_in += (row0+ty)*c_width+col;
-  g_ybar += by*c_width+col;
-  g_vhat += bx*c_height+row0+tx;
-
-#pragma unroll
-  for (int i = 0; i < WS-(WS%SOW); i+=SOW) {
+__global__ void algSAT_stage1(const int c_width, const int c_height,
+                              const float *g_in, float *g_ybar, float *g_vhat)
+{
+
+  const int tx = threadIdx.x, ty = threadIdx.y, bx = blockIdx.x,
+            by = blockIdx.y, col = bx * WS + tx, row0 = by * WS;
+
+  __shared__ float s_block[WS][WS + 1];
+
+  float(*bdata)[WS + 1] = (float(*)[WS + 1]) & s_block[ty][tx];
+
+  g_in += (row0 + ty) * c_width + col;
+  g_ybar += by * c_width + col;
+  g_vhat += bx * c_height + row0 + tx;
+
+#pragma unroll
+  for (int i = 0; i < WS - (WS % SOW); i += SOW)
+  {
     **bdata = *g_in;
     bdata += SOW;
     g_in += SOW * c_width;
   }
-  if( ty < WS%SOW ) {
+  if (ty < WS % SOW)
+  {
     **bdata = *g_in;
   }
 
   __syncthreads();
 
-  if( ty == 0 ) {
-
-    {   // calculate ybar -----------------------
-      float (*bdata)[WS+1] = (float (*)[WS+1]) &s_block[0][tx];
+  if (ty == 0)
+  {
+
+    {  // calculate ybar -----------------------
+      float(*bdata)[WS + 1] = (float(*)[WS + 1]) & s_block[0][tx];
 
       float prev = **bdata;
       ++bdata;


    g_ybar += c_width;
    y = *g_ybar += y;

  }
}



texture<float, 2, cudaReadModeElementType> textureReadDepth;

inline unsigned int divUp(unsigned int a, unsigned int b)
{
  return (a % b != 0) ? (a / b + 1) : (a / b);
}

#define WS 32   // Warp size (defines b x b block size where b = WS)
#define HWS 16  // Half Warp Size
#define DW 8    // Default number of warps (computational block height)
#define CHW 7   // Carry-heavy number of warps
                // (computational block height for some kernels)
#define OW   6  // Optimized number of warps (computational block height for
                // some kernels)
#define DNB  6  // Default number of blocks per SM (minimum blocks per SM
                // launch bounds)
#define ONB 5   // Optimized number of blocks per SM (minimum blocks per SM
                // for some kernels)
#define MTS 192 // Maximum number of threads per block with 8 blocks per SM
#define MBO 8   // Maximum number of blocks per SM using optimize or maximum
                // warps
#define CHB 7   // Carry-heavy number of blocks per SM using default number of
                // warps
#define MW 6    // Maximum number of warps per block with 8 blocks per SM (with
                // all warps computing)
#define SOW 5   // Dual-scheduler optimized number of warps per block (with 
                // 8 blocks per SMand to use the dual scheduler with 1
                // computing warp)
#define MBH 3   // Maximum number of blocks per SM using half-warp size

/**
 *  @ingroup gpu
