 			// We gotta make sure that we don't get too 'coarse'
 			if (w < 1 || h < 1) return;
 
-			// The data cost of each node on the current level uses the sum of
+			// The data cost of each node GINF_BP_BLOCK_WIDTHon the current level uses the sum of
 			// the data costs of four nodes below this level
 			data[t] = new Matrix<T>(3, w, h, grid->getNumLabels());
 			for (int y = 0; y < data[t - 1]->getSize(1); y++) {

	__global__ void gpuBpGetBeliefKernel(GpuGrid<T> *grid, T *msgs, T *from, int *result) {
		GINF_DECL_X; GINF_DECL_Y;
		MatDim dim(grid->getWidth(), grid->getHeight());
		MatDim dimMsgs(grid->getWidth(), grid->getHeight(), GINF_NUM_DIR);
		
		if (dim.isValidStrict(x, y)) {
			// Collect msgs from neighbours
			gpuBpCollect(grid, grid->getWidth(), grid->getHeight(), msgs, from, x, y);
		
			// Pick the label that minimizes the cost
			T minCost = grid->getDataCost(x, y, 0) + from[dim.idx(x, y, 0)];
			int minLabel = 0;
			
			for (int i = 1; i < grid->getNumLabels(); i++) {
				T cost = grid->getDataCost(x, y, i) + from[dim.idx(x, y, i)];
				if (cost < minCost) {
					minCost = cost;
					minLabel = i;

			for (int d = 0; d < GINF_NUM_DIR; d++) {
				for (int i = 0; i < grid->getNumLabels(); i++) {
					// Initialize message with node on the above level
					msgs[dim(x, y, d, i)] = prev[dimPrev(x / 2, y / 2, d, i)];
				}
			}
		}
