 /*!
  * Copyright 2020-2021 by XGBoost Contributors
  */
+#include <thrust/iterator/transform_iterator.h>
 #include <thrust/reduce.h>
-#include <thrust/iterator/transform_iterator.h>
+
 #include <algorithm>
 #include <ctgmath>
 #include <limits>
 
+#include "../../common/device_helpers.cuh"
+#include "../../data/ellpack_page.cuh"
+#include "histogram.cuh"
+#include "row_partitioner.cuh"
 #include "xgboost/base.h"
-#include "row_partitioner.cuh"
-
-#include "histogram.cuh"
-
-#include "../../data/ellpack_page.cuh"
-#include "../../common/device_helpers.cuh"
 
 namespace xgboost {
 namespace tree {

   * rounding is calcuated as exp(m), see the rounding factor calcuation for
   * details.
   */
  GradientSumT to_fixed_point = GradientSumT(
      T(1) / to_floating_point.GetGrad(), T(1) / to_floating_point.GetHess());

  return {histogram_rounding, to_fixed_point, to_floating_point};
}

template HistRounding<GradientPairPrecise>
CreateRoundingFactor(common::Span<GradientPair const> gpair);
template HistRounding<GradientPair>
CreateRoundingFactor(common::Span<GradientPair const> gpair);

template <typename GradientSumT, bool use_shared_memory_histograms>
__global__ void SharedMemHistKernel(EllpackDeviceAccessor matrix,
                                    FeatureGroupsAccessor feature_groups,
                                    common::Span<const RowPartitioner::RowIndexT> d_ridx,
                                    GradientSumT* __restrict__ d_node_hist,
                                    const GradientPair* __restrict__ d_gpair,
                                    HistRounding<GradientSumT> const rounding) {
  using SharedSumT = typename HistRounding<GradientSumT>::SharedSumT;
  using T = typename GradientSumT::ValueT;

  extern __shared__ char smem[];
  FeatureGroup group = feature_groups[blockIdx.y];
  SharedSumT *smem_arr = reinterpret_cast<SharedSumT *>(smem);
  if (use_shared_memory_histograms) {
    dh::BlockFill(smem_arr, group.num_bins, SharedSumT());
    __syncthreads();
  }
  int feature_stride = matrix.is_dense ? group.num_features : matrix.row_stride;
  size_t n_elements = feature_stride * d_ridx.size();
  for (auto idx : dh::GridStrideRange(static_cast<size_t>(0), n_elements)) {
    int ridx = d_ridx[idx / feature_stride];
    int gidx = matrix.gidx_iter[ridx * matrix.row_stride + group.start_feature +
                                idx % feature_stride];
    if (gidx != matrix.NumBins()) {
      // If we are not using shared memory, accumulate the values directly into
      // global memory
      gidx = use_shared_memory_histograms ? gidx - group.start_bin : gidx;
      if (use_shared_memory_histograms) {
        auto adjusted = rounding.ToFixedPoint(d_gpair[ridx]);
        dh::AtomicAddGpair(smem_arr + gidx, adjusted);
      } else {
        GradientSumT truncated{
            TruncateWithRoundingFactor<T>(rounding.rounding.GetGrad(),
                                          d_gpair[ridx].GetGrad()),
            TruncateWithRoundingFactor<T>(rounding.rounding.GetHess(),
                                          d_gpair[ridx].GetHess()),
        };
        dh::AtomicAddGpair(d_node_hist + gidx, truncated);
      }
    }
  }

  if (use_shared_memory_histograms) {
    // Write shared memory back to global memory
    __syncthreads();
    for (auto i : dh::BlockStrideRange(0, group.num_bins)) {
      auto truncated = rounding.ToFloatingPoint(smem_arr[i]);
      dh::AtomicAddGpair(d_node_hist + group.start_bin + i, truncated);
    }
  }
}


   */
  GradientSumT to_floating_point =
      histogram_rounding /
      T(IntT(1) << (sizeof(typename GradientSumT::ValueT) * 8 - 2));  // keep 1 for sign bit
  /**
   * Factor for converting gradients from floating-point to fixed-point. For
   * f64:
