Hendrik/xgboost
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

WARP_SIZE

Browse Source
This commit is contained in:
Hendrik Groove 2024-10-21 21:35:25 +02:00
parent 768c8b298c
commit 9659d0e7bd
1 changed files with 1 additions and 9 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
src/tree/gpu_hist/evaluate_splits.cu
View File

@ -21,7 +21,7 @@
#define WAVEFRONT_SIZE __AMDGCN_WAVEFRONT_SIZE #define WAVEFRONT_SIZE __AMDGCN_WAVEFRONT_SIZE
#endif #endif
#define WARP_SIZE WAVEFRONT_SIZE #define WARP_SIZE 32
#endif #endif
namespace xgboost::tree { 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, __device__ DeviceSplitCandidate operator+(const DeviceSplitCandidate &a,
const DeviceSplitCandidate &b) { const DeviceSplitCandidate &b) {
return b.loss_chg > a.loss_chg ? b : a; return b.loss_chg > a.loss_chg ? b : a;