From 748a67a3c488b6d9e307a09442b1b3216e40f942 Mon Sep 17 00:00:00 2001 From: Arnav Prabhu Date: Fri, 14 Aug 2026 21:21:57 -0500 Subject: [PATCH 1/3] [SYCL] Restore the nd_range default constructor PR #22908 removed nd_range's default constructor to align with the SYCL 2020 specification, which does not define one. That broke code which relies on default-constructing an nd_range and assigning it later, most notably CuTe's launch_policy used by PyTorch's XPU flash-attention kernels. Restore the constructor as a backwards-compatibility extension and add a test that checks it zero-initializes the constituent ranges. Fixes #22940 --- sycl/include/sycl/nd_range.hpp | 6 ++++++ sycl/test/basic_tests/nd_range.cpp | 8 ++++++++ 2 files changed, 14 insertions(+) diff --git a/sycl/include/sycl/nd_range.hpp b/sycl/include/sycl/nd_range.hpp index 9d1a0f51408ce..3486112738248 100644 --- a/sycl/include/sycl/nd_range.hpp +++ b/sycl/include/sycl/nd_range.hpp @@ -62,6 +62,12 @@ template class nd_range { nd_range(nd_range &&rhs) = default; nd_range &operator=(const nd_range &rhs) = default; nd_range &operator=(nd_range &&rhs) = default; + + // nd_range has no default constructor in SYCL 2020, but one is kept as an + // extension for backwards compatibility with code that default-constructs + // and later assigns an nd_range. + nd_range() = default; + ~nd_range() = default; // Common hidden friend functions for by-value semantics diff --git a/sycl/test/basic_tests/nd_range.cpp b/sycl/test/basic_tests/nd_range.cpp index 39e1479f31481..3dc89c8276196 100644 --- a/sycl/test/basic_tests/nd_range.cpp +++ b/sycl/test/basic_tests/nd_range.cpp @@ -58,4 +58,12 @@ int main() { assert(three_dim_nd_range.get_group_range() == sycl::range<3>(2, 2, 2)); assert(three_dim_nd_range.get_offset() == sycl::id<3>(0, 0, 0)); cout << "three_dim_nd_range passed " << endl; + + // nd_range keeps a default constructor (not in SYCL 2020) for backwards + // compatibility. Check that it zero-initializes the constituent ranges. + sycl::nd_range<3> default_nd_range; + assert(default_nd_range.get_global_range() == sycl::range<3>(0, 0, 0)); + assert(default_nd_range.get_local_range() == sycl::range<3>(0, 0, 0)); + assert(default_nd_range.get_offset() == sycl::id<3>(0, 0, 0)); + cout << "default nd_range passed " << endl; } From 7c19c4de9b66b959859c29838b7c1d07b53c7e5c Mon Sep 17 00:00:00 2001 From: Arnav Prabhu Date: Mon, 17 Aug 2026 17:19:40 -0500 Subject: [PATCH 2/3] [SYCL] Address nd_range review feedback --- sycl/include/sycl/nd_range.hpp | 3 --- sycl/test/basic_tests/nd_range.cpp | 3 +-- 2 files changed, 1 insertion(+), 5 deletions(-) diff --git a/sycl/include/sycl/nd_range.hpp b/sycl/include/sycl/nd_range.hpp index 3486112738248..00993cb324e68 100644 --- a/sycl/include/sycl/nd_range.hpp +++ b/sycl/include/sycl/nd_range.hpp @@ -63,9 +63,6 @@ template class nd_range { nd_range &operator=(const nd_range &rhs) = default; nd_range &operator=(nd_range &&rhs) = default; - // nd_range has no default constructor in SYCL 2020, but one is kept as an - // extension for backwards compatibility with code that default-constructs - // and later assigns an nd_range. nd_range() = default; ~nd_range() = default; diff --git a/sycl/test/basic_tests/nd_range.cpp b/sycl/test/basic_tests/nd_range.cpp index 3dc89c8276196..6ba3fa0093cbc 100644 --- a/sycl/test/basic_tests/nd_range.cpp +++ b/sycl/test/basic_tests/nd_range.cpp @@ -59,8 +59,7 @@ int main() { assert(three_dim_nd_range.get_offset() == sycl::id<3>(0, 0, 0)); cout << "three_dim_nd_range passed " << endl; - // nd_range keeps a default constructor (not in SYCL 2020) for backwards - // compatibility. Check that it zero-initializes the constituent ranges. + // Check that a default-constructed nd_range has zeroed ranges and offset. sycl::nd_range<3> default_nd_range; assert(default_nd_range.get_global_range() == sycl::range<3>(0, 0, 0)); assert(default_nd_range.get_local_range() == sycl::range<3>(0, 0, 0)); From a0aa87b9edc71a9093850309c80d6829c6302342 Mon Sep 17 00:00:00 2001 From: "Kornev, Nikita" Date: Mon, 24 Aug 2026 13:19:51 +0200 Subject: [PATCH 3/3] add-guards --- sycl/include/sycl/nd_range.hpp | 2 ++ sycl/test/basic_tests/nd_range.cpp | 2 ++ 2 files changed, 4 insertions(+) diff --git a/sycl/include/sycl/nd_range.hpp b/sycl/include/sycl/nd_range.hpp index 00993cb324e68..04c9fcc99396d 100644 --- a/sycl/include/sycl/nd_range.hpp +++ b/sycl/include/sycl/nd_range.hpp @@ -63,7 +63,9 @@ template class nd_range { nd_range &operator=(const nd_range &rhs) = default; nd_range &operator=(nd_range &&rhs) = default; +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES nd_range() = default; +#endif // __INTEL_PREVIEW_BREAKING_CHANGES ~nd_range() = default; diff --git a/sycl/test/basic_tests/nd_range.cpp b/sycl/test/basic_tests/nd_range.cpp index 6ba3fa0093cbc..dcce69be4f86a 100644 --- a/sycl/test/basic_tests/nd_range.cpp +++ b/sycl/test/basic_tests/nd_range.cpp @@ -59,10 +59,12 @@ int main() { assert(three_dim_nd_range.get_offset() == sycl::id<3>(0, 0, 0)); cout << "three_dim_nd_range passed " << endl; +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES // Check that a default-constructed nd_range has zeroed ranges and offset. sycl::nd_range<3> default_nd_range; assert(default_nd_range.get_global_range() == sycl::range<3>(0, 0, 0)); assert(default_nd_range.get_local_range() == sycl::range<3>(0, 0, 0)); assert(default_nd_range.get_offset() == sycl::id<3>(0, 0, 0)); cout << "default nd_range passed " << endl; +#endif // __INTEL_PREVIEW_BREAKING_CHANGES }