diff --git a/src/common/common.h b/src/common/common.h
index 04482a107..128776d96 100644
--- a/src/common/common.h
+++ b/src/common/common.h
@@ -39,6 +39,41 @@
 
 #endif  // defined(__CUDACC__)
 
+namespace dh {
+#if defined(__CUDACC__)
+/*
+ * Error handling  functions
+ */
+#define safe_cuda(ans) ThrowOnCudaError((ans), __FILE__, __LINE__)
+
+inline cudaError_t ThrowOnCudaError(cudaError_t code, const char *file, int line)
+{
+  if (code != cudaSuccess) {
+    LOG(FATAL) << thrust::system_error(code, thrust::cuda_category(),
+                                       std::string{file} + ": " +  // NOLINT
+                                       std::to_string(line)).what();
+  }
+  return code;
+}
+
+#elif defined(__HIP_PLATFORM_AMD__)
+/*
+ * Error handling  functions
+ */
+#define safe_cuda(ans) ThrowOnCudaError((ans), __FILE__, __LINE__)
+
+inline hipError_t ThrowOnCudaError(hipError_t code, const char *file, int line)
+{
+  if (code != hipSuccess) {
+    LOG(FATAL) << thrust::system_error(code, thrust::hip_category(),
+                                       std::string{file} + ": " +  // NOLINT
+                                       std::to_string(line)).what();
+  }
+  return code;
+}
+#endif
+}  // namespace dh
+
 namespace xgboost {
 namespace common {
 /*!
diff --git a/src/common/device_helpers.hip.h b/src/common/device_helpers.hip.h
index 36c783b49..31eb1197e 100644
--- a/src/common/device_helpers.hip.h
+++ b/src/common/device_helpers.hip.h
@@ -59,41 +59,6 @@
 
 #endif  // defined(XGBOOST_USE_RMM) && XGBOOST_USE_RMM == 1
 
-namespace dh {
-#if defined(__CUDACC__)
-/*
- * Error handling  functions
- */
-#define safe_cuda(ans) ThrowOnCudaError((ans), __FILE__, __LINE__)
-
-inline cudaError_t ThrowOnCudaError(cudaError_t code, const char *file, int line)
-{
-  if (code != cudaSuccess) {
-    LOG(FATAL) << thrust::system_error(code, thrust::cuda_category(),
-                                       std::string{file} + ": " +  // NOLINT
-                                       std::to_string(line)).what();
-  }
-  return code;
-}
-
-#elif defined(__HIP_PLATFORM_AMD__)
-/*
- * Error handling  functions
- */
-#define safe_cuda(ans) ThrowOnCudaError((ans), __FILE__, __LINE__)
-
-inline hipError_t ThrowOnCudaError(hipError_t code, const char *file, int line)
-{
-  if (code != hipSuccess) {
-    LOG(FATAL) << thrust::system_error(code, thrust::hip_category(),
-                                       std::string{file} + ": " +  // NOLINT
-                                       std::to_string(line)).what();
-  }
-  return code;
-}
-#endif
-}  // namespace dh
-
 namespace dh {
 
 // FIXME(jiamingy): Remove this once we get rid of cub submodule.
diff --git a/src/common/hist_util.cu b/src/common/hist_util.cu
index 08ef98ea1..7e92433b9 100644
--- a/src/common/hist_util.cu
+++ b/src/common/hist_util.cu
@@ -19,7 +19,13 @@
 #include <vector>
 
 #include "categorical.h"
+
+#if defined(XGBOOST_USE_CUDA)
 #include "device_helpers.cuh"
+#elif defined(XGBOOST_USE_HIP)
+#include "device_helpers.hip.h"
+#endif
+
 #include "hist_util.cuh"
 #include "hist_util.h"
 #include "math.h"  // NOLINT
@@ -113,18 +119,35 @@ void SortByWeight(dh::device_vector<float>* weights,
                   dh::device_vector<Entry>* sorted_entries) {
   // Sort both entries and wegihts.
   dh::XGBDeviceAllocator<char> alloc;
+
+#if defined(XGBOOST_USE_CUDA)
   thrust::sort_by_key(thrust::cuda::par(alloc), sorted_entries->begin(),
                       sorted_entries->end(), weights->begin(),
                       detail::EntryCompareOp());
+#elif defined(XGBOOST_USE_HIP)
+  thrust::sort_by_key(thrust::hip::par(alloc), sorted_entries->begin(),
+                      sorted_entries->end(), weights->begin(),
+                      detail::EntryCompareOp());
+#endif
 
   // Scan weights
   dh::XGBCachingDeviceAllocator<char> caching;
+
+#if defined(XGBOOST_USE_CUDA)
   thrust::inclusive_scan_by_key(thrust::cuda::par(caching),
                                 sorted_entries->begin(), sorted_entries->end(),
                                 weights->begin(), weights->begin(),
                                 [=] __device__(const Entry& a, const Entry& b) {
                                   return a.index == b.index;
                                 });
+#elif defined(XGBOOST_USE_HIP)
+  thrust::inclusive_scan_by_key(thrust::hip::par(caching),
+                                sorted_entries->begin(), sorted_entries->end(),
+                                weights->begin(), weights->begin(),
+                                [=] __device__(const Entry& a, const Entry& b) {
+                                  return a.index == b.index;
+                                });
+#endif
 }
 
 void RemoveDuplicatedCategories(
@@ -192,8 +215,14 @@ void ProcessBatch(int device, MetaInfo const &info, const SparsePage &page,
     sorted_entries = dh::device_vector<Entry>(host_data.begin() + begin,
                                               host_data.begin() + end);
   }
+
+#if defined(XGBOOST_USE_CUDA)
   thrust::sort(thrust::cuda::par(alloc), sorted_entries.begin(),
                sorted_entries.end(), detail::EntryCompareOp());
+#elif defined(XGBOOST_USE_HIP)
+  thrust::sort(thrust::hip::par(alloc), sorted_entries.begin(),
+               sorted_entries.end(), detail::EntryCompareOp());
+#endif
 
   HostDeviceVector<SketchContainer::OffsetT> cuts_ptr;
   dh::caching_device_vector<size_t> column_sizes_scan;
diff --git a/src/common/hist_util.cuh b/src/common/hist_util.cuh
index ef179b4b0..a027d856f 100644
--- a/src/common/hist_util.cuh
+++ b/src/common/hist_util.cuh
@@ -89,7 +89,7 @@ void GetColumnSizesScan(int device, size_t num_columns, size_t num_cuts_per_feat
                          cuts_ptr->DevicePointer());
   thrust::exclusive_scan(thrust::hip::par(alloc), column_sizes_scan->begin(),
                          column_sizes_scan->end(), column_sizes_scan->begin());
-#else
+#elif defined(XGBOOST_USE_CUDA)
   thrust::exclusive_scan(thrust::cuda::par(alloc), cut_ptr_it,
                          cut_ptr_it + column_sizes_scan->size(),
                          cuts_ptr->DevicePointer());
@@ -198,7 +198,7 @@ void ProcessSlidingWindow(AdapterBatch const &batch, MetaInfo const &info,
 #if defined(XGBOOST_USE_HIP)
   thrust::sort(thrust::hip::par(alloc), sorted_entries.begin(),
                sorted_entries.end(), detail::EntryCompareOp());
-#else
+#elif defined(XGBOOST_USE_CUDA)
   thrust::sort(thrust::cuda::par(alloc), sorted_entries.begin(),
                sorted_entries.end(), detail::EntryCompareOp());
 #endif
@@ -229,7 +229,7 @@ void ProcessWeightedSlidingWindow(Batch batch, MetaInfo const& info,
 
 #if defined(XGBOOST_USE_HIP)
   dh::safe_cuda(hipSetDevice(device));
-#else
+#elif defined(XGBOOST_USE_CUDA)
   dh::safe_cuda(cudaSetDevice(device));
 #endif
 
@@ -272,7 +272,7 @@ void ProcessWeightedSlidingWindow(Batch batch, MetaInfo const& info,
                                  batch_iter + begin,
                                  d_temp_weights.data(),  // output
                                  is_valid);
-#else
+#elif defined(XGBOOST_USE_CUDA)
     auto retit = thrust::copy_if(thrust::cuda::par(alloc),
                                  weight_iter + begin, weight_iter + end,
                                  batch_iter + begin,
@@ -295,7 +295,7 @@ void ProcessWeightedSlidingWindow(Batch batch, MetaInfo const& info,
                                  batch_iter + begin,
                                  d_temp_weights.data(),  // output
                                  is_valid);
-#else
+#elif defined(XGBOOST_USE_CUDA)
     auto retit = thrust::copy_if(thrust::cuda::par(alloc),
                                  weight_iter + begin, weight_iter + end,
                                  batch_iter + begin,
diff --git a/src/common/hist_util.hip b/src/common/hist_util.hip
index e69de29bb..86eb989b3 100644
--- a/src/common/hist_util.hip
+++ b/src/common/hist_util.hip
@@ -0,0 +1,4 @@
+
+#if defined(XGBOOST_USE_HIP)
+#include "hist_util.cu"
+#endif