 	int num_compute_warps = 1;
 	int total_threads = (num_compute_warps + dma_warps)*WARP_SIZE;
 
-	dma_ld_test<ALIGNMENT,ALIGN_OFFSET>
-		<<<1,total_threads,shared_buffer_size*sizeof(float),0>>>
-		(d_idata, d_odata, element_size, element_count, src_stride,
-		 dst_stride, shared_buffer_size, num_compute_warps*WARP_SIZE,
-		 dma_warps*WARP_SIZE);
+	int elem_bytes = element_size*sizeof(float);
+	if (elem_bytes <= 64)
+	{
+		dma_ld_test<ALIGNMENT,ALIGN_OFFSET,64>
+			<<<1,total_threads,shared_buffer_size*sizeof(float),0>>>
+			(d_idata, d_odata, element_size, element_count, src_stride,
+			 dst_stride, shared_buffer_size, num_compute_warps*WARP_SIZE,
+			 dma_warps*WARP_SIZE);
+	}
+	else if (elem_bytes <= 128)
+	{
+		dma_ld_test<ALIGNMENT,ALIGN_OFFSET,128>
+			<<<1,total_threads,shared_buffer_size*sizeof(float),0>>>
+			(d_idata, d_odata, element_size, element_count, src_stride,
+			 dst_stride, shared_buffer_size, num_compute_warps*WARP_SIZE,
+			 dma_warps*WARP_SIZE);
+	}
+	else if (elem_bytes <= 256)
+	{
+		dma_ld_test<ALIGNMENT,ALIGN_OFFSET,256>
+			<<<1,total_threads,shared_buffer_size*sizeof(float),0>>>
+			(d_idata, d_odata, element_size, element_count, src_stride,
+			 dst_stride, shared_buffer_size, num_compute_warps*WARP_SIZE,
+			 dma_warps*WARP_SIZE);
+	}
+	else if (elem_bytes <= 512)
+	{
+		dma_ld_test<ALIGNMENT,ALIGN_OFFSET,512>
+			<<<1,total_threads,shared_buffer_size*sizeof(float),0>>>
+			(d_idata, d_odata, element_size, element_count, src_stride,
+			 dst_stride, shared_buffer_size, num_compute_warps*WARP_SIZE,
+			 dma_warps*WARP_SIZE);
+	}
+	else if (elem_bytes <= 1024)
+	{
+		dma_ld_test<ALIGNMENT,ALIGN_OFFSET,1024>
+			<<<1,total_threads,shared_buffer_size*sizeof(float),0>>>
+			(d_idata, d_odata, element_size, element_count, src_stride,
+			 dst_stride, shared_buffer_size, num_compute_warps*WARP_SIZE,
+			 dma_warps*WARP_SIZE);
+	}
+	else if (elem_bytes <= 2048)
+	{
+		dma_ld_test<ALIGNMENT,ALIGN_OFFSET,2048>
+			<<<1,total_threads,shared_buffer_size*sizeof(float),0>>>
+			(d_idata, d_odata, element_size, element_count, src_stride,
+			 dst_stride, shared_buffer_size, num_compute_warps*WARP_SIZE,
+			 dma_warps*WARP_SIZE);
+	}
+	else if (elem_bytes <= 4096)
+	{
+		dma_ld_test<ALIGNMENT,ALIGN_OFFSET,4096>
+			<<<1,total_threads,shared_buffer_size*sizeof(float),0>>>
+			(d_idata, d_odata, element_size, element_count, src_stride,
+			 dst_stride, shared_buffer_size, num_compute_warps*WARP_SIZE,
+			 dma_warps*WARP_SIZE);
+	}
+	else if (elem_bytes <= 8192)
+	{
+		dma_ld_test<ALIGNMENT,ALIGN_OFFSET,8192>
+			<<<1,total_threads,shared_buffer_size*sizeof(float),0>>>
+			(d_idata, d_odata, element_size, element_count, src_stride,
+			 dst_stride, shared_buffer_size, num_compute_warps*WARP_SIZE,
+			 dma_warps*WARP_SIZE);
+	}
+	else
+	{
+		dma_ld_test<ALIGNMENT,ALIGN_OFFSET,0>
+			<<<1,total_threads,shared_buffer_size*sizeof(float),0>>>
+			(d_idata, d_odata, element_size, element_count, src_stride,
+			 dst_stride, shared_buffer_size, num_compute_warps*WARP_SIZE,
+			 dma_warps*WARP_SIZE);
+	}
 	CUDA_SAFE_CALL( cudaThreadSynchronize());
 
 	CUDA_SAFE_CALL( cudaMemcpy (h_odata, d_odata, output_size*sizeof(float), cudaMemcpyDeviceToHost));

	success04_2 = success04_2 && run_all_experiments<4,2>(8192,32,16);
	success04_3 = success04_3 && run_all_experiments<4,3>(8192,32,16);

	fprintf(stdout,"\nResults:\n");
	fprintf(stdout,"\tAlignment16-Offset0: %s\n",(success16_0?"SUCCESS":"FAILURE"));
	fprintf(stdout,"\tAlignment08-Offset0: %s\n",(success08_0?"SUCCESS":"FAILURE"));

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

template<int ALIGNMENT, int ALIGN_OFFSET, int MAX_BYTES_PER_ELMT>
__global__ void __launch_bounds__(1024,1)
dma_ld_test ( float *idata, float *odata, int element_size /*in number of floats*/, int num_elements, int src_stride/*bytes*/, int dst_stride/*bytes*/, int buffer_size /*number of floats*/, int num_compute_threads, int num_dma_threads_per_ld)
{
	extern __shared__ float buffer[];	

	cudaDMAStrided<MAX_BYTES_PER_ELMT,ALIGNMENT>
	  dma0 (1, num_dma_threads_per_ld, num_compute_threads,
		num_compute_threads, element_size*sizeof(float),
		num_elements, src_stride*sizeof(float), dst_stride*sizeof(float));
