 	void gpuDecodeHbp(Grid<T> *grid, int numLevels, int numItersPerLevel, Matrix<int> *result) {
 		Gpu gpu;
 		
-		// Create matrices to store data costs at each level
-		Matrix<T> **data = new Matrix<T>* [numLevels];
+		// We need to keep track of the widths and heights of each level
+		int *widths = new int[numLevels], *heights = new int[numLevels];
 		
 		// Allocate memory on the device
 		int *dResult = gpu.createRawMat(result);

		}
	}
	
	/*
	template <typename T>
	__global__ void gpuBpCalcDataCost(GpuGrid<T> *grid, int w, int h, int pw, int ph, T *data, T *prev) {
		GINF_DECL_X; GINF_DECL_Y;
		MatDim dim(w, h, grid->getNumLabels());
		MatDim dimPrev(pw, ph, grid->getNumLabels());
		
		if (dim.isValid(x / 2, y / 2)) {
			for (int i = 0; i < grid->getNumLabels(); i++) {
				data[dim.idx(x / 2, y / 2, i)] += prev[dimPrev.idx(x, y, i)];
			}
		}
	}*/
	
	template <typename T>
	__global__ void gpuBpInitMsgs(GpuGrid<T> *grid, int w, int h, int pw, int ph, T *msgs, T *prev) {
		GINF_DECL_X; GINF_DECL_Y;	
		MatDim dim(w, h, GINF_NUM_DIR);
		MatDim dimPrev(pw, ph, GINF_NUM_DIR);
		
		if (dim.isValidStrict(x, y)) {
			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)];
				}
			}
		}

		GpuGrid<T> *dGrid = gpu.createGrid(grid);
		
		// At the finest level, the data costs correspond to the original problem
		widths[0] = grid->getWidth();
		heights[0] = grid->getHeight();

		// Copy the original data costs to the device
		dData[0] = gpu.createRawMat(grid->getDataCosts());

		// Prepare block sizes
		dim3 blockSize(GINF_BP_BLOCK_WIDTH, GINF_BP_BLOCK_HEIGHT);
