From 5ea7624b24dea80e3bc5c5de56175441d40c25cc Mon Sep 17 00:00:00 2001 From: Dmitry Razdoburdin <> Date: Fri, 26 Jan 2024 06:47:33 -0800 Subject: [PATCH 1/3] Add the basics of partition builder implementation and the related test --- plugin/sycl/common/partition_builder.h | 110 ++++++++++++++++++ tests/cpp/CMakeLists.txt | 34 +++++- tests/cpp/plugin/test_sycl_multiclass_obj.cc | 4 + .../cpp/plugin/test_sycl_partition_builder.cc | 91 +++++++++++++++ tests/cpp/plugin/test_sycl_predictor.cc | 10 +- tests/cpp/plugin/test_sycl_regression_obj.cc | 4 + 6 files changed, 250 insertions(+), 3 deletions(-) create mode 100644 plugin/sycl/common/partition_builder.h create mode 100644 tests/cpp/plugin/test_sycl_partition_builder.cc diff --git a/plugin/sycl/common/partition_builder.h b/plugin/sycl/common/partition_builder.h new file mode 100644 index 000000000000..eb5d8f665a92 --- /dev/null +++ b/plugin/sycl/common/partition_builder.h @@ -0,0 +1,110 @@ +/*! + * Copyright 2017-2023 XGBoost contributors + */ +#ifndef PLUGIN_SYCL_COMMON_PARTITION_BUILDER_H_ +#define PLUGIN_SYCL_COMMON_PARTITION_BUILDER_H_ + +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wtautological-constant-compare" +#pragma GCC diagnostic ignored "-W#pragma-messages" +#include +#pragma GCC diagnostic pop +#include + +#include +#include +#include + +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wtautological-constant-compare" +#include "../../../src/common/column_matrix.h" +#pragma GCC diagnostic pop + +#include "../data.h" + +#include + +namespace xgboost { +namespace sycl { +namespace common { + +// The builder is required for samples partition to left and rights children for set of nodes +class PartitionBuilder { + public: + static constexpr size_t maxLocalSums = 256; + static constexpr size_t subgroupSize = 16; + + + template + void Init(::sycl::queue* qu, size_t n_nodes, Func funcNTaks) { + qu_ = qu; + nodes_offsets_.resize(n_nodes+1); + result_rows_.resize(2 * n_nodes); + n_nodes_ = n_nodes; + + + nodes_offsets_[0] = 0; + for (size_t i = 1; i < n_nodes+1; ++i) { + nodes_offsets_[i] = nodes_offsets_[i-1] + funcNTaks(i-1); + } + + if (data_.Size() < nodes_offsets_[n_nodes]) { + data_.Resize(qu, nodes_offsets_[n_nodes]); + } + } + + size_t GetSubgroupSize() { + return subgroupSize; + } + + + size_t GetNLeftElems(int nid) const { + return result_rows_[2 * nid]; + } + + + size_t GetNRightElems(int nid) const { + return result_rows_[2 * nid + 1]; + } + + void SetNLeftElems(int nid, size_t val) { + result_rows_[2 * nid] = val; + } + + + void SetNRightElems(int nid, size_t val) { + result_rows_[2 * nid + 1] = val; + } + + xgboost::common::Span GetData(int nid) { + return { data_.Data() + nodes_offsets_[nid], nodes_offsets_[nid + 1] - nodes_offsets_[nid] }; + } + + void MergeToArray(size_t nid, + size_t* data_result, + ::sycl::event* event) { + size_t n_nodes_total = GetNLeftElems(nid) + GetNRightElems(nid); + if (n_nodes_total > 0) { + const size_t* data = data_.Data() + nodes_offsets_[nid]; + qu_->memcpy(data_result, data, sizeof(size_t) * n_nodes_total, *event); + } + } + + protected: + std::vector nodes_offsets_; + std::vector result_rows_; + std::vector<::sycl::event> nodes_events_; + size_t n_nodes_; + + USMVector parts_size_; + USMVector data_; + + ::sycl::queue* qu_; +}; + +} // namespace common +} // namespace sycl +} // namespace xgboost + + +#endif // PLUGIN_SYCL_COMMON_PARTITION_BUILDER_H_ diff --git a/tests/cpp/CMakeLists.txt b/tests/cpp/CMakeLists.txt index 08862feee79a..20923519ac49 100644 --- a/tests/cpp/CMakeLists.txt +++ b/tests/cpp/CMakeLists.txt @@ -14,8 +14,38 @@ if(USE_CUDA) endif() file(GLOB_RECURSE SYCL_TEST_SOURCES "plugin/test_sycl_*.cc") -if(NOT PLUGIN_SYCL) - list(REMOVE_ITEM TEST_SOURCES ${SYCL_TEST_SOURCES}) +list(REMOVE_ITEM TEST_SOURCES ${SYCL_TEST_SOURCES}) + +if(PLUGIN_SYCL) + set(CMAKE_CXX_COMPILER "icpx") + file(GLOB_RECURSE SYCL_TEST_SOURCES "plugin/test_sycl_*.cc") + add_library(plugin_sycl_test OBJECT ${SYCL_TEST_SOURCES}) + + target_include_directories(plugin_sycl_test + PRIVATE + ${gtest_SOURCE_DIR}/include + ${xgboost_SOURCE_DIR}/include + ${xgboost_SOURCE_DIR}/dmlc-core/include + ${xgboost_SOURCE_DIR}/rabit/include) + + target_compile_definitions(plugin_sycl_test PUBLIC -DXGBOOST_USE_SYCL=1) + + target_link_libraries(plugin_sycl_test PUBLIC -fsycl) + + set_target_properties(plugin_sycl_test PROPERTIES + COMPILE_FLAGS -fsycl + CXX_STANDARD 17 + CXX_STANDARD_REQUIRED ON + POSITION_INDEPENDENT_CODE ON) + if(USE_OPENMP) + find_package(OpenMP REQUIRED) + set_target_properties(plugin_sycl_test PROPERTIES + COMPILE_FLAGS "-fsycl -qopenmp") + endif() + # Get compilation and link flags of plugin_sycl and propagate to testxgboost + target_link_libraries(testxgboost PUBLIC plugin_sycl_test) + # Add all objects of plugin_sycl to testxgboost + target_sources(testxgboost INTERFACE $) endif() if(PLUGIN_FEDERATED) diff --git a/tests/cpp/plugin/test_sycl_multiclass_obj.cc b/tests/cpp/plugin/test_sycl_multiclass_obj.cc index d809ecad3fc1..d306337ac599 100644 --- a/tests/cpp/plugin/test_sycl_multiclass_obj.cc +++ b/tests/cpp/plugin/test_sycl_multiclass_obj.cc @@ -2,7 +2,11 @@ * Copyright 2018-2023 XGBoost contributors */ #include +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wtautological-constant-compare" +#pragma GCC diagnostic ignored "-W#pragma-messages" #include +#pragma GCC diagnostic pop #include "../objective/test_multiclass_obj.h" diff --git a/tests/cpp/plugin/test_sycl_partition_builder.cc b/tests/cpp/plugin/test_sycl_partition_builder.cc new file mode 100644 index 000000000000..aaf5ba3114dc --- /dev/null +++ b/tests/cpp/plugin/test_sycl_partition_builder.cc @@ -0,0 +1,91 @@ +/** + * Copyright 2020-2023 by XGBoost contributors + */ +#include + +#include +#include +#include + +#include "../../../plugin/sycl/common/partition_builder.h" +#include "../../../plugin/sycl/device_manager.h" +#include "../helpers.h" + +namespace xgboost::sycl::common { + +TEST(SyclPartitionBuilder, BasicTest) { + constexpr size_t kNodes = 5; + // Number of rows for each node + std::vector rows = { 5, 5, 10, 1, 2 }; + + DeviceManager device_manager; + auto qu = device_manager.GetQueue(DeviceOrd::SyclDefault()); + PartitionBuilder builder; + builder.Init(&qu, kNodes, [&](size_t i) { + return rows[i]; + }); + + // We test here only the basics, thus syntetic partition builder is adopted + // Number of rows to go left for each node. + std::vector rows_for_left_node = { 2, 0, 7, 1, 2 }; + + size_t first_row_id = 0; + for(size_t nid = 0; nid < kNodes; ++nid) { + size_t n_rows_nodes = rows[nid]; + + auto rid_buff = builder.GetData(nid); + size_t rid_buff_size = rid_buff.size(); + auto* rid_buff_ptr = rid_buff.data(); + + size_t n_left = rows_for_left_node[nid]; + size_t n_right = rows[nid] - n_left; + + qu.submit([&](::sycl::handler& cgh) { + cgh.parallel_for<>(::sycl::range<1>(n_left), [=](::sycl::id<1> pid) { + int row_id = first_row_id + pid[0]; + rid_buff_ptr[pid[0]] = row_id; + }); + }); + qu.wait(); + first_row_id += n_left; + + // We are storing indexes for the right side in the tail of the array to save some memory + qu.submit([&](::sycl::handler& cgh) { + cgh.parallel_for<>(::sycl::range<1>(n_right), [=](::sycl::id<1> pid) { + int row_id = first_row_id + pid[0]; + rid_buff_ptr[rid_buff_size - pid[0] - 1] = row_id; + }); + }); + qu.wait(); + first_row_id += n_right; + + builder.SetNLeftElems(nid, n_left); + builder.SetNRightElems(nid, n_right); + } + + ::sycl::event event; + std::vector v(*std::max_element(rows.begin(), rows.end())); + size_t row_id = 0; + for(size_t nid = 0; nid < kNodes; ++nid) { + builder.MergeToArray(nid, v.data(), &event); + qu.wait(); + + // Check that row_id for left side are correct + for(size_t j = 0; j < rows_for_left_node[nid]; ++j) { + ASSERT_EQ(v[j], row_id++); + } + + // Check that row_id for right side are correct + for(size_t j = 0; j < rows[nid] - rows_for_left_node[nid]; ++j) { + ASSERT_EQ(v[rows[nid] - j - 1], row_id++); + } + + // Check that number of left/right rows are correct + size_t n_left = builder.GetNLeftElems(nid); + size_t n_right = builder.GetNRightElems(nid); + ASSERT_EQ(n_left, rows_for_left_node[nid]); + ASSERT_EQ(n_right, (rows[nid] - rows_for_left_node[nid])); + } +} + +} // namespace xgboost::common diff --git a/tests/cpp/plugin/test_sycl_predictor.cc b/tests/cpp/plugin/test_sycl_predictor.cc index f82a9f33d5f8..d5b3a5e5cd9a 100755 --- a/tests/cpp/plugin/test_sycl_predictor.cc +++ b/tests/cpp/plugin/test_sycl_predictor.cc @@ -2,11 +2,19 @@ * Copyright 2017-2023 XGBoost contributors */ #include +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wtautological-constant-compare" +#pragma GCC diagnostic ignored "-W#pragma-messages" #include +#pragma GCC diagnostic pop +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wtautological-constant-compare" #include "../../../src/data/adapter.h" -#include "../../../src/data/proxy_dmatrix.h" #include "../../../src/gbm/gbtree.h" +#pragma GCC diagnostic pop + +#include "../../../src/data/proxy_dmatrix.h" #include "../../../src/gbm/gbtree_model.h" #include "../filesystem.h" // dmlc::TemporaryDirectory #include "../helpers.h" diff --git a/tests/cpp/plugin/test_sycl_regression_obj.cc b/tests/cpp/plugin/test_sycl_regression_obj.cc index 66b4ea508477..349415390268 100644 --- a/tests/cpp/plugin/test_sycl_regression_obj.cc +++ b/tests/cpp/plugin/test_sycl_regression_obj.cc @@ -2,7 +2,11 @@ * Copyright 2017-2019 XGBoost contributors */ #include +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wtautological-constant-compare" +#pragma GCC diagnostic ignored "-W#pragma-messages" #include +#pragma GCC diagnostic pop #include #include "../helpers.h" From 4d700bbe9215608d8165aa16fe09e24eff35e530 Mon Sep 17 00:00:00 2001 From: Dmitry Razdoburdin <> Date: Fri, 26 Jan 2024 06:52:24 -0800 Subject: [PATCH 2/3] remove some dead code --- plugin/sycl/common/partition_builder.h | 17 ++++------------- tests/cpp/plugin/test_sycl_partition_builder.cc | 2 +- 2 files changed, 5 insertions(+), 14 deletions(-) diff --git a/plugin/sycl/common/partition_builder.h b/plugin/sycl/common/partition_builder.h index eb5d8f665a92..4d1e794a4561 100644 --- a/plugin/sycl/common/partition_builder.h +++ b/plugin/sycl/common/partition_builder.h @@ -31,10 +31,6 @@ namespace common { // The builder is required for samples partition to left and rights children for set of nodes class PartitionBuilder { public: - static constexpr size_t maxLocalSums = 256; - static constexpr size_t subgroupSize = 16; - - template void Init(::sycl::queue* qu, size_t n_nodes, Func funcNTaks) { qu_ = qu; @@ -53,11 +49,6 @@ class PartitionBuilder { } } - size_t GetSubgroupSize() { - return subgroupSize; - } - - size_t GetNLeftElems(int nid) const { return result_rows_[2 * nid]; } @@ -67,11 +58,12 @@ class PartitionBuilder { return result_rows_[2 * nid + 1]; } + // For test purposes only void SetNLeftElems(int nid, size_t val) { result_rows_[2 * nid] = val; } - + // For test purposes only void SetNRightElems(int nid, size_t val) { result_rows_[2 * nid + 1] = val; } @@ -82,18 +74,17 @@ class PartitionBuilder { void MergeToArray(size_t nid, size_t* data_result, - ::sycl::event* event) { + ::sycl::event event) { size_t n_nodes_total = GetNLeftElems(nid) + GetNRightElems(nid); if (n_nodes_total > 0) { const size_t* data = data_.Data() + nodes_offsets_[nid]; - qu_->memcpy(data_result, data, sizeof(size_t) * n_nodes_total, *event); + qu_->memcpy(data_result, data, sizeof(size_t) * n_nodes_total, event); } } protected: std::vector nodes_offsets_; std::vector result_rows_; - std::vector<::sycl::event> nodes_events_; size_t n_nodes_; USMVector parts_size_; diff --git a/tests/cpp/plugin/test_sycl_partition_builder.cc b/tests/cpp/plugin/test_sycl_partition_builder.cc index aaf5ba3114dc..db5ad1179e12 100644 --- a/tests/cpp/plugin/test_sycl_partition_builder.cc +++ b/tests/cpp/plugin/test_sycl_partition_builder.cc @@ -67,7 +67,7 @@ TEST(SyclPartitionBuilder, BasicTest) { std::vector v(*std::max_element(rows.begin(), rows.end())); size_t row_id = 0; for(size_t nid = 0; nid < kNodes; ++nid) { - builder.MergeToArray(nid, v.data(), &event); + builder.MergeToArray(nid, v.data(), event); qu.wait(); // Check that row_id for left side are correct From 1230d33a70a54861e12c34731002740e21bd99d4 Mon Sep 17 00:00:00 2001 From: Dmitry Razdoburdin <> Date: Fri, 26 Jan 2024 07:37:52 -0800 Subject: [PATCH 3/3] update copyright year --- plugin/sycl/common/partition_builder.h | 2 +- tests/cpp/plugin/test_sycl_partition_builder.cc | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/plugin/sycl/common/partition_builder.h b/plugin/sycl/common/partition_builder.h index 4d1e794a4561..37d1af241ab1 100644 --- a/plugin/sycl/common/partition_builder.h +++ b/plugin/sycl/common/partition_builder.h @@ -1,5 +1,5 @@ /*! - * Copyright 2017-2023 XGBoost contributors + * Copyright 2017-2024 XGBoost contributors */ #ifndef PLUGIN_SYCL_COMMON_PARTITION_BUILDER_H_ #define PLUGIN_SYCL_COMMON_PARTITION_BUILDER_H_ diff --git a/tests/cpp/plugin/test_sycl_partition_builder.cc b/tests/cpp/plugin/test_sycl_partition_builder.cc index db5ad1179e12..90bc757eb1b0 100644 --- a/tests/cpp/plugin/test_sycl_partition_builder.cc +++ b/tests/cpp/plugin/test_sycl_partition_builder.cc @@ -1,5 +1,5 @@ /** - * Copyright 2020-2023 by XGBoost contributors + * Copyright 2020-2024 by XGBoost contributors */ #include