Adding SYGraph primitives to oneDAL#3482
Open
olegkkruglov wants to merge 61 commits into
Open
Conversation
olegkkruglov
force-pushed
the
sygraph-primitives
branch
from
2c429d3 to
6a75234
Compare
olegkkruglov
requested review from
Alexandr-Solovev,
ahuber21,
avolkov-intel,
david-cortes-intel,
ethanglaser,
homksei,
icfaust and
napetrov
as code owners
olegkkruglov
force-pushed
the
sygraph-primitives
branch
from
0d63b37 to
1b40b33
Compare
Contributor
There was a problem hiding this comment.
Pull request overview
This PR extends oneDAL's backend primitives with SYGraph-oriented frontier support for GPU graph traversal, adding the core bitmap-frontier/advance/graph pieces plus tests and build integration for the new backend module.
Changes:
- Adds new frontier primitives: bitmap storage, frontier container/view, CSR graph wrapper, and header-only advance kernel support.
- Adds DPC++ tests for frontier basics, advance behavior, BFS behavior, and graph-generation helpers.
- Registers the new primitive in Bazel/public headers and updates shared backend/license-checker support.
Reviewed changes
Copilot reviewed 14 out of 14 changed files in this pull request and generated 14 comments.
Show a summary per file
| File | Description |
|---|---|
cpp/oneapi/dal/backend/primitives/frontier/test/utils.hpp |
Test helper utilities for device info and random CSR graph generation. |
cpp/oneapi/dal/backend/primitives/frontier/test/frontier_dpc.cpp |
Basic DPC++ unit test for frontier insert/check/offset behavior. |
cpp/oneapi/dal/backend/primitives/frontier/test/bfs_dpc.cpp |
End-to-end BFS-style test built on the new primitives. |
cpp/oneapi/dal/backend/primitives/frontier/test/advance_dpc.cpp |
DPC++ unit test for the advance primitive. |
cpp/oneapi/dal/backend/primitives/frontier/graph.hpp |
New CSR graph view/container used by frontier primitives. |
cpp/oneapi/dal/backend/primitives/frontier/frontier.hpp |
Frontier public interface and device-view definitions. |
cpp/oneapi/dal/backend/primitives/frontier/frontier_dpc.cpp |
DPC++ implementation of frontier operations and active-frontier computation. |
cpp/oneapi/dal/backend/primitives/frontier/BUILD |
Bazel module and test registration for the new frontier package. |
cpp/oneapi/dal/backend/primitives/frontier/bitset.hpp |
Low-level bitmap helper used by the frontier implementation. |
cpp/oneapi/dal/backend/primitives/frontier/advance.hpp |
Header-only SYCL advance kernel and launch logic. |
cpp/oneapi/dal/backend/primitives/frontier.hpp |
Umbrella include for the new frontier primitive APIs. |
cpp/oneapi/dal/backend/primitives/BUILD |
Adds frontier to aggregated primitive modules/tests. |
cpp/oneapi/dal/backend/common.hpp |
Adds a backend helper for querying max subgroup count. |
.github/.licenserc.yaml |
Updates license-header matching rules for new copyright lines. |
Comment on lines
+20
to
+24
| #include "oneapi/dal/common.hpp" | ||
| #include "oneapi/dal/graph/detail/common.hpp" | ||
| #include "oneapi/dal/graph/detail/container.hpp" | ||
| #include "oneapi/dal/backend/primitives/ndarray.hpp" | ||
|
|
Comment on lines
+72
to
+75
| csr_graph_view(std::uint64_t num_nodes, | ||
| vertex_t* row_ptr, | ||
| edge_t* col_indices, | ||
| weight_t* weights) |
Comment on lines
+121
to
+125
| csr_graph(sycl::queue& queue, | ||
| std::vector<VertexT> row_ptr, | ||
| std::vector<EdgeT> col_indices, | ||
| std::vector<WeightT> weights = {}, | ||
| sycl::usm::alloc alloc = sycl::usm::alloc::shared) |
| _mlb_layer.get_mutable_data(), | ||
| offsets_pointer, | ||
| offsets_size_pointer, | ||
| static_cast<std::uint64_t>(_data_layer.get_count()) }; |
| } | ||
|
|
||
| inline pr::ndview<std::uint32_t, 1> get_offsets() const { | ||
| return _offsets.slice(1, _offsets.get_count()); |
Comment on lines
+85
to
+87
| SYCL_EXTERNAL inline std::uint32_t get_degree(const vertex_t vertex) const { | ||
| ONEDAL_ASSERT(vertex < _num_nodes, "Vertex index out of bounds"); | ||
| return static_cast<std::uint32_t>(_row_ptr[vertex + 1] - _row_ptr[vertex]); |
Comment on lines
+59
to
+67
| auto sum_reduction = sycl::reduction(empty_buff_ptr, sycl::plus<>()); | ||
| auto* const f_ptr = _mlb_layer.get_mutable_data(); | ||
|
|
||
| cgh.parallel_for(range, sum_reduction, [=](sycl::id<1> idx, auto& sum_v) { | ||
| sum_v += f_ptr[idx]; | ||
| }); | ||
| }); | ||
| auto empty_sum = empty_buff.at_device(_queue, 0, { e }); | ||
| return empty_sum == 0; |
| using std::swap; | ||
| swap(f1._data_layer, f2._data_layer); | ||
| swap(f1._mlb_layer, f2._mlb_layer); | ||
| swap(f1._offsets, f2._offsets); |
Comment on lines
+312
to
+314
| auto offset_size = in.get_offsets_size().at_device(q, 0, { to_wait }); | ||
| global_size = offset_size * element_bitsize; | ||
|
|
| const sycl::local_accessor<T, 1> workgroup_reduce; | ||
| const sycl::local_accessor<uint32_t, 1> workgroup_reduce_tail; | ||
| const sycl::local_accessor<uint32_t, 1> workgroup_ids; | ||
| const LambdaT functor; |
olegkkruglov
force-pushed
the
sygraph-primitives
branch
2 times, most recently
from
a034a96 to
8947a09
Compare
…move obsolete test header
…eck and clear operations
…consistency; update related functionality; fixed bug for SIGSEGV
…name printing and frontier checks
- Removed the existing frontier_dpc.hpp file to streamline the codebase. - Introduced new test files for advance operation, BFS, and basic frontier operations. - Implemented comprehensive tests to validate the functionality of the frontier data structure. - Enhanced the frontier class with additional methods for better performance and usability. - Ensured compatibility with SYCL and improved device memory management.
…y frontier_context API
This was referenced May 29, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
This PR is created to finalize merging of #3289. Contains all content of #3289 and fixes to make CI green.
Currently LinuxMakeDPCPP(avx2) is red because of terminated build command. Tries to reproduce any issues locally were not successful.
Checklist:
Completeness and readability
Testing
Performance