From 5148ff78d3d03334a7cd120bf7eea97498ec9f67 Mon Sep 17 00:00:00 2001 From: Jiaming Yuan Date: Sat, 11 Feb 2023 02:04:07 +0800 Subject: [PATCH] Rename ranking utils to threading utils. These utilities were written for ranking but are used for other places like AUC. Also, they are about assigning CUDA threads instead of being ranking specific. --- ...{ranking_utils.cuh => threading_utils.cuh} | 40 +++++++++++-------- src/metric/auc.cu | 4 +- ...nking_utils.cu => test_threading_utils.cu} | 7 ++-- 3 files changed, 29 insertions(+), 22 deletions(-) rename src/common/{ranking_utils.cuh => threading_utils.cuh} (68%) rename tests/cpp/common/{test_ranking_utils.cu => test_threading_utils.cu} (92%) diff --git a/src/common/ranking_utils.cuh b/src/common/threading_utils.cuh similarity index 68% rename from src/common/ranking_utils.cuh rename to src/common/threading_utils.cuh index f63e38cba58c..c21d312d2e03 100644 --- a/src/common/ranking_utils.cuh +++ b/src/common/threading_utils.cuh @@ -1,13 +1,17 @@ -/*! - * Copyright 2021 by XGBoost Contributors +/** + * Copyright 2021-2023 by XGBoost Contributors */ -#ifndef XGBOOST_COMMON_RANKING_UTILS_H_ -#define XGBOOST_COMMON_RANKING_UTILS_H_ +#ifndef XGBOOST_COMMON_THREADING_UTILS_CUH_ +#define XGBOOST_COMMON_THREADING_UTILS_CUH_ + +#include // std::min +#include // std::size_t -#include -#include "xgboost/base.h" -#include "device_helpers.cuh" -#include "./math.h" +#include "./math.h" // Sqr +#include "common.h" +#include "device_helpers.cuh" // LaunchN +#include "xgboost/base.h" // XGBOOST_DEVICE +#include "xgboost/span.h" // Span namespace xgboost { namespace common { @@ -15,10 +19,10 @@ namespace common { * \param n Number of items (length of the base) * \param h hight */ -XGBOOST_DEVICE inline size_t DiscreteTrapezoidArea(size_t n, size_t h) { - n -= 1; // without diagonal entries - h = std::min(n, h); // Specific for ranking. - size_t total = ((n - (h - 1)) + n) * h / 2; +XGBOOST_DEVICE inline std::size_t DiscreteTrapezoidArea(std::size_t n, std::size_t h) { + n -= 1; // without diagonal entries + h = std::min(n, h); // Used for ranking, h <= n + std::size_t total = ((n - (h - 1)) + n) * h / 2; return total; } @@ -29,12 +33,14 @@ XGBOOST_DEVICE inline size_t DiscreteTrapezoidArea(size_t n, size_t h) { * Equivalent to loops like: * * \code - * for (size i = 0; i < h; ++i) { - * for (size_t j = i + 1; j < n; ++j) { + * for (std::size_t i = 0; i < h; ++i) { + * for (std::size_t j = i + 1; j < n; ++j) { * do_something(); * } * } * \endcode + * + * with h <= n */ template inline size_t @@ -79,6 +85,6 @@ XGBOOST_DEVICE inline void UnravelTrapeziodIdx(size_t i_idx, size_t n, j = idx - n_elems + i + 1; } -} // namespace common -} // namespace xgboost -#endif // XGBOOST_COMMON_RANKING_UTILS_H_ +} // namespace common +} // namespace xgboost +#endif // XGBOOST_COMMON_THREADING_UTILS_CUH_ diff --git a/src/metric/auc.cu b/src/metric/auc.cu index 7673e7b29d55..788d9a5698a5 100644 --- a/src/metric/auc.cu +++ b/src/metric/auc.cu @@ -12,8 +12,8 @@ #include #include "../collective/device_communicator.cuh" -#include "../common/optional_weight.h" // OptionalWeights -#include "../common/ranking_utils.cuh" +#include "../common/optional_weight.h" // OptionalWeights +#include "../common/threading_utils.cuh" // UnravelTrapeziodIdx,SegmentedTrapezoidThreads #include "auc.h" #include "xgboost/data.h" #include "xgboost/span.h" diff --git a/tests/cpp/common/test_ranking_utils.cu b/tests/cpp/common/test_threading_utils.cu similarity index 92% rename from tests/cpp/common/test_ranking_utils.cu rename to tests/cpp/common/test_threading_utils.cu index 8d240e41ff77..f7160b1b56f9 100644 --- a/tests/cpp/common/test_ranking_utils.cu +++ b/tests/cpp/common/test_threading_utils.cu @@ -1,13 +1,14 @@ /** - * Copyright 2021 by XGBoost Contributors + * Copyright 2021-2023 by XGBoost Contributors */ #include -#include "../../../src/common/ranking_utils.cuh" +#include // thrust::copy + #include "../../../src/common/device_helpers.cuh" +#include "../../../src/common/threading_utils.cuh" namespace xgboost { namespace common { - TEST(SegmentedTrapezoidThreads, Basic) { size_t constexpr kElements = 24, kGroups = 3; dh::device_vector offset_ptr(kGroups + 1, 0);