 			}
 			
 			// Write the result
-			result[dim.idx(x, y)] = minLabel;
+			result[dim(x, y)] = minLabel;
 		}
 	}
 	

		Gpu gpu;
		
		// Create matrices to store data costs at each level
            	Matrix<T> **data = new Matrix<T>* [numLevels];
		
		// Allocate memory on the device
		int *dResult = gpu.createRawMat(result);


namespace ginf {
	template <typename T>
	__global__ void gpuBpCollectKernel(GpuGrid<T> *grid, int w, int h, T *msgs, T *from, T *data) {
		GINF_DECL_X; GINF_DECL_Y;
	
		MatDim dim(w, h);
		MatDim dimMsgs(w, h, GINF_NUM_DIR);

		if (dim.isValidStrict(x, y)) {
			// For each state
			for (int i = 0; i < grid->getNumLabels(); ++i) {
				from[dim(x, y, i)] = data[dim(x, y, i)];
				
				// For each neighbour in each direction, sum up the messages
				for (int d = 0; d < GINF_NUM_DIR; d++) {
					from[dim(x, y, i)] += msgs[dimMsgs(x, y, d, i)];
				}
			}
		}
	}
	
	template <typename T>
	__global__ void gpuBpKernel(int iter, GpuGrid<T> *grid, int w, int h, T *msgs, T *from, T *data, int d) {
		GINF_DECL_ALT_X(iter); GINF_DECL_Y;
		
		MatDim dim(w, h, GINF_NUM_DIR);
		MatDim dimMsgs(w, h, GINF_NUM_DIR);

		if (dim.isValidStrict(x, y)) {
			// Create a vector to store the message. This is stored in local memory, and so writing to it
			// is coalesced.
			T out[GINF_MAX_NUM_LABELS];
			
			// Find the linear scaling constant
			T scale = grid->getSmoothnessCost(0, 1);
			
			// Calculate message vector in the forward direction
			T minh = out[0] = from[dim(x, y, 0)] - msgs[dimMsgs(x, y, d, 0)];
			for (int i = 1; i < grid->getNumLabels(); i++) {
				out[i] = GINF_MIN(from[dim(x, y, i)] - msgs[dimMsgs(x, y, d, i)], out[i - 1] + scale);
				minh = GINF_MIN(minh, out[i]);
			}
				
			// Truncate the messages using truncate constant
			minh += grid->getSmoothnessCost(0, grid->getNumLabels() - 1);
				
			// Compute the final out vector, and sum up all the elements to perform normalization
			T val = 0;
			for (int i = grid->getNumLabels() - 2; i >= 0; i--) {
				out[i] = GINF_MIN(GINF_MIN(out[i], out[i + 1] + scale), minh);
				val += out[i];
			}

			// Find the normalizing constant
			val /= grid->getNumLabels();

			// Update the message
			int nx = x + dDirX[d], ny = y + dDirY[d];
			for (int i = 0; i < grid->getNumLabels(); i++) {
				// Update using the normalized message values
				msgs[dim(nx, ny, GINF_OPP_DIR(d), i)] = out[i] - val;
			}
		}
	}
