 template<typename ELMT_TYPE, int RADIUS, bool CORNERS>
 __host__ bool check_halo(ELMT_TYPE *idata, ELMT_TYPE *odata, int dimx, int dimy, int pitch)
 {
-	return false;
+	bool pass = true;
+	// First check the rows
+	{
+		// Check the top rows first
+		for (int i=0; i<RADIUS; i++)
+		{
+			// Compute the start of the row
+			ELMT_TYPE *src_row_start = idata + (i*pitch) + (CORNERS ? 0 : RADIUS);				
+			ELMT_TYPE *dst_row_start = odata + (i*(dimx+2*RADIUS)) + (CORNERS ? 0 : RADIUS);
+			for (int j=0; j<(dimx+(CORNERS ? 2*RADIUS : 0)); j++)
+			{
+				if (!(equal<ELMT_TYPE>(src_row_start[j],dst_row_start[j])))
+				{
+					fprintf(stdout,"Error on top row %d at element %d\n",i,j);
+					pass = false;
+					return pass;
+				}
+			}
+		}
+		// Check the bottom rows next
+		for (int i=0; i<RADIUS; i++)
+		{
+			// Compute the start of the row
+			ELMT_TYPE *src_row_start = idata + (RADIUS+dimy+i)*pitch + (CORNERS ? 0 : RADIUS);
+			ELMT_TYPE *dst_row_start = odata + (RADIUS+dimy+i)*(dimx+2*RADIUS) + (CORNERS ? 0 : RADIUS);
+			for (int j=0; j<(dimx+(CORNERS ? 2*RADIUS : 0)); j++)
+			{
+				if (!(equal<ELMT_TYPE>(src_row_start[j],dst_row_start[j])))
+				{
+					fprintf(stdout,"Error on bottom row %d at element %d\n",i,j);
+					pass = false;
+					return pass;
+				}
+			}
+		}
+	}
+	// Then check the side panels
+	{
+		// Check the left panel
+		for (int i=0; i<dimy; i++)
+		{
+			ELMT_TYPE *src_row_start = idata + (RADIUS+i)*pitch;
+			ELMT_TYPE *dst_row_start = odata + (RADIUS+i)*(dimx+2*RADIUS);
+			for (int j=0; j<RADIUS; j++)
+			{
+				if (!(equal<ELMT_TYPE>(src_row_start[j],dst_row_start[j])))
+				{
+					fprintf(stdout,"Error on left panel at row %d at element %d\n",i,j);
+					pass = false;
+					return pass;
+				}
+			}
+		}
+		// Check the right panel
+		for (int i=0; i<dimy; i++)
+		{
+			ELMT_TYPE *src_row_start = idata + (RADIUS+i)*pitch + RADIUS + dimx;
+			ELMT_TYPE *dst_row_start = odata + (RADIUS+i)*(dimx+2*RADIUS) + RADIUS + dimx;
+			for (int j=0; j<RADIUS; j++)
+			{
+				if (!(equal<ELMT_TYPE>(src_row_start[j],dst_row_start[j])))
+				{
+					fprintf(stdout,"Error on right panel at row %d at element %d\n",i,j);
+					pass = false;
+					return pass;
+				}
+			}
+		}
+	}
+	return pass;
 }
 
 template<typename ELMT_TYPE>

		 num_compute_warps*WARP_SIZE, dma_warps*WARP_SIZE);
	CUDA_SAFE_CALL( cudaThreadSynchronize());

	CUDA_SAFE_CALL( cudaMemcpy (h_odata, d_odata, shared_buffer_size*sizeof(float), cudaMemcpyDeviceToHost));

	// Print out the arrays
	printf("Source array:\n");
	print_2d_array<ELMT_TYPE>(&h_idata[ALIGN_OFFSET],dimy+2*RADIUS,pitch);
	printf("Result array:\n");
	print_2d_array<ELMT_TYPE>(&h_odata[ALIGN_OFFSET],dimy+2*RADIUS,dimx+2*RADIUS);

	// Check the result
	bool pass = check_halo<ELMT_TYPE, RADIUS, CORNERS>(h_idata, h_odata, dimx, dimy, pitch);
	//if (!pass)
	{
		fprintf(stdout,"Result - %s\n",(pass?"SUCCESS":"FAILURE"));
		fflush(stdout);
	}

template<>
__host__ __device__ __forceinline__ float get_val(float4 &t) { return t.x; }

template<typename TYPE>
__host__ __forceinline__ bool equal(TYPE one, TYPE two) { return false; }

template<>
__host__ __forceinline__ bool equal(float one, float two) { return (one==two); }

template<>
__host__ __forceinline__ bool equal(float2 one, float2 two) { return (one.x==two.x)&&(one.y==two.y); }

template<>
__host__ __forceinline__ bool equal(float4 one, float4 two) { return (one.x==two.x)&&(one.y==two.y)&&(one.z==two.z)&&(one.w==two.w); }

// I hate global variables, but whatever
long total_experiments = 0;
