Avoid the use of restrict in the evaluator.

[trivialfis]

It somehow caches the written value in the registry, resulting in an incorrect value in the second iteration of the scan.

No performance loss is observed.

[trivialfis]

cc @rongou .

[Copilot]

Pull Request Overview

This PR removes the __restrict__ qualifier from evaluator methods to avoid caching the written best_split pointer in the GPU register and ensure correct behavior across iterations.

• Dropped __restrict__ on DeviceSplitCandidate* best_split parameters in four evaluator routines
• Verified no observable performance regression after removal

Comments suppressed due to low confidence (4)

src/tree/gpu_hist/evaluate_splits.cu:113

• [nitpick] Pointer declarator spacing is inconsistent; consider standardizing to either DeviceSplitCandidate* best_split or DeviceSplitCandidate *best_split across all methods for consistency.

  __device__ __forceinline__ void Numerical(DeviceSplitCandidate * best_split) {

src/tree/gpu_hist/evaluate_splits.cu:148

• [nitpick] Pointer declarator spacing is inconsistent; consider standardizing to either DeviceSplitCandidate* best_split or DeviceSplitCandidate *best_split across all methods for consistency.

  __device__ __forceinline__ void OneHot(DeviceSplitCandidate *best_split) {

src/tree/gpu_hist/evaluate_splits.cu:184

• [nitpick] Pointer declarator spacing is inconsistent; consider standardizing to either DeviceSplitCandidate* best_split or DeviceSplitCandidate *best_split across all methods for consistency.

                                                  DeviceSplitCandidate *best_split) {

src/tree/gpu_hist/evaluate_splits.cu:207

• [nitpick] Pointer declarator spacing is inconsistent; consider standardizing to either DeviceSplitCandidate* best_split or DeviceSplitCandidate *best_split across all methods for consistency.

  __device__ __forceinline__ void Partition(DeviceSplitCandidate *best_split,