Skip to content

Add a workaround to custom_inclusive_scan_over_group #2275

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
ndgrigorian merged 3 commits into from
Apr 9, 2026
Merged

Add a workaround to custom_inclusive_scan_over_group #2275

Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
4326d02
Add a w/a to custom_inclusive_scan_over_group
antonwolfy Apr 9, 2026
b5f77f4
Apply pre-commit formatting rule
antonwolfy Apr 9, 2026
b495d8e
Use relaxed memory order
antonwolfy Apr 9, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
18 changes: 15 additions & 3 deletions dpctl/tensor/libtensor/include/utils/sycl_utils.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -268,9 +268,21 @@ T custom_inclusive_scan_over_group(GroupT &&wg,
const bool in_range = (lane_id < n_aggregates);
const bool in_bounds = in_range && (lane_id > 0 || large_wg);

T __scan_val = (in_bounds)
? local_mem_acc[(offset + lane_id) * max_sgSize - 1]
: identity;
// Here is a bug where IGC incorrectly optimized the below code:
// T __scan_val = (in_bounds)
// ? local_mem_acc[(offset + lane_id) * max_sgSize - 1]
// : identity;
// That causes `__scan_val` is not initialized with `identity` value:
// wgs = 256, max_sgSize = 16 => n_aggregates = 16
// wi = 0: in_range = 1, in_bounds = 0 => __scan_val = identity
// The w/s adds SYCL atomic fence, since the explicit memory fence
// prevents reordering/elimination, while it will add slight overhead.
T __scan_val = identity;
sycl::atomic_fence(sycl::memory_order::relaxed,
sycl::memory_scope::work_item);
if (in_bounds) {
__scan_val = local_mem_acc[(offset + lane_id) * max_sgSize - 1];
}
for (std::uint32_t step = 1; step < sgSize; step *= 2) {
const bool advanced_lane = (lane_id >= step);
const std::uint32_t src_lane_id =
Expand Down
Loading