From 9659d0e7bda73a55337118357a95e11fc4135eaa Mon Sep 17 00:00:00 2001
From: Hendrik Groove <hendrik.groove@dw.com>
Date: Mon, 21 Oct 2024 21:35:25 +0200
Subject: [PATCH] WARP_SIZE

---
 src/tree/gpu_hist/evaluate_splits.cu | 10 +---------
 1 file changed, 1 insertion(+), 9 deletions(-)

diff --git a/src/tree/gpu_hist/evaluate_splits.cu b/src/tree/gpu_hist/evaluate_splits.cu
index b49f389b2..efdeead67 100644
--- a/src/tree/gpu_hist/evaluate_splits.cu
+++ b/src/tree/gpu_hist/evaluate_splits.cu
@@ -21,7 +21,7 @@
 #define WAVEFRONT_SIZE __AMDGCN_WAVEFRONT_SIZE
 #endif
 
-#define WARP_SIZE WAVEFRONT_SIZE
+#define WARP_SIZE 32
 #endif
 
 namespace xgboost::tree {
@@ -332,14 +332,6 @@ __global__ __launch_bounds__(kBlockSize) void EvaluateSplitsKernel(
   }
 }
 
-template __global__ void EvaluateSplitsKernel<64>(
-    bst_feature_t max_active_features,
-    common::Span<const EvaluateSplitInputs> d_inputs,
-    const EvaluateSplitSharedInputs shared_inputs,
-    common::Span<bst_feature_t> sorted_idx,
-    const TreeEvaluator::SplitEvaluator<GPUTrainingParam> evaluator,
-    common::Span<DeviceSplitCandidate> out_candidates);
-
 __device__ DeviceSplitCandidate operator+(const DeviceSplitCandidate &a,
                                           const DeviceSplitCandidate &b) {
   return b.loss_chg > a.loss_chg ? b : a;