 
 			// Loop through each pair of labels
 			T minCost = grid->getLabelingCost(result);
-			for (int alpha = 1; alpha < grid->getNumLabels(); alpha++) {
-				for (int beta = 0; beta < alpha; beta++) {
+			for (int alpha = 0; alpha < grid->getNumLabels(); alpha++) {
+				printf("%lf,%d\n", TIMER_END, grid->getLabelingCost(result));
+			
+				for (int beta = alpha + 1; beta < grid->getNumLabels(); beta++) {
 					// Now, construct the graph with alpha and beta as terminal nodes
 					gpuGraphCutConstructGraphKernel<<<gridSize, blockSize>>>(dGrid, dResult, dHeight, dExcess, dResidue, alpha, beta);
 
 					// We need to keep this data synced between CPU and GPU
 					// every time we perform global relabeling
 					gpu.copyRawMat(&height, dHeight);
-					gpu.copyRawMat(&excess, dExcess);
 					gpu.copyRawMat(&residue, dResidue);
 
 					// Run push-relabel until convergence
 					int iters = 0;
 					while (true) {
+						// Mark each tile as active or inactive
+						gpu.clearMat(dActive, isActive.getTotalSize());
+						gpuFindActiveTiles<<<gridSize, blockSize>>>(dGrid, dHeight, dExcess, dActive);
+						gpu.copyRawMat(&isActive, dActive);
+						
+						// Build a list of active tiles
+						if (buildActiveTileList(&isActive, activeList) == 0) break;
+						
 						// Global relabeling optimization on the CPU
-						if (graphCutGlobalRelabel(grid, &height, &excess, &residue)) {
-							// No more active nodes, we're done
-							break;
-						}
+						graphCutGlobalRelabel(grid, &height, &residue);
 						
 						// Copy relabeled height data back to the GPU
 						gpu.copyRawMat(dHeight, &height);
 
 						// Run normal push and relabel on the GPU
 						for (int i = 0; i < GINF_GRAPHCUT_NUM_ITERS_PER_GLOBAL_RELABEL; i++) {
-							gpuGraphCutPushKernel<<<gridSize, blockSize>>>(iters & 1, dGrid, dHeight, dExcess, dResidue);
-							cudaThreadSynchronize();
+							gpuGraphCutPushKernel<<<gridSize, blockSize>>>(dGrid, dHeight, dExcess, dResidue);
+							gpuGraphCutRelabelKernel<<<gridSize, blockSize>>>(dGrid, dHeight, dExcess, dResidue);
 							
-							gpuGraphCutRelabelKernel<<<gridSize, blockSize>>>(iters & 1, dGrid, dHeight, dExcess, dResidue);
-							cudaThreadSynchronize();
 							iters++;
-							
 						}
 						
 						// Copy graph data back to CPU for global relabeling
 						gpu.copyRawMat(&height, dHeight);
-						gpu.copyRawMat(&excess, dExcess);
 						gpu.copyRawMat(&residue, dResidue);
 					}
 

		MatDim dim(grid->getWidth(), grid->getHeight());

		// Check if (x, y) is in the alpha-beta partition
		if (dim.isValid(x, y) && f[dim.idx(x, y)] == alpha || f[dim.idx(x, y)] == beta) {
			int alphaCost = 0, betaCost = 0;
			for (int d = 0; d < 4; d++) {
				int nx = x + dDirX[d], ny = y + dDirY[d];
				if (dim.isValid(nx, ny)) {
					// Check if this neighbour (nx, ny) is in the alpha-beta partition
					if (f[dim.idx(nx, ny)] == alpha || f[dim.idx(nx, ny)] == beta) {
						// Capacity of this edge is V(alpha, beta)
						capacity[dim.idx(x, y, d)] = grid->getSmoothnessCost(alpha, beta);
					} else {
						// There is no edge between these two nodes
						capacity[dim.idx(x, y, d)] = 0;

						alphaCost += grid->getSmoothnessCost(alpha, f[dim.idx(nx, ny)]);
						betaCost += grid->getSmoothnessCost(beta, f[dim.idx(nx, ny)]);
					}
				} else {
					// No edge here
					capacity[dim.idx(x, y, d)] = 0;
				}
			}

			// Calculate the capacity of t-links
			capacity[dim.idx(x, y, GINF_DIR_SOURCE)] = grid->getDataCost(x, y, alpha) + alphaCost;
			capacity[dim.idx(x, y, GINF_DIR_SINK)] = grid->getDataCost(x, y, beta) + betaCost;
			
			// Saturate t-link from source
			excess[dim.idx(x, y)] = capacity[dim.idx(x, y, GINF_DIR_SOURCE)];
		
			// Reset height to 0
			height[dim.idx(x, y)] = 0;
		}
	}

	template <typename T>

		T *dExcess = gpu.createRawMat(&excess),
		  *dResidue = gpu.createRawMat(&residue);
		int *dResult = gpu.createRawMat(result),
		  *dHeight = gpu.createRawMat(&height),
		  *dActive = gpu.createRawMat(&isActive);
		
		GpuGrid<T> *dGrid = gpu.createGrid(grid);
		
		// Loop until there are no more label changes
		bool success = true;
		for (int t = 0; t < numIters && success; t++) {
