 	bool success04_2 = true;
 	bool success04_3 = true;
 
-	run_experiment<16,0,float,4,true>(32,8,64,1);
-
 	// An 8192*sizeof(float) element is almost all of shared memory (can only fit one)
 	// We probably want to test up to 32 elements per stride if we're going to have up to 16 dma warps
-	//success16_0 = success16_0 && run_all_experiments<16,0>(8192,32,16);
-	//success08_0 = success08_0 && run_all_experiments<8,0>(8192,32,16);
-	//success08_2 = success08_2 && run_all_experiments<8,2>(8192,32,16);
-	//success04_0 = success04_0 && run_all_experiments<4,0>(8192,32,16);
-	//success04_1 = success04_1 && run_all_experiments<4,1>(8192,32,16);
-	//success04_2 = success04_2 && run_all_experiments<4,2>(8192,32,16);
-	//success04_3 = success04_3 && run_all_experiments<4,3>(8192,32,16);
+	//success16_0 = success16_0 && run_all_experiments<16,0>(16,64,64);
+	//success08_0 = success08_0 && run_all_experiments<8,0>(16,64,64);
+	//success08_2 = success08_2 && run_all_experiments<8,1>(16,64,64);
+	success04_0 = success04_0 && run_all_experiments<4,0>(16,64,64);
+	success04_1 = success04_1 && run_all_experiments<4,1>(16,64,64);
+	success04_2 = success04_2 && run_all_experiments<4,2>(16,64,64);
+	success04_3 = success04_3 && run_all_experiments<4,3>(16,64,64);
+	
+	//success16_0 = success16_0 && run_experiment<4,0,float4,5,false>(8,8,32,2);
 
 	fprintf(stdout,"\nResults:\n");
 	fprintf(stdout,"\tAlignment16-Offset0: %s\n",(success16_0?"SUCCESS":"FAILURE"));

#include "cuda.h"
#include "cuda_runtime.h"

#define CUDADMA_DEBUG_ON
#include "../../include/cudaDMA.h"

#define WARP_SIZE 32


	// Allocate the inpute data
	int input_size = (pitch*(dimy+2*RADIUS)+ALIGN_OFFSET);
	ELMT_TYPE *h_idata = (ELMT_TYPE*)malloc(input_size*sizeof(ELMT_TYPE));
	for (int i=0; i<input_size; i++)
		initialize<ELMT_TYPE>(h_idata[i],(rand()%10));
	// Allocate device memory and copy down
	ELMT_TYPE *d_idata;
	CUDA_SAFE_CALL( cudaMalloc( (void**)&d_idata, input_size*sizeof(ELMT_TYPE)));
	CUDA_SAFE_CALL( cudaMemcpy( d_idata, h_idata, input_size*sizeof(ELMT_TYPE), cudaMemcpyHostToDevice));

	// Allocate the output size
	ELMT_TYPE *h_odata = (ELMT_TYPE*)malloc(shared_buffer_size*sizeof(ELMT_TYPE));
	for (int i=0; i<shared_buffer_size; i++)
		initialize<ELMT_TYPE>(h_odata[i],0.0f);
	// Allocate device memory and copy down
	ELMT_TYPE *d_odata;
	CUDA_SAFE_CALL( cudaMalloc( (void**)&d_odata, shared_buffer_size*sizeof(ELMT_TYPE)));
	CUDA_SAFE_CALL( cudaMemcpy( d_odata, h_odata, shared_buffer_size*sizeof(ELMT_TYPE), cudaMemcpyHostToDevice));

	int num_compute_warps = 1;
	int total_threads = (num_compute_warps + dma_warps)*WARP_SIZE;
