[backport] Remove __restrict__. (dmlc#11669)

Author: trivialfis

src/tree/gpu_hist/evaluate_splits.cu

@@ -110,7 +110,7 @@ class EvaluateSplitAgent {
     return gpair;
   }
 
-  __device__ __forceinline__ void Numerical(DeviceSplitCandidate *__restrict__ best_split) {
+  __device__ __forceinline__ void Numerical(DeviceSplitCandidate *best_split) {
     for (int scan_begin = gidx_begin; scan_begin < gidx_end; scan_begin += kBlockSize) {
       bool thread_active = (scan_begin + threadIdx.x) < gidx_end;
       GradientPairInt64 bin = thread_active ? LoadGpair(node_histogram + scan_begin + threadIdx.x)
@@ -147,7 +147,7 @@ class EvaluateSplitAgent {
     }
   }
 
-  __device__ __forceinline__ void OneHot(DeviceSplitCandidate *__restrict__ best_split) {
+  __device__ __forceinline__ void OneHot(DeviceSplitCandidate *best_split) {
     for (int scan_begin = gidx_begin; scan_begin < gidx_end; scan_begin += kBlockSize) {
       bool thread_active = (scan_begin + threadIdx.x) < gidx_end;
 

tests/cpp/common/test_device_vector.cu

@@ -114,7 +114,9 @@ TEST(TestVirtualMem, Version) {
 #if defined(xgboost_IS_WIN)
   ASSERT_FALSE(pinned.IsVm());
 #else  // defined(xgboost_IS_WIN)
-  if (major >= 12 && minor >= 5) {
+  if (major > 12) {
+    ASSERT_TRUE(pinned.IsVm());
+  } else if (major >= 12 && minor >= 5) {
     ASSERT_TRUE(pinned.IsVm());
   } else {
     ASSERT_FALSE(pinned.IsVm());