diff --git a/src/linear/updater_gpu_coordinate.cu b/src/linear/updater_gpu_coordinate.cu
index b63c1317e..eb2ffd1ee 100644
--- a/src/linear/updater_gpu_coordinate.cu
+++ b/src/linear/updater_gpu_coordinate.cu
@@ -11,7 +11,13 @@
 
 #include "coordinate_common.h"
 #include "../common/common.h"
+
+#if defined(XGBOOST_USE_CUDA)
 #include "../common/device_helpers.cuh"
+#elif defined(XGBOOST_USE_HIP)
+#include "../common/device_helpers.hip.h"
+#endif
+
 #include "../common/timer.h"
 #include "./param.h"
 
@@ -60,7 +66,12 @@ class GPUCoordinateUpdater : public LinearUpdater {  // NOLINT
       return;
     }
 
+#if defined(XGBOOST_USE_CUDA)
     dh::safe_cuda(cudaSetDevice(ctx_->gpu_id));
+#elif defined(XGBOOST_USE_HIP)
+    dh::safe_cuda(hipSetDevice(ctx_->gpu_id));
+#endif
+
     // The begin and end indices for the section of each column associated with
     // this device
     std::vector<std::pair<bst_uint, bst_uint>> column_segments;
@@ -86,10 +97,18 @@ class GPUCoordinateUpdater : public LinearUpdater {  // NOLINT
     for (size_t fidx = 0; fidx < batch.Size(); fidx++) {
       auto col = page[fidx];
       auto seg = column_segments[fidx];
+
+#if defined(XGBOOST_USE_CUDA)
       dh::safe_cuda(cudaMemcpy(
           data_.data().get() + row_ptr_[fidx],
           col.data() + seg.first,
           sizeof(Entry) * (seg.second - seg.first), cudaMemcpyHostToDevice));
+#elif defined(XGBOOST_USE_HIP)
+      dh::safe_cuda(hipMemcpy(
+          data_.data().get() + row_ptr_[fidx],
+          col.data() + seg.first,
+          sizeof(Entry) * (seg.second - seg.first), hipMemcpyHostToDevice));
+#endif
     }
   }
 
@@ -170,7 +189,12 @@ class GPUCoordinateUpdater : public LinearUpdater {  // NOLINT
 
   // This needs to be public because of the __device__ lambda.
   GradientPair GetBiasGradient(int group_idx, int num_group) {
+#if defined(XGBOOST_USE_CUDA)
     dh::safe_cuda(cudaSetDevice(ctx_->gpu_id));
+#elif defined(XGBOOST_USE_HIP)
+    dh::safe_cuda(hipSetDevice(ctx_->gpu_id));
+#endif
+
     auto counting = thrust::make_counting_iterator(0ull);
     auto f = [=] __device__(size_t idx) {
       return idx * num_group + group_idx;
@@ -194,7 +218,12 @@ class GPUCoordinateUpdater : public LinearUpdater {  // NOLINT
 
   // This needs to be public because of the __device__ lambda.
   GradientPair GetGradient(int group_idx, int num_group, int fidx) {
+#if defined(XGBOOST_USE_CUDA)
     dh::safe_cuda(cudaSetDevice(ctx_->gpu_id));
+#elif defined(XGBOOST_USE_HIP)
+    dh::safe_cuda(hipSetDevice(ctx_->gpu_id));
+#endif
+
     common::Span<xgboost::Entry> d_col = dh::ToSpan(data_).subspan(row_ptr_[fidx]);
     size_t col_size = row_ptr_[fidx + 1] - row_ptr_[fidx];
     common::Span<GradientPair> d_gpair = dh::ToSpan(gpair_);
@@ -227,10 +256,17 @@ class GPUCoordinateUpdater : public LinearUpdater {  // NOLINT
   }
 
   void UpdateGpair(const std::vector<GradientPair> &host_gpair) {
+#if defined(XGBOOST_USE_CUDA)
     dh::safe_cuda(cudaMemcpyAsync(
         gpair_.data().get(),
         host_gpair.data(),
         gpair_.size() * sizeof(GradientPair), cudaMemcpyHostToDevice));
+#elif defined(XGBOOST_USE_HIP)
+    dh::safe_cuda(hipMemcpyAsync(
+        gpair_.data().get(),
+        host_gpair.data(),
+        gpair_.size() * sizeof(GradientPair), hipMemcpyHostToDevice));
+#endif
   }
 
   // training parameter
diff --git a/src/linear/updater_gpu_coordinate.hip b/src/linear/updater_gpu_coordinate.hip
index e69de29bb..b973a568f 100644
--- a/src/linear/updater_gpu_coordinate.hip
+++ b/src/linear/updater_gpu_coordinate.hip
@@ -0,0 +1,4 @@
+
+#if defined(XGBOOST_USE_HIP)
+#include "updater_gpu_coordinate.cu"
+#endif