   return err;
 }
 
+// JFC: problem here ncols a non-multiple of 16, and nrows < 32. 
+
 template<class T>
 int maxmini_cols(T *in, T *out, int *outi, int nrows, int ncols, T minv, int dir) {
   int nc1, nc2;
   setinds(ncols, nc1, nc2);
-  dim3 grid(1, nc1, nc2);
+  dim3 grid(nc1, nc2, 1);    
   int ny = min(32, 1+nrows/32);
   dim3 tblock(32, ny, 1);
   __maxmini_cols<T><<<grid,tblock>>>(in, out, outi, nrows, ncols, minv, dir);

  __shared__ int maxi[32];
  T vmax, vtmp;
  int imax, itmp, i, k;
  int bid = blockIdx.y + blockIdx.z * blockDim.y;

  if (bid < ncols) {
    vmax = maxminv;

        outi[bid] = imax;
      }
    }
    __syncthreads();
  }
}

