   *
   *  File: reduce.cuh
   *  Created: Feb 12, 2013
-  *  Modified: Mon 18 Feb 2013 01:04:37 PM PST
+  *  Modified: Sun 24 Feb 2013 06:51:54 PM PST
   *
   *  Author: Abhinav Sarje <asarje@lbl.gov>
   */

			data_t temp_result;
			cudaMemcpy(&temp_result, &output[0], sizeof(data_t), cudaMemcpyDeviceToHost);
			result = op(result, temp_result);
			//std::cout << grid << ". result: " << result << std::endl;
		} // for

		output = base_output;

	} // reduce_block()


	template <typename iterator_t, typename data_t, typename functor_t>
	data_t reduce(iterator_t start, iterator_t end, data_t init, functor_t op) {

		unsigned int num_subtiles = NUM_SUBTILES_;
		unsigned int n = end - start;
		unsigned int block_size = BLOCK_DIM_REDUCE_;
		unsigned int n_grid_max = MAX_GRID_SIZE_ * block_size * num_subtiles;
		unsigned int num_grids = ceil((float) n / n_grid_max);

		data_t *output = NULL, *base_output = NULL;
		unsigned int output_size = MAX_GRID_SIZE_;
		cudaMalloc((void**) &output, output_size * sizeof(data_t));
		base_output = output;

		unsigned int d_shmem_size = block_size * sizeof(data_t);

		data_t result = init;
		for(unsigned int grid = 0; grid < num_grids; ++ grid) {	// for each grid
