 	} // reduce_block()
 
 
-	// this performs better than thrust::reduce
 	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_;
-		unsigned int n_grid_max = GRID_SIZE_ * block_size * num_subtiles;
+		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 = GRID_SIZE_;
+		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);
-
-		//std::cout << "num_grids: " << num_grids << ", n: " << n
-		//			<< ", d_shmem_size: " << d_shmem_size << std::endl;
 
 		data_t result = init;
 		for(unsigned int grid = 0; grid < num_grids; ++ grid) {	// for each grid

  *
  *  File: reduce.cuh
  *  Created: Feb 12, 2013
  *  Modified: Mon 18 Feb 2013 01:04:37 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);
		} // for

		output = base_output;
