Use single precision in gain calculation, use pointers instead of span. (#8051)

src/tree/gpu_hist/evaluate_splits.cuh

@@ -68,7 +68,7 @@ class GPUHistEvaluator {
   // storage for sorted index of feature histogram, used for sort based splits.
   dh::device_vector<bst_feature_t> cat_sorted_idx_;
   // cached input for sorting the histogram, used for sort based splits.
-  using SortPair = thrust::tuple<uint32_t, double>;
+  using SortPair = thrust::tuple<uint32_t, float>;
   dh::device_vector<SortPair> sort_input_;
   // cache for feature index
   dh::device_vector<bst_feature_t> feature_idx_;

src/tree/gpu_hist/evaluator.cu

@@ -89,7 +89,7 @@ common::Span<bst_feature_t const> GPUHistEvaluator<GradientSumT>::SortHistogram(
                                                           input.gradient_histogram[j]);
                         return thrust::make_tuple(i, lw);
                       }
-                      return thrust::make_tuple(i, 0.0);
+                      return thrust::make_tuple(i, 0.0f);
                     });
   // Sort an array segmented according to
   // - nodes