[ROCm] re-add support for ROCm builds

Previously #6086 added ROCm support but after numerous rebases it lost
critical changes. This PR restores the ROCm build.

There are many source file changes but most were automated using the
following:

```bash
for f in `grep -rl '#ifdef USE_CUDA'`
do
    sed -i 's@#ifdef USE_CUDA@#if defined(USE_CUDA) || defined(USE_ROCM)@g' $f
done

for f in `grep -rl '#endif  // USE_CUDA'`
do
    sed -i 's@#endif  // USE_CUDA@#endif  // USE_CUDA || USE_ROCM@g' $f
done
```

src/objective/cuda/cuda_regression_objective.cpp

@@ -4,7 +4,7 @@
  * license information.
  */
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include "cuda_regression_objective.hpp"
 
@@ -106,4 +106,4 @@ void CUDARegressionQuantileloss::Init(const Metadata& metadata, data_size_t num_
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM

src/objective/cuda/cuda_regression_objective.cu

@@ -4,7 +4,7 @@
  * license information.
  */
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include "cuda_regression_objective.hpp"
 #include <LightGBM/cuda/cuda_algorithms.hpp>
@@ -478,4 +478,4 @@ void CUDARegressionQuantileloss::LaunchGetGradientsKernel(const double* score, s
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM

src/objective/cuda/cuda_regression_objective.hpp

@@ -7,7 +7,7 @@
 #ifndef LIGHTGBM_OBJECTIVE_CUDA_CUDA_REGRESSION_OBJECTIVE_HPP_
 #define LIGHTGBM_OBJECTIVE_CUDA_CUDA_REGRESSION_OBJECTIVE_HPP_
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #define GET_GRADIENTS_BLOCK_SIZE_REGRESSION (1024)
 
@@ -163,5 +163,5 @@ class CUDARegressionQuantileloss : public CUDARegressionObjectiveInterface<Regre
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM
 #endif  // LIGHTGBM_OBJECTIVE_CUDA_CUDA_REGRESSION_OBJECTIVE_HPP_

src/objective/objective_function.cpp

@@ -18,7 +18,7 @@
 namespace LightGBM {
 
 ObjectiveFunction* ObjectiveFunction::CreateObjectiveFunction(const std::string& type, const Config& config) {
-  #ifdef USE_CUDA
+  #if defined(USE_CUDA) || defined(USE_ROCM)
   if (config.device_type == std::string("cuda") &&
       config.data_sample_strategy != std::string("goss") &&
       config.boosting != std::string("rf")) {
@@ -64,7 +64,7 @@ ObjectiveFunction* ObjectiveFunction::CreateObjectiveFunction(const std::string&
       return nullptr;
     }
   } else {
-  #endif  // USE_CUDA
+  #endif  // USE_CUDA || USE_ROCM
     if (type == std::string("regression")) {
       return new RegressionL2loss(config);
     } else if (type == std::string("regression_l1")) {
@@ -100,9 +100,9 @@ ObjectiveFunction* ObjectiveFunction::CreateObjectiveFunction(const std::string&
     } else if (type == std::string("custom")) {
       return nullptr;
     }
-  #ifdef USE_CUDA
+  #if defined(USE_CUDA) || defined(USE_ROCM)
   }
-  #endif  // USE_CUDA
+  #endif  // USE_CUDA || USE_ROCM
   Log::Fatal("Unknown objective type name: %s", type.c_str());
   return nullptr;
 }

src/treelearner/cuda/cuda_best_split_finder.cpp

@@ -4,7 +4,7 @@
  * license information.
  */
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include <algorithm>
 
@@ -396,4 +396,4 @@ void CUDABestSplitFinder::SetUsedFeatureByNode(const std::vector<int8_t>& is_fea
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM

src/treelearner/cuda/cuda_best_split_finder.cu

@@ -5,7 +5,7 @@
  * Modifications Copyright(C) 2023 Advanced Micro Devices, Inc. All rights reserved.
  */
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include "cuda_best_split_finder.hpp"
 
@@ -933,7 +933,11 @@ __global__ void FindBestSplitsDiscretizedForLeafKernel(
   if (is_feature_used_bytree[inner_feature_index]) {
     if (task->is_categorical) {
       __threadfence();  // ensure store issued before trap
+#if defined(USE_CUDA)
       asm("trap;");
+#elif defined(USE_ROCM)
+      __builtin_trap();
+#endif
     } else {
       if (!task->reverse) {
         if (use_16bit_bin) {
@@ -2239,4 +2243,4 @@ void CUDABestSplitFinder::LaunchInitCUDARandomKernel() {
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM

src/treelearner/cuda/cuda_best_split_finder.hpp

@@ -7,7 +7,7 @@
 #ifndef LIGHTGBM_TREELEARNER_CUDA_CUDA_BEST_SPLIT_FINDER_HPP_
 #define LIGHTGBM_TREELEARNER_CUDA_CUDA_BEST_SPLIT_FINDER_HPP_
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include <LightGBM/bin.h>
 #include <LightGBM/dataset.h>
@@ -240,5 +240,5 @@ class CUDABestSplitFinder {
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM
 #endif  // LIGHTGBM_TREELEARNER_CUDA_CUDA_BEST_SPLIT_FINDER_HPP_

src/treelearner/cuda/cuda_data_partition.cpp

@@ -4,7 +4,7 @@
  * license information.
  */
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include <algorithm>
 #include <memory>
@@ -376,4 +376,4 @@ void CUDADataPartition::ReduceLeafGradStat(
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM

src/treelearner/cuda/cuda_data_partition.cu

@@ -5,7 +5,7 @@
  * Modifications Copyright(C) 2023 Advanced Micro Devices, Inc. All rights reserved.
  */
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include "cuda_data_partition.hpp"
 
@@ -1120,4 +1120,4 @@ void CUDADataPartition::LaunchReduceLeafGradStat(
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM

src/treelearner/cuda/cuda_data_partition.hpp

@@ -6,7 +6,7 @@
 #ifndef LIGHTGBM_TREELEARNER_CUDA_CUDA_DATA_PARTITION_HPP_
 #define LIGHTGBM_TREELEARNER_CUDA_CUDA_DATA_PARTITION_HPP_
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include <LightGBM/bin.h>
 #include <LightGBM/meta.h>
@@ -392,5 +392,5 @@ class CUDADataPartition {
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM
 #endif  // LIGHTGBM_TREELEARNER_CUDA_CUDA_DATA_PARTITION_HPP_

src/treelearner/cuda/cuda_gradient_discretizer.cu

@@ -4,7 +4,7 @@
  * license information.
  */
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include <algorithm>
 
@@ -168,4 +168,4 @@ void CUDAGradientDiscretizer::DiscretizeGradients(
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM

src/treelearner/cuda/cuda_gradient_discretizer.hpp

@@ -7,7 +7,7 @@
 #ifndef LIGHTGBM_TREELEARNER_CUDA_CUDA_GRADIENT_DISCRETIZER_HPP_
 #define LIGHTGBM_TREELEARNER_CUDA_CUDA_GRADIENT_DISCRETIZER_HPP_
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include <LightGBM/bin.h>
 #include <LightGBM/meta.h>
@@ -114,5 +114,5 @@ class CUDAGradientDiscretizer: public GradientDiscretizer {
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM
 #endif  // LIGHTGBM_TREELEARNER_CUDA_CUDA_GRADIENT_DISCRETIZER_HPP_

src/treelearner/cuda/cuda_histogram_constructor.cpp

@@ -4,7 +4,7 @@
  * license information.
  */
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include "cuda_histogram_constructor.hpp"
 
@@ -185,4 +185,4 @@ void CUDAHistogramConstructor::ResetConfig(const Config* config) {
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM

src/treelearner/cuda/cuda_histogram_constructor.cu

@@ -5,7 +5,7 @@
  * Modifications Copyright(C) 2023 Advanced Micro Devices, Inc. All rights reserved.
  */
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include "cuda_histogram_constructor.hpp"
 
@@ -959,4 +959,4 @@ void CUDAHistogramConstructor::LaunchSubtractHistogramKernel(
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM

src/treelearner/cuda/cuda_histogram_constructor.hpp

@@ -6,7 +6,7 @@
 #ifndef LIGHTGBM_TREELEARNER_CUDA_CUDA_HISTOGRAM_CONSTRUCTOR_HPP_
 #define LIGHTGBM_TREELEARNER_CUDA_CUDA_HISTOGRAM_CONSTRUCTOR_HPP_
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include <LightGBM/cuda/cuda_row_data.hpp>
 #include <LightGBM/cuda/cuda_utils.hu>
@@ -192,5 +192,5 @@ class CUDAHistogramConstructor {
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM
 #endif  // LIGHTGBM_TREELEARNER_CUDA_CUDA_HISTOGRAM_CONSTRUCTOR_HPP_

src/treelearner/cuda/cuda_leaf_splits.cpp

@@ -4,7 +4,7 @@
  * license information.
  */
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include "cuda_leaf_splits.hpp"
 
@@ -75,4 +75,4 @@ void CUDALeafSplits::Resize(const data_size_t num_data) {
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM

src/treelearner/cuda/cuda_leaf_splits.cu

@@ -6,7 +6,7 @@
  */
 
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include "cuda_leaf_splits.hpp"
 #include <LightGBM/cuda/cuda_algorithms.hpp>
@@ -244,4 +244,4 @@ void CUDALeafSplits::LaunchInitValuesKernel(
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM

src/treelearner/cuda/cuda_leaf_splits.hpp

@@ -6,7 +6,7 @@
 #ifndef LIGHTGBM_TREELEARNER_CUDA_CUDA_LEAF_SPLITS_HPP_
 #define LIGHTGBM_TREELEARNER_CUDA_CUDA_LEAF_SPLITS_HPP_
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include <LightGBM/cuda/cuda_utils.hu>
 #include <LightGBM/bin.h>
@@ -175,5 +175,5 @@ class CUDALeafSplits {
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM
 #endif  // LIGHTGBM_TREELEARNER_CUDA_CUDA_LEAF_SPLITS_HPP_

src/treelearner/cuda/cuda_single_gpu_tree_learner.cpp

@@ -4,7 +4,7 @@
  * license information.
  */
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include "cuda_single_gpu_tree_learner.hpp"
 
@@ -609,4 +609,4 @@ void CUDASingleGPUTreeLearner::RenewDiscretizedTreeLeaves(CUDATree* cuda_tree) {
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM

src/treelearner/cuda/cuda_single_gpu_tree_learner.cu

@@ -5,7 +5,7 @@
  * Modifications Copyright(C) 2023 Advanced Micro Devices, Inc. All rights reserved.
  */
 
-#ifdef USE_CUDA
+#if defined(USE_CUDA) || defined(USE_ROCM)
 
 #include "cuda_single_gpu_tree_learner.hpp"
 
@@ -292,4 +292,4 @@ void CUDASingleGPUTreeLearner::LaunchCalcLeafValuesGivenGradStat(
 
 }  // namespace LightGBM
 
-#endif  // USE_CUDA
+#endif  // USE_CUDA || USE_ROCM