From e866e4e7542d8da0f64d81f49c0bf3ca065253dd Mon Sep 17 00:00:00 2001 From: Artur Gainullin Date: Wed, 25 Mar 2026 19:17:03 +0100 Subject: [PATCH] [SYCL] Use unsigned int limit for -fsycl-id-queries-fit-in-int Change the fit-in-int optimization to use UINT_MAX instead of INT_MAX as the limit for SYCL ID queries. This allows the full unsigned range for work-item IDs, which are non-negative. I.e. this optimization is based on the assumption that values will fit in i32, but since these values are unsigned there is no reason to check for signed integer limits. --- sycl/doc/UsersManual.md | 4 ++-- .../proposed/sycl_ext_oneapi_range_type.asciidoc | 4 ++-- sycl/include/sycl/detail/defines.hpp | 2 +- sycl/include/sycl/detail/id_queries_fit_in_int.hpp | 10 +++++----- sycl/test-e2e/Basic/range_offset_fit_in_int.cpp | 11 +++++++---- 5 files changed, 17 insertions(+), 14 deletions(-) diff --git a/sycl/doc/UsersManual.md b/sycl/doc/UsersManual.md index 2829d51f0797a..fffd7b6530f1f 100644 --- a/sycl/doc/UsersManual.md +++ b/sycl/doc/UsersManual.md @@ -183,8 +183,8 @@ and not recommended to use in production environment. **`-f[no-]sycl-id-queries-fit-in-int`** - Assume/Do not assume that SYCL ID queries fit within MAX_INT. It assumes - that these values fit within MAX_INT: + Assume/Do not assume that SYCL ID queries fit within MAX_UINT. It assumes + that these values fit within MAX_UINT: * id class get() member function and operator[] * item class get_id() member function and operator[] * nd_item class get_global_id()/get_global_linear_id() member functions diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_range_type.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_range_type.asciidoc index 1105d81aa6906..ec3e38b73a124 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_range_type.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_range_type.asciidoc @@ -178,8 +178,8 @@ implementations that may support this extension. If a translation unit is compiled with the `-fsycl-id-queries-fit-in-int` option, all kernels and `SYCL_EXTERNAL` functions without an explicitly -specified `range_type` property are compiled as-if `range_type` was -specified as a property of that kernel or function. +specified `range_type` property are compiled as-if `range_type` +was specified as a property of that kernel or function. == Implementation notes diff --git a/sycl/include/sycl/detail/defines.hpp b/sycl/include/sycl/detail/defines.hpp index 8b10a8a6d1017..62387c879e73d 100644 --- a/sycl/include/sycl/detail/defines.hpp +++ b/sycl/include/sycl/detail/defines.hpp @@ -12,7 +12,7 @@ #if __SYCL_ID_QUERIES_FIT_IN_INT__ && __has_builtin(__builtin_assume) #include -#define __SYCL_ASSUME_INT(x) __builtin_assume((x) <= INT_MAX) +#define __SYCL_ASSUME_INT(x) __builtin_assume((x) <= UINT_MAX) #else #define __SYCL_ASSUME_INT(x) #if __SYCL_ID_QUERIES_FIT_IN_INT__ && !__has_builtin(__builtin_assume) diff --git a/sycl/include/sycl/detail/id_queries_fit_in_int.hpp b/sycl/include/sycl/detail/id_queries_fit_in_int.hpp index 97abbd1ea1cfb..c7fe0f9157e6f 100644 --- a/sycl/include/sycl/detail/id_queries_fit_in_int.hpp +++ b/sycl/include/sycl/detail/id_queries_fit_in_int.hpp @@ -8,7 +8,7 @@ // // Our SYCL implementation has a special mode (introduced for performance // reasons) in which it assume that all result of all id queries (i.e. global -// sizes, work-group sizes, local id, global id, etc.) fit within MAX_INT. +// sizes, work-group sizes, local id, global id, etc.) fit within UINT_MAX. // // This header contains corresponding helper functions related to this mode. // @@ -35,7 +35,7 @@ namespace detail { #if __SYCL_ID_QUERIES_FIT_IN_INT__ constexpr static const char *Msg = - "Provided range and/or offset does not fit in int. Pass " + "Provided range and/or offset does not fit in unsigned int. Pass " "`-fno-sycl-id-queries-fit-in-int' to remove this limit."; template @@ -43,14 +43,14 @@ typename std::enable_if_t::value || std::is_same::value> checkValueRangeImpl(ValT V) { static constexpr size_t Limit = - static_cast((std::numeric_limits::max)()); + static_cast((std::numeric_limits::max)()); if (V > Limit) throw sycl::exception(make_error_code(errc::nd_range), Msg); } inline void checkMulOverflow(size_t a, size_t b) { #ifndef _MSC_VER - int Product; + unsigned int Product; if (__builtin_mul_overflow(a, b, &Product)) { throw sycl::exception(make_error_code(errc::nd_range), Msg); } @@ -64,7 +64,7 @@ inline void checkMulOverflow(size_t a, size_t b) { inline void checkMulOverflow(size_t a, size_t b, size_t c) { #ifndef _MSC_VER - int Product; + unsigned int Product; if (__builtin_mul_overflow(a, b, &Product) || __builtin_mul_overflow(Product, c, &Product)) { throw sycl::exception(make_error_code(errc::nd_range), Msg); diff --git a/sycl/test-e2e/Basic/range_offset_fit_in_int.cpp b/sycl/test-e2e/Basic/range_offset_fit_in_int.cpp index 95c9fa38dd597..16f3725930dfc 100644 --- a/sycl/test-e2e/Basic/range_offset_fit_in_int.cpp +++ b/sycl/test-e2e/Basic/range_offset_fit_in_int.cpp @@ -8,9 +8,9 @@ namespace S = sycl; -constexpr char Msg[] = "Provided range and/or offset does not fit in int. " - "Pass `-fno-sycl-id-queries-fit-in-int' to " - "remove this limit."; +constexpr char Msg[] = + "Provided range and/or offset does not fit in unsigned int. Pass " + "`-fno-sycl-id-queries-fit-in-int' to remove this limit."; void checkRangeException(S::exception &E) { std::cerr << E.what() << std::endl; @@ -33,7 +33,10 @@ void test() { S::queue Queue(EH); - static constexpr size_t OutOfLimitsSize = static_cast(INT_MAX) + 1; + // Test is valid only if size_t is larger than unsigned int, we don't run the + // testing on platforms where this condition is not met. + static_assert(sizeof(size_t) > sizeof(unsigned int)); + static constexpr size_t OutOfLimitsSize = static_cast(UINT_MAX) + 1; S::range<2> RangeOutOfLimits{OutOfLimitsSize, 1}; S::range<2> RangeInLimits{1, 1};