diff --git a/plugin/sycl/common/partition_builder.h b/plugin/sycl/common/partition_builder.h new file mode 100644 index 000000000000..b942e85441fb --- /dev/null +++ b/plugin/sycl/common/partition_builder.h @@ -0,0 +1,115 @@ +/*! + * 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 "../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) { + 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]); + } + prefix_sums_.Resize(qu, maxLocalSums); + } + + + xgboost::common::Span GetData(int nid) { + return { data_.Data() + nodes_offsets_[nid], nodes_offsets_[nid + 1] - nodes_offsets_[nid] }; + } + + + xgboost::common::Span GetPrefixSums() { + return { prefix_sums_.Data(), prefix_sums_.Size() }; + } + + + size_t GetLocalSize(const xgboost::common::Range1d& range) { + size_t range_size = range.end() - range.begin(); + size_t local_subgroups = range_size / (maxLocalSums * subgroupSize) + + !!(range_size % (maxLocalSums * subgroupSize)); + return subgroupSize * local_subgroups; + } + + + size_t GetSubgroupSize() { + return subgroupSize; + } + + size_t* GetResultRowsPtr() { + return result_rows_.data(); + } + + + size_t GetNLeftElems(int nid) const { + // return result_left_rows_[nid]; + return result_rows_[2 * nid]; + } + + + size_t GetNRightElems(int nid) const { + // return result_right_rows_[nid]; + return result_rows_[2 * nid + 1]; + } + + + ::sycl::event MergeToArray(::sycl::queue* qu, size_t node_in_set, + size_t* data_result, + ::sycl::event priv_event) { + size_t n_nodes_total = GetNLeftElems(node_in_set) + GetNRightElems(node_in_set); + if (n_nodes_total > 0) { + const size_t* data = data_.Data() + nodes_offsets_[node_in_set]; + return qu->memcpy(data_result, data, sizeof(size_t) * n_nodes_total, priv_event); + } else { + return ::sycl::event(); + } + } + + protected: + std::vector nodes_offsets_; + std::vector result_rows_; + size_t n_nodes_; + + USMVector data_; + USMVector prefix_sums_; +}; + +} // namespace common +} // namespace sycl +} // namespace xgboost + + +#endif // PLUGIN_SYCL_COMMON_PARTITION_BUILDER_H_ diff --git a/plugin/sycl/common/row_set.h b/plugin/sycl/common/row_set.h index dd05e9e59133..7ed9e4c5e531 100644 --- a/plugin/sycl/common/row_set.h +++ b/plugin/sycl/common/row_set.h @@ -120,121 +120,6 @@ class RowSetCollection { std::vector elem_of_each_node_; }; - -// 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]); - } - prefix_sums_.Resize(&qu, maxLocalSums); - } - - - xgboost::common::Span GetData(int nid) { - return { data_.Data() + nodes_offsets_[nid], nodes_offsets_[nid + 1] - nodes_offsets_[nid] }; - } - - - xgboost::common::Span GetPrefixSums() { - return { prefix_sums_.Data(), prefix_sums_.Size() }; - } - - - size_t GetLocalSize(const xgboost::common::Range1d& range) { - size_t range_size = range.end() - range.begin(); - size_t local_subgroups = range_size / (maxLocalSums * subgroupSize) + - !!(range_size % (maxLocalSums * subgroupSize)); - return subgroupSize * local_subgroups; - } - - - size_t GetSubgroupSize() { - return subgroupSize; - } - - // void SetNLeftElems(int nid, size_t n_left) { - // result_left_rows_[nid] = n_left; - // } - - - // void SetNRightElems(int nid, size_t n_right) { - // result_right_rows_[nid] = n_right; - // } - - // ::sycl::event SetNLeftRightElems(::sycl::queue& qu, const USMVector& parts_size, - // const std::vector<::sycl::event>& priv_events) { - // auto event = qu.submit([&](::sycl::handler& cgh) { - // cgh.depends_on(priv_events); - // cgh.parallel_for<>(::sycl::range<1>(n_nodes_), [=](::sycl::item<1> nid) { - // const size_t node_in_set = nid.get_id(0); - // result_left_rows_[node_in_set] = parts_size[2 * node_in_set]; - // result_right_rows_[node_in_set] = parts_size[2 * node_in_set + 1]; - // }); - // }); - // return event; - // } - - - size_t* GetResultRowsPtr() { - return result_rows_.data(); - } - - - size_t GetNLeftElems(int nid) const { - // return result_left_rows_[nid]; - return result_rows_[2 * nid]; - } - - - size_t GetNRightElems(int nid) const { - // return result_right_rows_[nid]; - return result_rows_[2 * nid + 1]; - } - - - ::sycl::event MergeToArray(::sycl::queue* qu, size_t node_in_set, - size_t* data_result, - ::sycl::event priv_event) { - size_t n_nodes_total = GetNLeftElems(node_in_set) + GetNRightElems(node_in_set); - if (n_nodes_total > 0) { - const size_t* data = data_.Data() + nodes_offsets_[node_in_set]; - return qu->memcpy(data_result, data, sizeof(size_t) * n_nodes_total, priv_event); - } else { - return ::sycl::event(); - } - } - - protected: - std::vector nodes_offsets_; - std::vector result_rows_; - size_t n_nodes_; - - USMVector data_; - USMVector prefix_sums_; - ::sycl::queue qu_; -}; - - } // namespace common } // namespace sycl } // namespace xgboost diff --git a/plugin/sycl/tree/updater_quantile_hist.cc b/plugin/sycl/tree/updater_quantile_hist.cc index 5a6cc96810b1..c6bb87c699ef 100644 --- a/plugin/sycl/tree/updater_quantile_hist.cc +++ b/plugin/sycl/tree/updater_quantile_hist.cc @@ -1363,7 +1363,7 @@ void QuantileHistMaker::Builder::ApplySplit( std::vector split_conditions; FindSplitConditions(nodes, *p_tree, gmat, &split_conditions); - partition_builder_.Init(qu_, n_nodes, [&](size_t node_in_set) { + partition_builder_.Init(&qu_, n_nodes, [&](size_t node_in_set) { const int32_t nid = nodes[node_in_set].nid; return row_set_collection_[nid].Size(); }); diff --git a/plugin/sycl/tree/updater_quantile_hist.h b/plugin/sycl/tree/updater_quantile_hist.h index d44503ea762f..8eb3b3d02c09 100644 --- a/plugin/sycl/tree/updater_quantile_hist.h +++ b/plugin/sycl/tree/updater_quantile_hist.h @@ -16,6 +16,7 @@ #include "../common/hist_util.h" #include "../common/row_set.h" +#include "../common/partition_builder.h" #include "split_evaluator.h" #include "../device_manager.h" diff --git a/tests/cpp/CMakeLists.txt b/tests/cpp/CMakeLists.txt index 2294ab75d4f3..f1a992785fe3 100644 --- a/tests/cpp/CMakeLists.txt +++ b/tests/cpp/CMakeLists.txt @@ -14,9 +14,39 @@ if (USE_CUDA) endif (USE_CUDA) file(GLOB_RECURSE SYCL_TEST_SOURCES "plugin/test_sycl_*.cc") -if (NOT PLUGIN_SYCL) - list(REMOVE_ITEM TEST_SOURCES ${SYCL_TEST_SOURCES}) -endif (NOT PLUGIN_SYCL) +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) target_include_directories(testxgboost PRIVATE ${xgboost_SOURCE_DIR}/plugin/federated) diff --git a/tests/cpp/plugin/test_sycl_multiclass_obj.cc b/tests/cpp/plugin/test_sycl_multiclass_obj.cc index fadfc6d41c96..ebfd8c41249e 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-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 "../objective/test_multiclass_obj.h" diff --git a/tests/cpp/plugin/test_sycl_predictor.cc b/tests/cpp/plugin/test_sycl_predictor.cc index 4bcb74b2b5ca..000e6b819e9d 100755 --- a/tests/cpp/plugin/test_sycl_predictor.cc +++ b/tests/cpp/plugin/test_sycl_predictor.cc @@ -2,9 +2,17 @@ * Copyright 2017-2020 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/gbm/gbtree.h" +#pragma GCC diagnostic pop #include "../../../src/data/proxy_dmatrix.h" #include "../../../src/gbm/gbtree_model.h" #include "../filesystem.h" // dmlc::TemporaryDirectory diff --git a/tests/cpp/plugin/test_sycl_regression_obj.cc b/tests/cpp/plugin/test_sycl_regression_obj.cc index 00041395f46f..6a40e7bf9e9c 100755 --- 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"