 	if (dma0.owns_this_thread())
 	{
 		// Get the pointers to the origin
-		//ELMT_TYPE *base_ptr = &(idata[ALIGN_OFFSET + RADIUS + RADIUS*pitch]);
-		//dma0.execute_dma(base_ptr, &(buffer[ALIGN_OFFSET + RADIUS + RADIUS*(dimx+2*RADIUS)]));
+#ifndef CUDADMA_DEBUG_ON
+		ELMT_TYPE *base_ptr = &(idata[ALIGN_OFFSET + RADIUS + RADIUS*pitch]);
+		dma0.execute_dma(base_ptr, &(buffer[ALIGN_OFFSET + RADIUS + RADIUS*(dimx+2*RADIUS)]));
+#else
 		dma0.wait_for_dma_start();
 		dma0.finish_async_dma();
+#endif
 	}
 	else
 	{

__global__ void __launch_bounds__(1024,1)
dma_ld_test ( ELMT_TYPE * idata, ELMT_TYPE * odata, int dimx, int dimy, int pitch, int buffer_size /*number of ELMT_TYPE*/, int num_compute_threads, int num_dma_threads_per_ld)
{
	extern __shared__ ELMT_TYPE buffer[];	

	cudaDMAHalo<ELMT_TYPE,RADIUS,CORNERS,ALIGNMENT>
	  dma0 (1, num_dma_threads_per_ld, num_compute_threads,

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;
