 	constexpr int gridsize = 1; // EY : 20170304 for number of blocks on a grid, it seems like it's 1 for block reduce, which wasn't explicitly said so, but seemed implied in the CUB documentation
 	constexpr int M_x1 = L1;
 	constexpr int d1   = 1;
+
+	constexpr int M_x2 = L1;
+	constexpr int d2   = 2;
 	
 	///////////// END of BLOCK, GRID DIMENSIONS ////////////////////////
 	////////////////////////////////////////////////////////////////////

	// boilerplate for interesting values to test out for reduce ///////
	////////////////////////////////////////////////////////////////////
	constexpr int L1 = 1024; // "TILE_SIZE"
	
	
	// Allocate problem device arrays
    float 	*dev_f = nullptr;  // NULL, otherwise for nullptr, you're going to need the -std=c++11 flag for compilation
	float   *dev_fmax = nullptr;
	float   *dev_fsum = nullptr;
	
	/* Allocate n floats on device */
    CubDebugExit(dev_allocator.DeviceAllocate((void**)&dev_f, sizeof(float) * L1 ));

	/* Allocate floats on device */
    CubDebugExit(dev_allocator.DeviceAllocate((void**)&dev_fmax, sizeof(float) ));

	CubDebugExit( cudaDeviceSynchronize());

	BlockSum1DKernel<M_x1,d1,cub::BLOCK_REDUCE_WARP_REDUCTIONS><<<gridsize,M_x1>>>(dev_f, dev_fsum);

	/* Copy device memory to host */
	CubDebugExit(cudaMemcpy(f_vec.data(), dev_f, sizeof(float) * L1 , cudaMemcpyDeviceToHost));
