 }  // anonymous namespace
 
 struct Clip : public thrust::unary_function<GradientPair, Pair> {
-  static XGBOOST_DEV_INLINE float Pclip(float v) {
-    return v > 0 ? v : 0;
-  }
-  static XGBOOST_DEV_INLINE float Nclip(float v) {
-    return v < 0 ? abs(v) : 0;
-  }
+  static XGBOOST_DEV_INLINE float Pclip(float v) { return v > 0 ? v : 0; }
+  static XGBOOST_DEV_INLINE float Nclip(float v) { return v < 0 ? abs(v) : 0; }
 
   XGBOOST_DEV_INLINE Pair operator()(GradientPair x) const {
     auto pg = Pclip(x.GetGrad());

  using T = typename GradientSumT::ValueT;
  dh::XGBCachingDeviceAllocator<char> alloc;

  thrust::device_ptr<GradientPair const> gpair_beg {gpair.data()};
  thrust::device_ptr<GradientPair const> gpair_end {gpair.data() + gpair.size()};
  auto beg = thrust::make_transform_iterator(gpair_beg, Clip());
  auto end = thrust::make_transform_iterator(gpair_end, Clip());
  Pair p = dh::Reduce(thrust::cuda::par(alloc), beg, end, Pair{}, thrust::plus<Pair>{});
  GradientPair positive_sum {p.first}, negative_sum {p.second};

  auto histogram_rounding = GradientSumT {
    CreateRoundingFactor<T>(std::max(positive_sum.GetGrad(), negative_sum.GetGrad()),
                            gpair.size()),
    CreateRoundingFactor<T>(std::max(positive_sum.GetHess(), negative_sum.GetHess()),
                            gpair.size()) };

  using IntT = typename HistRounding<GradientSumT>::SharedSumT::ValueT;


                            common::Span<GradientPair const> gpair,
                            common::Span<const uint32_t> d_ridx,
                            common::Span<GradientSumT> histogram,
                            HistRounding<GradientSumT> rounding, bool force_global_memory) {
  // decide whether to use shared memory
  int device = 0;
  dh::safe_cuda(cudaGetDevice(&device));
  // opt into maximum shared memory for the kernel if necessary
  size_t max_shared_memory = dh::MaxSharedMemoryOptin(device);

  size_t smem_size =
      sizeof(typename HistRounding<GradientSumT>::SharedSumT) * feature_groups.max_group_bins;
  bool shared = !force_global_memory && smem_size <= max_shared_memory;
  smem_size = shared ? smem_size : 0;

  constexpr int kBlockThreads = 1024;
  constexpr int kItemsPerThread = 8;
  constexpr int kItemsPerTile = kBlockThreads * kItemsPerThread;

  auto runit = [&](auto kernel) {
    if (shared) {
      dh::safe_cuda(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize,
                                         max_shared_memory));
    }

    // determine the launch configuration
    int num_groups = feature_groups.NumGroups();
    int n_mps = 0;
    dh::safe_cuda(cudaDeviceGetAttribute(&n_mps, cudaDevAttrMultiProcessorCount, device));
    int n_blocks_per_mp = 0;
    dh::safe_cuda(cudaOccupancyMaxActiveBlocksPerMultiprocessor(&n_blocks_per_mp, kernel,
                                                                kBlockThreads, smem_size));
    // This gives the number of blocks to keep the device occupied
    // Use this as the maximum number of blocks
    unsigned grid_size = n_blocks_per_mp * n_mps;

    // Otherwise launch blocks such that each block has a minimum amount of work to do
    // There are fixed costs to launching each block, e.g. zeroing shared memory
    // The below amount of minimum work was found by experimentation
    constexpr int kMinItemsPerBlock = kItemsPerTile;
    int columns_per_group = common::DivRoundUp(matrix.row_stride, feature_groups.NumGroups());
    // Average number of matrix elements processed by each group
    std::size_t items_per_group = d_ridx.size() * columns_per_group;

    // Allocate number of blocks such that each block has about kMinItemsPerBlock work
    // Up to a maximum where the device is saturated
    grid_size =
        min(grid_size,
            unsigned(common::DivRoundUp(items_per_group, kMinItemsPerBlock)));

    dh::LaunchKernel {dim3(grid_size, num_groups),
        static_cast<uint32_t>(kBlockThreads), smem_size}(
        kernel, matrix, feature_groups, d_ridx, histogram.data(), gpair.data(), rounding);
  };

  if (shared) {
    runit(SharedMemHistKernel<GradientSumT, true, kBlockThreads, kItemsPerThread>);
  } else {
    runit(SharedMemHistKernel<GradientSumT, false, kBlockThreads, kItemsPerThread>);
  }

  dh::safe_cuda(cudaGetLastError());
}

template void BuildGradientHistogram<GradientPairPrecise>(
    EllpackDeviceAccessor const& matrix, FeatureGroupsAccessor const& feature_groups,
    common::Span<GradientPair const> gpair, common::Span<const uint32_t> ridx,
    common::Span<GradientPairPrecise> histogram, HistRounding<GradientPairPrecise> rounding,
    bool force_global_memory);

}  // namespace tree
