    * both evaluation targets as local const expressions
    * @todo verify if the bug still exists, if not, remove this enum
    * representation.
-   */
-
-  enum
-  {
-    // Number of scan lanes
-    log_scan_lanes = _LOG_SCAN_LANES,
-    scan_lanes = 1 << log_scan_lanes,
-
-    // Number of partials per lange
-    log_partials_per_lane = _LOG_ACTIVE_THREADS,
-    partials_per_lane = 1 << log_partials_per_lane,
-
-    // Number of raking threads
-    log_raking_threads = _LOG_RAKING_THREADS,
-    raking_threads = 1 << log_raking_threads,
-
-    // Number of raking threads per lane
-    log_raking_threads_per_lane =
-      log_raking_threads - log_scan_lanes, // must be positive!
-    raking_threads_per_lane = 1 << log_raking_threads_per_lane,
-
-    // Partials to be raked per raking thread
-    log_partials_per_segment =
-      log_partials_per_lane - log_raking_threads_per_lane,
-    partials_per_segment = 1 << log_partials_per_segment,
-
-    // Number of partials that we can put in one stripe across the shared memory
-    // banks
-    log_partials_per_bank_array =
-      util::math::log2(util::properties::shared_memory_banks()) +
-      util::math::log2(util::properties::shared_memory_bank_stride()) -
-      util::math::log2(sizeof(value_t)),
-    partials_per_bank_array = 1 << log_partials_per_bank_array,
-
-    log_segments_per_bank_array =
-      util::limits::max(0,
-                        log_partials_per_bank_array - log_partials_per_segment),
-    segments_per_bank_array = 1 << log_segments_per_bank_array,
-
-    // Whether or not one warp of raking threads can rake entirely in one stripe
-    // across the shared memory banks
-    no_padding =
-      (log_segments_per_bank_array >=
+   * @note (06/02/2020) removed enum
+   */
+
+  // Number of scan lanes
+  auto log_scan_lanes = _LOG_SCAN_LANES;
+  auto scan_lanes = 1 << log_scan_lanes;
+
+  // Number of partials per lange
+  auto log_partials_per_lane = _LOG_ACTIVE_THREADS;
+  auto partials_per_lane = 1 << log_partials_per_lane;
+
+  // Number of raking threads
+  auto log_raking_threads = _LOG_RAKING_THREADS;
+  auto raking_threads = 1 << log_raking_threads;
+
+  // Number of raking threads per lane
+  auto log_raking_threads_per_lane =
+    log_raking_threads - log_scan_lanes; // must be positive!
+  auto raking_threads_per_lane = 1 << log_raking_threads_per_lane;
+
+  // Partials to be raked per raking thread
+  auto log_partials_per_segment =
+    log_partials_per_lane - log_raking_threads_per_lane;
+  auto partials_per_segment = 1 << log_partials_per_segment;
+
+  // Number of partials that we can put in one stripe across the shared memory
+  // banks
+  auto log_partials_per_bank_array =
+    util::math::log2(util::properties::shared_memory_banks()) +
+    util::math::log2(util::properties::shared_memory_bank_stride()) -
+    util::math::log2(sizeof(value_t));
+  auto partials_per_bank_array = 1 << log_partials_per_bank_array;
+
+  auto log_segments_per_bank_array =
+    util::limits::max(0,
+                      log_partials_per_bank_array - log_partials_per_segment);
+  auto segments_per_bank_array = 1 << log_segments_per_bank_array;
+
+  // Whether or not one warp of raking threads can rake entirely in one stripe
+  // across the shared memory banks
+  auto no_padding =
+    (log_segments_per_bank_array >=
+     util::math::log2(util::properties::maximum_threads_per_warp()));
+
+  // Number of raking segments we can have without padding (i.e., a "row")
+  auto log_segments_per_row =
+    (no_padding) ? log_raking_threads : // All raking threads (segments)
+      util::limits::min(
+        log_raking_threads_per_lane,
+        log_segments_per_bank_array); // Up to as many segments per lane (all
+                                      // lanes must have same amount of
+                                      // padding to have constant lane stride)
+
+  auto segments_per_row = 1 << log_segments_per_row;
+
+  // Number of partials per row
+  auto log_partials_per_row = log_segments_per_row + log_partials_per_segment;
+  auto partials_per_row = 1 << log_partials_per_row;
+
+  // Number of partials that we must use to "pad out" one memory bank
+  auto log_bank_padding_partials = util::limits::max(
+    0,
+    util::math::log2(util::properties::shared_memory_bank_stride()) -
+      util::math::log2(sizeof(value_t)));
+  auto bank_padding_partials = 1 << log_bank_padding_partials;
+
+  // Number of partials that we must use to "pad out" a lane to one memory
+  // bank
+  auto lane_padding_partials =
+    util::limits::max(0, partials_per_bank_array - partials_per_lane);
+
+  // Number of partials (including padding) per "row"
+  auto padded_partials_per_row =
+    (no_padding)
+      ? partials_per_row
+      : partials_per_row + lane_padding_partials + bank_padding_partials;
+
+  // Number of rows in a grid
+  auto log_rows = log_raking_threads - log_segments_per_row;
+  auto rows = 1 << log_rows;
+
+  // Number of rows per lane (always at least one)
+  auto log_rows_per_lane =
+    util::limits::max(0, log_raking_threads_per_lane - log_segments_per_row);
+  auto rows_per_lane = 1 << log_rows_per_lane;
+
+  // Padded stride between lanes (in partials)
+  auto lane_stride =
+    (no_padding) ? partials_per_lane : rows_per_lane * padded_partials_per_row;
+
+  // Number of elements needed to back this level of the raking grid
+  auto raking_elements = rows * padded_partials_per_row;
+
+  // If there are prefix dependences between lanes, a secondary raking grid
+  // type will be needed in the event we have more than one warp of raking
+  // threads
+
+  typedef
+    typename util::meta::_if<(util::numeric_traits<value_t>::REPRESENTATION ==
+                              util::representation_t::not_a_number),
+                             value_t,
+                             volatile value_t>::type_t warp_scan_t;
+
+  typedef typename util::meta::_if<
+    _DEPENDENT_LANES &&
+      (log_raking_threads >
        util::math::log2(util::properties::maximum_threads_per_warp())),
-
-    // Number of raking segments we can have without padding (i.e., a "row")
-    log_segments_per_row =
-      (no_padding) ? log_raking_threads : // All raking threads (segments)
-        util::limits::min(
-          log_raking_threads_per_lane,
-          log_segments_per_bank_array), // Up to as many segments per lane (all
-                                        // lanes must have same amount of
-                                        // padding to have constant lane stride)
-
-    segments_per_row = 1 << log_segments_per_row,
-
-    // Number of partials per row
-    log_partials_per_row = log_segments_per_row + log_partials_per_segment,
-    partials_per_row = 1 << log_partials_per_row,
-
-    // Number of partials that we must use to "pad out" one memory bank
-    log_bank_padding_partials = util::limits::max(
-      0,
-      util::math::log2(util::properties::shared_memory_bank_stride()) -
-        util::math::log2(sizeof(value_t))),
-    bank_padding_partials = 1 << log_bank_padding_partials,
-
-    // Number of partials that we must use to "pad out" a lane to one memory
-    // bank
-    lane_padding_partials =
-      util::limits::max(0, partials_per_bank_array - partials_per_lane),
-
-    // Number of partials (including padding) per "row"
-    padded_partials_per_row =
-      (no_padding)
-        ? partials_per_row
-        : partials_per_row + lane_padding_partials + bank_padding_partials,
-
-    // Number of rows in a grid
-    log_rows = log_raking_threads - log_segments_per_row,
-    rows = 1 << log_rows,
-
-    // Number of rows per lane (always at least one)
-    log_rows_per_lane =
-      util::limits::max(0, log_raking_threads_per_lane - log_segments_per_row),
-    rows_per_lane = 1 << log_rows_per_lane,
-
-    // Padded stride between lanes (in partials)
-    lane_stride = (no_padding) ? partials_per_lane
-                               : rows_per_lane * padded_partials_per_row,
-
-    // Number of elements needed to back this level of the raking grid
-    raking_elements = rows * padded_partials_per_row,
-  }; // enum
+    raking<               // Secondary grid
+                          // GUNROCK_CUDA_ARCH,
+      log_raking_threads, // Depositing threads (the primary raking threads)
+      0, // 1 lane (the primary raking threads only make one deposit)
+      util::math::log2(
+        util::properties::maximum_threads_per_warp()), // Raking threads (1
+                                                       // warp)
+      false,    // There is only one lane, so there are no inter-lane prefix
+                // dependences
+      value_t>, // Partial type
+    util::meta::null_t> // No secondary grid
+    ::type_t secondary_grid_t;
+
+  /**
+   * Utility class for totaling the SMEM elements needed for an raking grid
+   * hierarchy
+   */
+  template<typename raking, int __dummy = 0>
+  struct total_raking_elements
+  {
+    // Recurse
+    auto VALUE = raking::raking_elements +
+                 total_raking_elements<typename raking::secondary_grid_t>;
+  };
+
+  template<int __dummy>
+  struct total_raking_elements<util::meta::null_t, __dummy>
+  {
+    // Terminate
+    auto VALUE = 0
+  };
+
+  // Total number of smem raking elements needed back this hierarchy
+  // of raking grids (may be reused for other purposes)
+  auto total_raking_elements = total_raking_elements<raking>::VALUE;
+
+  /**
+   * Type of pointer for inserting partials into lanes, e.g.,
+   * lane_partial[LANE][0] = ...
+   */
+  typedef value_t (*lane_partial_t)[lane_stride];
+
+  /**
+   * Type of pointer for raking across lane segments
+   */
+  typedef value_t* raking_segment_t;
+
+  // Returns the location in the smem grid where the calling thread can
+  // insert/extract its partial for raking reduction/scan into the first lane.
+  // Positions in subsequent lanes can be obtained via increments of
+  // lane_stride.
+  static GUNROCK_HOST_DEVICE lane_partial_t my_lane_partial(value_t* smem)
+  {
+    int row = threadIdx.x >> log_partials_per_row;
+    int col = threadIdx.x & (partials_per_row - 1);
+
+    return (lane_partial_t)(smem + (row * padded_partials_per_row) + col);
+  }
+
+  // Returns the location in the smem grid where the calling thread can begin
+  // serial raking/scanning
+  static GUNROCK_HOST_DEVICE raking_segment_t my_raking_segment(value_t* smem)
+  {
+    int row = threadIdx.x >> log_segments_per_row;
+    int col = (threadIdx.x & (segments_per_row - 1))
+              << log_partials_per_segment;
+    return (raking_segment_t)(smem + (row * padded_partials_per_row) + col);
+  }
+
+  /**
+   * Displays configuration to standard out
+   */
+  static GUNROCK_HOST_DEVICE void print()
+  {
+    printf("SCAN_LANES: %d\n"
+           "PARTIALS_PER_LANE: %d\n"
+           "RAKING_THREADS: %d\n"
+           "RAKING_THREADS_PER_LANE: %d\n"
+           "PARTIALS_PER_SEG: %d\n"
+           "PARTIALS_PER_BANK_ARRAY: %d\n"
+           "SEGS_PER_BANK_ARRAY: %d\n"
+           "NO_PADDING: %d\n"
+           "SEGS_PER_ROW: %d\n"
+           "PARTIALS_PER_ROW: %d\n"
+           "BANK_PADDING_PARTIALS: %d\n"
+           "LANE_PADDING_PARTIALS: %d\n"
+           "PADDED_PARTIALS_PER_ROW: %d\n"
+           "ROWS: %d\n"
+           "ROWS_PER_LANE: %d\n"
+           "LANE_STRIDE: %d\n"
+           "RAKING_ELEMENTS: %d\n",
+           scan_lanes,
+           partials_per_lane,
+           raking_threads,
+           raking_threads_per_lane,
+           partials_per_segment,
+           partials_per_bank_array,
+           segments_per_bank_array,
+           no_padding,
+           segments_per_row,
+           partials_per_row,
+           bank_padding_partials,
+           lane_padding_partials,
+           padded_partials_per_row,
+           rows,
+           rows_per_lane,
+           lane_stride,
+           raking_elements);
+  }
+
+}; // struct raking
+
+} // namespace grid
+
+namespace details {
+
+/**
+ * Operational details for threads working in an raking grid
+ */
+template<typename grid::raking,
+         typename secondary_raking_t = typename grid::raking::secondary_grid_t>
+struct raking;
+
+/**
+ * Operational details for threads working in an raking grid (specialized for
+ * one-level raking grid)
+ */
+template<typename grid::raking>
+struct raking<grid::raking, util::meta::null_t> : grid::raking
+{
+  auto queue_rsvn_thread = 0;
+  auto cumulative_thread = grid::raking::raking_threads - 1;
+  auto warp_threads = util::properties::maximum_threads_per_warp();
+
+  typedef typename grid::raking::value_t value_t; // Partial type
+  typedef typename grid::raking::warpscan_t (
+    *warpscan_storage_t)[warp_threads]; // Warpscan storage type
+  typedef typename util::meta::null_t
+    secondary_raking_t; // Type of next-level grid raking details
+
+  /**
+   * Smem pool backing raking grid lanes
+   */
+  value_t* smem_pool;
+
+  /**
+   * Warpscan storage
+   */
+  warpscan_storage_t warpscan;
+
+  /**
+   * The location in the smem grid where the calling thread can insert/extract
+   * its partial for raking reduction/scan into the first lane.
+   */
+  typename grid::raking::lane_partial_t lane_partial;
+
+  /**
+   * Returns the location in the smem grid where the calling thread can begin
+   * serial raking/scanning
+   */
+  typename grid::raking::raking_segment_t raking_segment;
+
+  // Constructor
+  GUNROCK_F_DEVICE raking(value_t* smem_pool)
+    : smem_pool(smem_pool)
+    , lane_partial(
+        grid::raking::my_lane_partial(smem_pool)) // set lane partial pointer
+  {
+    if (threadIdx.x < grid::raking::raking_threads) {
+      // Set raking segment pointer
+      raking_segment = grid::raking::my_raking_segment(smem_pool);
+    }
+  }
+
+  // Constructor
+  GUNROCK_F_DEVICE raking(value_t* smem_pool, warpscan_storage_t warpscan)
+    : smem_pool(smem_pool)
+    , warpscan(warpscan)
+    , lane_partial(
+        grid::raking::my_lane_partial(smem_pool)) // set lane partial pointer
+  {
+    if (threadIdx.x < grid::raking::raking_threads) {
+      // Set raking segment pointer
+      raking_segment = grid::raking::my_raking_segment(smem_pool);
+    }
+  }
+
+  // Constructor
+  GUNROCK_F_DEVICE raking(value_t* smem_pool,
+                          warpscan_storage_t warpscan,
+                          value_t warpscan_identity)
+    : smem_pool(smem_pool)
+    , warpscan(warpscan)
+    , lane_partial(
+        grid::raking::my_lane_partial(smem_pool)) // set lane partial pointer
+  {
+    if (threadIdx.x < grid::raking::RAKING_THREADS) {
+      // Initialize first half of warpscan storage to identity
+      warpscan[0][threadIdx.x] = warpscan_identity;
+
+      // Set raking segment pointer
+      raking_segment = grid::raking::my_raking_segment(smem_pool);
+    }
+  }
+
+  // Return the cumulative partial left in the final warpscan cell
+  GUNROCK_F_DEVICE value_t cumulative_partial() const
+  {
+    return warpscan[1][cumulative_thread];
+  }
+
+  // Return the queue reservation in the first warpscan cell
+  GUNROCK_F_DEVICE value_t queue_reservation() const
+  {
+    return warpscan[1][queue_rsvn_thread];
+  }
+
+  GUNROCK_F_DEVICE value_t* get_smem_pool() { return smem_pool; }
+
+}; // struct raking
+
+/**
+ * Operational details for threads working in a hierarchical raking grid
+ */
+template<typename grid::raking, typename secondary_raking_t>
+struct raking : grid::raking
+{
+
+  auto cumulative_thread = grid::raking::raking_threads - 1;
+  auto warp_threads = util::properties::maximum_threads_per_warp();
+
+  typedef typename grid::raking::value_t value_t; // Partial type
+  typedef typename grid::raking::warpscan_t (
+    *warpscan_storage_t)[warp_threads]; // Warpscan storage type
+  typedef raking<secondary_raking_t>
+    secondary_raking_details_t; // Type of next-level grid raking details
+
+  /**
+   * The location in the smem grid where the calling thread can insert/extract
+   * its partial for raking reduction/scan into the first lane.
+   */
+  typename grid::raking::lane_partial_t lane_partial;
+
+  /**
+   * Returns the location in the smem grid where the calling thread can begin
+   * serial raking/scanning
+   */
+  typename grid::raking::raking_segment_t raking_segment;
+
+  /**
+   * Secondary-level grid details
+   */
+  secondary_raking_details_t secondary_raking;
+
+  // Constructor
+  GUNROCK_F_DEVICE raking(value_t* smem_pool)
+    : lane_partial(grid::raking::my_lane_partial(smem_pool))
+    , // set lane partial pointer
+    secondary_raking(smem_pool + grid::raking::raking_elements)
+  {
+    if (threadIdx.x < grid::raking::raking_threads) {
+      // Set raking segment pointer
+      raking_segment = grid::raking::my_raking_segment(smem_pool);
+    }
+  }
+
+  // Constructor
+  GUNROCK_F_DEVICE raking(value_t* smem_pool,
+                          warpscan_storage_t warpscan)
+    : lane_partial(grid::raking::my_lane_partial(smem_pool))
+    , // set lane partial pointer
+    secondary_raking(smem_pool + grid::raking::raking_elements, warpscan)
+  {
+    if (threadIdx.x < grid::raking::raking_threads) {
+      // Set raking segment pointer
+      raking_segment = grid::raking::my_raking_segment(smem_pool);
+    }
+  }
+
+  // Constructor
+  GUNROCK_F_DEVICE raking(value_t* smem_pool,
+                          warpscan_storage_t warpscan,
+                          value_t warpscan_identity)
+    : lane_partial(grid::raking::my_lane_partial(smem_pool))
+    , // set lane partial pointer
+    secondary_raking(smem_pool + grid::raking::raking_elements,
+                     warpscan,
+                     warpscan_identity)
+  {
+    if (threadIdx.x < grid::raking::raking_threads) {
+      // Set raking segment pointer
+      raking_segment = grid::raking::my_raking_segment(smem_pool);
+    }
+  }
+
+  // Return the cumulative partial left in the final warpscan cell
+  GUNROCK_F_DEVICE value_t cumulative_partial() const
+  {
+    return secondary_raking.cumulative_partial();
+  }
+
+  // Return the queue reservation in the first warpscan cell
+  GUNROCK_F_DEVICE value_t queue_reservation() const
+  {
+    return secondary_raking.queue_reservation();
+  }
+
+  /**
+   *
+   */
+  GUNROCK_F_DEVICE value_t* get_smem_pool()
+  {
+    return secondary_raking.get_smem_pool();
+  }
 };
 
-} // namespace grid
-
+} // namespace details
 } // namespace device
 } // namespace raking
 

 */
namespace grid {

template<architecture_t GUNROCK_CUDA_ARCH,
         unsigned _LOG_ACTIVE_THREADS, // Number of threads placing a lane
                                       // partial (i.e., the number of partials
                                       // per lane)
         unsigned _LOG_SCAN_LANES,     // Number of scan lanes
         unsigned _LOG_RAKING_THREADS, // Number of threads used for raking
                                       // (typically 1 warp)
         bool _DEPENDENT_LANES, // If there are prefix dependences between lanes
                                // (i.e., downsweeping will incorporate
                                // aggregates from previous lanes)
         typename raking_value_t // Type of items we will be reducing/scanning
         >
struct rake
{
  // Type of items we will be reducing/scanning
  typedef raking_value_t value_t;

 */

#pragma once
#include <gunrock/util/device_properties.hxx>
#include <gunrock/util/limits.hxx>
#include <gunrock/util/meta.hxx>
#include <gunrock/util/numeric_traits.hxx>
