Separate EvaluateFeature implementations.

src/common/device_helpers.cuh

@@ -1939,4 +1939,25 @@ class CUDAStream {
   CUDAStreamView View() const { return CUDAStreamView{stream_}; }
   void Sync() { this->View().Sync(); }
 };
+
+// Force nvcc to load data as constant
+template <typename T>
+class LDGIterator {
+  typedef typename cub::UnitWord<T>::DeviceWord DeviceWordT;
+  static constexpr std::size_t kNumWords = sizeof(T) / sizeof(DeviceWordT);
+
+  const T* ptr;
+
+ public:
+  XGBOOST_DEVICE LDGIterator(const T* ptr) : ptr(ptr) {}
+  __device__ T operator[](std::size_t idx) const {
+    DeviceWordT tmp[kNumWords];
+#pragma unroll
+    for (int i = 0; i < kNumWords; i++) {
+      tmp[i] = __ldg(reinterpret_cast<const DeviceWordT*>(ptr + idx) + i);
+    }
+    return *reinterpret_cast<const T*>(tmp);
+  }
+};
+
 }  // namespace dh