 	free(temp);
 	free(temp_old);
 	
+	free(bl_norm_L2);
+	
 	cudaDeviceReset();
 	
 	return 0;

	
	// number of cells in x and y directions
	// including unused boundary cells
	uint num_rows = 128 + 2;
	uint num_cols = 128 + 2;
	uint size_temp = num_rows * num_cols;
	uint size = (num_rows - 2) * (num_cols - 2);
	


__global__ void red_kernel (uint rowmax, uint colmax, const Real * aP,
														const Real * aW, const Real * aE, const Real * aS,
														const Real * aN, const Real * b, Real * temp, Real * bl_norm_L2)
{	
	uint row = 1 + (blockIdx.y * blockDim.y) + threadIdx.y;
	uint col = 1 + (blockIdx.x * blockDim.x) + threadIdx.x;

	// store residual for block
	__shared__ Real res_cache[block_size];
	
	res_cache[threadIdx.y] = 0.0;
		
	// only red cell if even
	if ((row + col) % 2 == 0) {
