 			}
 		
 			// Push to neighbours
-			for (int d = 0; d < 4; d++) {
+			for (int d = 0; d < GINF_NUM_DIR; d++) {
 				if (residue[dim.idx(x, y, d)] && height[dim.idx(x, y)] == height[dim.idx(x + dDirX[d], y + dDirY[d])] + 1) {
-					gpuGraphCutPushToPixel(iter, grid, excess, residue, d);
-				}
-			}
-		}
-	}
-	
-	template <typename T>
-	__global__ void gpuGraphCutRelabelKernel(bool iter, GpuGrid<T> *grid, int *height, T *excess, T *residue) {
-		GINF_DECL_ALT_X(iter); GINF_DECL_Y;
+					gpuGraphCutPushToPixel(grid, excess, residue, d);
+				}
+			}
+		}
+	}
+	
+	template <typename T>
+	__global__ void gpuGraphCutRelabelKernel(GpuGrid<T> *grid, int *height, T *excess, T *residue) {
+		GINF_DECL_X; GINF_DECL_Y;
 		MatDim dim(grid->getWidth(), grid->getHeight());
 
 		if (dim.isValid(x, y) && excess[dim.idx(x, y)] && height[dim.idx(x, y)] < grid->getNumNodes() + 2) {

	}
	
	template <typename T>
	__global__ void gpuGraphCutPushKernel(bool iter, GpuGrid<T> *grid, int *height, T *excess, T *residue) {
		GINF_DECL_ALT_X(iter); GINF_DECL_Y;
		MatDim dim(grid->getWidth(), grid->getHeight());

		if (dim.isValid(x, y) && excess[dim.idx(x, y)] && height[dim.idx(x, y)] < grid->getNumNodes() + 2) {

		// Start with the intial labeling
		result->copyFrom(initial);
		
		// Prepare block sizes
		dim3 blockSize(GINF_GRAPHCUT_BLOCK_WIDTH, GINF_GRAPHCUT_BLOCK_HEIGHT);
		dim3 gridSize(GINF_DIM(grid->getWidth(), blockSize.x), GINF_DIM(grid->getHeight(), blockSize.y));
		
		// Boolean matrix indicating whether a node is reachable from the source in the residue graph
		Matrix<int> reachable(2, grid->getWidth(), grid->getHeight());
		
		// Boolean matrix indicate whether a block contains an active node or not
		Matrix<int> isActive(2, gridSize.x, gridSize.y);
		
		// List of active blocks (xy coordinates)
		int2 activeList[gridSize.x * gridSize.y];

		// Create a temp matrix to store new labelings
		Matrix<int> newResult(2, grid->getWidth(), grid->getHeight());
