From 7553a78d81b221c00c78897b0fe7504f47afd4c4 Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Thu, 9 Jul 2026 10:30:07 -0400 Subject: [PATCH 1/9] Add macro configuration for SYCL support --- include/boost/int128/detail/config.hpp | 40 +++++++++++++++++++++----- 1 file changed, 33 insertions(+), 7 deletions(-) diff --git a/include/boost/int128/detail/config.hpp b/include/boost/int128/detail/config.hpp index 4a93c2c3..2d31b5f1 100644 --- a/include/boost/int128/detail/config.hpp +++ b/include/boost/int128/detail/config.hpp @@ -5,8 +5,17 @@ #ifndef BOOST_INT128_DETAIL_CONFIG_HPP #define BOOST_INT128_DETAIL_CONFIG_HPP -// Use 128-bit integers -#if defined(BOOST_HAS_INT128) || (defined(__SIZEOF_INT128__) && !defined(_MSC_VER)) && !defined(BOOST_INT128_NO_BUILTIN_INT128) +// The SYCL device target (spir64) has no native 128-bit integer, so force the portable +// code path on the device pass. This mirrors a user-supplied BOOST_INT128_NO_BUILTIN_INT128 +// and keeps host/device selection consistent even though __x86_64__ stays defined on device. +#if defined(__SYCL_DEVICE_ONLY__) && !defined(BOOST_INT128_NO_BUILTIN_INT128) +# define BOOST_INT128_NO_BUILTIN_INT128 +#endif + +// Use 128-bit integers. +// The SYCL device target (spir64) has no native 128-bit integer, so on the device pass +// we fall back to the portable path (the same one used on platforms without __int128). +#if (defined(BOOST_HAS_INT128) || (defined(__SIZEOF_INT128__) && !defined(_MSC_VER)) && !defined(BOOST_INT128_NO_BUILTIN_INT128)) && !defined(__SYCL_DEVICE_ONLY__) #define BOOST_INT128_HAS_INT128 @@ -33,7 +42,7 @@ using builtin_u128 = unsigned __int128; } // namespace int128 } // namespace boost -#elif __has_include(<__msvc_int128.hpp>) && _MSVC_LANG >= 202002L +#elif __has_include(<__msvc_int128.hpp>) && _MSVC_LANG >= 202002L && !defined(__SYCL_DEVICE_ONLY__) #ifndef BOOST_INT128_BUILD_MODULE #include <__msvc_int128.hpp> @@ -170,7 +179,8 @@ using builtin_u128 = std::_Unsigned128; #endif // Platform macros // Hardware 128-bit by 64-bit unsigned division via the x86-64 DIV instruction -#if defined(__x86_64__) && (defined(__GNUC__) || defined(__clang__)) && !defined(_MSC_VER) && !defined(__CUDA_ARCH__) +// Excluded on the CUDA and SYCL device passes (the device target is not x86-64) +#if defined(__x86_64__) && (defined(__GNUC__) || defined(__clang__)) && !defined(_MSC_VER) && !defined(__CUDA_ARCH__) && !defined(__SYCL_DEVICE_ONLY__) # define BOOST_INT128_HAS_X86_64_DIVQ #endif @@ -290,10 +300,26 @@ using builtin_u128 = std::_Unsigned128; # endif #endif +// GPU device support. CUDA is auto-detected via __CUDACC__ (opt-in with +// BOOST_INT128_ENABLE_CUDA). SYCL is fully opt-in via BOOST_INT128_ENABLE_SYCL; +// must be included before so SYCL_EXTERNAL exists. #if defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA) -# define BOOST_INT128_HOST_DEVICE __host__ __device__ -#else -# define BOOST_INT128_HOST_DEVICE +# define BOOST_INT128_CUDA_ENABLED __host__ __device__ +# define BOOST_INT128_HAS_GPU_SUPPORT +#elif defined(BOOST_INT128_ENABLE_SYCL) +# define BOOST_INT128_SYCL_ENABLED SYCL_EXTERNAL +# define BOOST_INT128_HAS_GPU_SUPPORT #endif +#ifndef BOOST_INT128_CUDA_ENABLED +# define BOOST_INT128_CUDA_ENABLED +#endif +#ifndef BOOST_INT128_SYCL_ENABLED +# define BOOST_INT128_SYCL_ENABLED +#endif + +// Exactly one sub-macro is ever non-empty; expands to "__host__ __device__" (CUDA), +// "SYCL_EXTERNAL" (SYCL), or nothing (host). +#define BOOST_INT128_HOST_DEVICE BOOST_INT128_CUDA_ENABLED BOOST_INT128_SYCL_ENABLED + #endif // BOOST_INT128_DETAIL_CONFIG_HPP From d30061dfa3accb787fde8d56883870271722eeca Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Thu, 9 Jul 2026 10:30:44 -0400 Subject: [PATCH 2/9] Replace CUDA macros with general GPU ones --- include/boost/int128/detail/common_mul.hpp | 4 ++-- include/boost/int128/detail/int128_imp.hpp | 6 +++--- include/boost/int128/detail/mini_from_chars.hpp | 4 ++-- include/boost/int128/detail/uint128_imp.hpp | 6 +++--- 4 files changed, 10 insertions(+), 10 deletions(-) diff --git a/include/boost/int128/detail/common_mul.hpp b/include/boost/int128/detail/common_mul.hpp index e0c1a8e1..d4848929 100644 --- a/include/boost/int128/detail/common_mul.hpp +++ b/include/boost/int128/detail/common_mul.hpp @@ -48,11 +48,11 @@ BOOST_INT128_HOST_DEVICE BOOST_INT128_FORCE_INLINE constexpr std::uint64_t umul( hi = static_cast(product >> 64U); return static_cast(product); - #elif defined(_M_AMD64) && !defined(__GNUC__) && !defined(__CUDA_ARCH__) + #elif defined(_M_AMD64) && !defined(__GNUC__) && !defined(__CUDA_ARCH__) && !defined(__SYCL_DEVICE_ONLY__) return _umul128(a, b, &hi); - #elif defined(_M_ARM64) && !defined(__CUDA_ARCH__) + #elif defined(_M_ARM64) && !defined(__CUDA_ARCH__) && !defined(__SYCL_DEVICE_ONLY__) hi = __umulh(a, b); return a * b; diff --git a/include/boost/int128/detail/int128_imp.hpp b/include/boost/int128/detail/int128_imp.hpp index 898ae962..0ad50e18 100644 --- a/include/boost/int128/detail/int128_imp.hpp +++ b/include/boost/int128/detail/int128_imp.hpp @@ -122,8 +122,8 @@ int128_t BOOST_INT128_HOST_DEVICE constexpr operator float() const noexcept; BOOST_INT128_HOST_DEVICE constexpr operator double() const noexcept; - // Long double does not exist on device - #if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)) + // Long double does not exist on the CUDA or SYCL (spir64) device + #if !defined(BOOST_INT128_HAS_GPU_SUPPORT) constexpr operator long double() const noexcept; #endif @@ -300,7 +300,7 @@ BOOST_INT128_HOST_DEVICE constexpr int128_t::operator double() const noexcept return static_cast(high) * detail::offset_value_v + static_cast(low); } -#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)) +#if !defined(BOOST_INT128_HAS_GPU_SUPPORT) constexpr int128_t::operator long double() const noexcept { diff --git a/include/boost/int128/detail/mini_from_chars.hpp b/include/boost/int128/detail/mini_from_chars.hpp index 92c2f27c..233b37be 100644 --- a/include/boost/int128/detail/mini_from_chars.hpp +++ b/include/boost/int128/detail/mini_from_chars.hpp @@ -23,7 +23,7 @@ namespace detail { namespace impl { -#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)) +#if !defined(BOOST_INT128_HAS_GPU_SUPPORT) BOOST_INT128_INLINE_CONSTEXPR unsigned char uchar_values[] = {255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, @@ -50,7 +50,7 @@ static_assert(sizeof(uchar_values) == 256, "uchar_values should represent all 25 // Convert characters for 0-9, A-Z, a-z to 0-35. Anything else is 255 BOOST_INT128_HOST_DEVICE BOOST_INT128_FORCE_INLINE constexpr auto digit_from_char(char val) noexcept -> unsigned char { - #if defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA) + #if defined(BOOST_INT128_HAS_GPU_SUPPORT) constexpr unsigned char uchar_values[] = {255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, diff --git a/include/boost/int128/detail/uint128_imp.hpp b/include/boost/int128/detail/uint128_imp.hpp index e1b8d73d..8ba30bf0 100644 --- a/include/boost/int128/detail/uint128_imp.hpp +++ b/include/boost/int128/detail/uint128_imp.hpp @@ -133,8 +133,8 @@ uint128_t BOOST_INT128_HOST_DEVICE constexpr operator float() const noexcept; BOOST_INT128_HOST_DEVICE constexpr operator double() const noexcept; - // long doubles do not exist on device - #if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)) + // long doubles do not exist on the CUDA or SYCL (spir64) device + #if !defined(BOOST_INT128_HAS_GPU_SUPPORT) constexpr operator long double() const noexcept; #endif @@ -302,7 +302,7 @@ BOOST_INT128_HOST_DEVICE constexpr uint128_t::operator double() const noexcept return static_cast(high) * detail::offset_value_v + static_cast(low); } -#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)) +#if !defined(BOOST_INT128_HAS_GPU_SUPPORT) constexpr uint128_t::operator long double() const noexcept { From a9373e56fa2abdda1673b8b40164d0899598ef7d Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Thu, 9 Jul 2026 10:31:20 -0400 Subject: [PATCH 3/9] Fix potential buffer overflows --- include/boost/int128/detail/mini_to_chars.hpp | 14 +++++++++----- include/boost/int128/fmt_format.hpp | 2 +- include/boost/int128/format.hpp | 2 +- include/boost/int128/iostream.hpp | 2 +- include/boost/int128/string.hpp | 2 +- 5 files changed, 13 insertions(+), 9 deletions(-) diff --git a/include/boost/int128/detail/mini_to_chars.hpp b/include/boost/int128/detail/mini_to_chars.hpp index 80e1fcd0..5f4d6c7c 100644 --- a/include/boost/int128/detail/mini_to_chars.hpp +++ b/include/boost/int128/detail/mini_to_chars.hpp @@ -12,7 +12,11 @@ namespace boost { namespace int128 { namespace detail { -#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)) +// A 128-bit integer needs up to 128 binary digits (base 2); allow for a leading sign and a +// null terminator so mini_to_chars is safe for every supported base (2, 8, 10, 16). +constexpr std::size_t mini_to_chars_buffer_size = 130; + +#if !defined(BOOST_INT128_HAS_GPU_SUPPORT) BOOST_INT128_INLINE_CONSTEXPR char lower_case_digit_table[] = { '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', @@ -30,9 +34,9 @@ static_assert(sizeof(upper_case_digit_table) == sizeof(char) * 16, "10 numbers, #endif // !__NVCC__ -BOOST_INT128_HOST_DEVICE constexpr char* mini_to_chars(char (&buffer)[64], uint128_t v, const int base, const bool uppercase) noexcept +BOOST_INT128_HOST_DEVICE constexpr char* mini_to_chars(char (&buffer)[mini_to_chars_buffer_size], uint128_t v, const int base, const bool uppercase) noexcept { - #if defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA) + #if defined(BOOST_INT128_HAS_GPU_SUPPORT) constexpr char lower_case_digit_table[] = { '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', 'a', 'b', 'c', 'd', 'e', 'f' @@ -44,7 +48,7 @@ BOOST_INT128_HOST_DEVICE constexpr char* mini_to_chars(char (&buffer)[64], uint1 }; #endif - char* last {buffer + 64U}; + char* last {buffer + sizeof(buffer)}; *--last = '\0'; if (v == 0U) @@ -97,7 +101,7 @@ BOOST_INT128_HOST_DEVICE constexpr char* mini_to_chars(char (&buffer)[64], uint1 return last; } -BOOST_INT128_HOST_DEVICE constexpr char* mini_to_chars(char (&buffer)[64], const int128_t v, const int base, const bool uppercase) noexcept +BOOST_INT128_HOST_DEVICE constexpr char* mini_to_chars(char (&buffer)[mini_to_chars_buffer_size], const int128_t v, const int base, const bool uppercase) noexcept { char* p {nullptr}; diff --git a/include/boost/int128/fmt_format.hpp b/include/boost/int128/fmt_format.hpp index e37f3634..2199f8b5 100644 --- a/include/boost/int128/fmt_format.hpp +++ b/include/boost/int128/fmt_format.hpp @@ -246,7 +246,7 @@ struct formatter template auto format(T v, FormatContext& ctx) const { - char buffer[64]; + char buffer[detail::mini_to_chars_buffer_size]; bool isneg {false}; boost::int128::uint128_t abs_v {}; diff --git a/include/boost/int128/format.hpp b/include/boost/int128/format.hpp index 417442c7..bc0e5bfa 100644 --- a/include/boost/int128/format.hpp +++ b/include/boost/int128/format.hpp @@ -235,7 +235,7 @@ struct formatter template auto format(T v, FormatContext& ctx) const { - char buffer[64]; + char buffer[boost::int128::detail::mini_to_chars_buffer_size]; bool isneg {false}; boost::int128::uint128_t abs_v {}; diff --git a/include/boost/int128/iostream.hpp b/include/boost/int128/iostream.hpp index d6c4c7c0..878e4610 100644 --- a/include/boost/int128/iostream.hpp +++ b/include/boost/int128/iostream.hpp @@ -103,7 +103,7 @@ BOOST_INT128_EXPORT template & os, const LibIntegerType& v) -> std::enable_if_t, std::basic_ostream&> { - char buffer[64U] {}; + char buffer[detail::mini_to_chars_buffer_size] {}; const auto flags {os.flags()}; int base {10}; diff --git a/include/boost/int128/string.hpp b/include/boost/int128/string.hpp index f7a4812b..ccf44fd9 100644 --- a/include/boost/int128/string.hpp +++ b/include/boost/int128/string.hpp @@ -20,7 +20,7 @@ namespace int128 { template auto to_string(const T& value) -> std::enable_if_t<(std::is_same::value || std::is_same::value), std::string> { - char buffer[64]; + char buffer[detail::mini_to_chars_buffer_size]; const auto last {detail::mini_to_chars(buffer, value, 10, false)}; return std::string{last, buffer + sizeof(buffer) - 1}; } From ec6cebd4f28f57ff8a5f4ed98420b758f7dabf2e Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Thu, 9 Jul 2026 10:31:42 -0400 Subject: [PATCH 4/9] Use newly adapted Boost.Charconv for SYCL support --- include/boost/int128/charconv.hpp | 20 ++++++++++++-------- 1 file changed, 12 insertions(+), 8 deletions(-) diff --git a/include/boost/int128/charconv.hpp b/include/boost/int128/charconv.hpp index 01b4fcb9..44183129 100644 --- a/include/boost/int128/charconv.hpp +++ b/include/boost/int128/charconv.hpp @@ -16,6 +16,10 @@ # define BOOST_CHARCONV_ENABLE_CUDA #endif +#if defined(BOOST_INT128_ENABLE_SYCL) && !defined(BOOST_CHARCONV_ENABLE_SYCL) +# define BOOST_CHARCONV_ENABLE_SYCL +#endif + #include #include #include @@ -44,23 +48,23 @@ struct make_signed { using type = int128::int128_t; }; template <> struct make_signed { using type = int128::int128_t; }; -#if defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA) +#if defined(BOOST_INT128_HAS_GPU_SUPPORT) template <> -__host__ __device__ constexpr int128::uint128_t get_max_value() +BOOST_INT128_HOST_DEVICE constexpr int128::uint128_t get_max_value() { return std::numeric_limits::max(); } template <> -__host__ __device__ constexpr int128::int128_t get_max_value() +BOOST_INT128_HOST_DEVICE constexpr int128::int128_t get_max_value() { return std::numeric_limits::max(); } -#endif // __NVCC__ +#endif -#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)) +#if !defined(BOOST_INT128_HAS_GPU_SUPPORT) BOOST_INT128_INLINE_CONSTEXPR int128::uint128_t int128_pow10[39] = { @@ -109,7 +113,7 @@ BOOST_INT128_INLINE_CONSTEXPR int128::uint128_t int128_pow10[39] = BOOST_INT128_HOST_DEVICE constexpr int num_digits(const int128::uint128_t& x) noexcept { - #if defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA) + #if defined(BOOST_INT128_HAS_GPU_SUPPORT) constexpr int128::uint128_t int128_pow10[39] = { @@ -187,7 +191,7 @@ BOOST_INT128_HOST_DEVICE constexpr int num_digits(const int128::uint128_t& x) no BOOST_INT128_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, const int128::uint128_t value, const int base = 10) noexcept { - #if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)) + #if !defined(BOOST_INT128_HAS_GPU_SUPPORT) if (base == 10) { @@ -201,7 +205,7 @@ BOOST_INT128_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* BOOST_INT128_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, const int128::int128_t value, const int base = 10) noexcept { - #if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)) + #if !defined(BOOST_INT128_HAS_GPU_SUPPORT) if (base == 10) { From f1ddca31ab63c4ce002f33594e00a7e9ca3213b9 Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Thu, 9 Jul 2026 10:32:04 -0400 Subject: [PATCH 5/9] Add SYCL test harness --- test/sycl_test.hpp | 318 +++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 318 insertions(+) create mode 100644 test/sycl_test.hpp diff --git a/test/sycl_test.hpp b/test/sycl_test.hpp new file mode 100644 index 00000000..d48925f5 --- /dev/null +++ b/test/sycl_test.hpp @@ -0,0 +1,318 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// Shared harness for the Boost.int128 SYCL tests. Each test runs an operation +// element-wise on the SYCL device over random inputs and verifies that the +// device results match a host recomputation of the same operation. + +#ifndef BOOST_INT128_TEST_SYCL_TEST_HPP +#define BOOST_INT128_TEST_SYCL_TEST_HPP + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace int128_sycl_test { + +constexpr int num_elements {20000}; +// Matches the buffer size mini_to_chars expects (large enough for base 2). +constexpr int buf_size {static_cast(boost::int128::detail::mini_to_chars_buffer_size)}; + +// Produces a full-range random value of the int128 type T from a 64-bit engine. This avoids +// pulling Boost.Random (whose headers are host-only) into the single-source SYCL translation +// unit, where the device pass would still have to parse them. +template +T random_value(std::mt19937_64& rng) +{ + using high_type = decltype(T{}.high); + return T{static_cast(rng()), static_cast(rng())}; +} + +// Returns a divisor that avoids the divide-by-zero and INT_MIN / -1 cases, so division +// stays well defined. Applied identically on host and device. +template +inline T safe_divisor(const T a, const T b) noexcept +{ + if (b == T{0}) + { + return T{1}; + } + if (std::numeric_limits::is_signed && a == (std::numeric_limits::min)() && b == T{-1}) + { + return T{1}; + } + return b; +} + +template +inline T safe_div(const T a, const T b) noexcept +{ + return a / safe_divisor(a, b); +} + +template +inline T safe_mod(const T a, const T b) noexcept +{ + return a % safe_divisor(a, b); +} + +// Generic runner. op(a, b, i) is evaluated on the device for every element and then +// re-evaluated on the host; any mismatch fails the test. op must depend only on +// device-enabled int128 facilities. +template +int run(Op op) +{ + sycl::queue q; + std::cout << "SYCL device: " + << q.get_device().get_info() << "\n"; + + InT* a {sycl::malloc_shared(num_elements, q)}; + InT* b {sycl::malloc_shared(num_elements, q)}; + OutT* out {sycl::malloc_shared(num_elements, q)}; + + std::mt19937_64 rng {42}; + for (int i {0}; i < num_elements; ++i) + { + a[i] = random_value(rng); + b[i] = random_value(rng); + } + + q.submit([&](sycl::handler& h) + { + h.parallel_for(sycl::range<1>(num_elements), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + out[i] = op(a[i], b[i], i); + }); + }).wait(); + + int failures {0}; + for (int i {0}; i < num_elements; ++i) + { + if (out[i] != op(a[i], b[i], i)) + { + if (failures < 5) + { + std::cerr << "Mismatch at element " << i << "\n"; + } + ++failures; + } + } + + sycl::free(a, q); + sycl::free(b, q); + sycl::free(out, q); + + if (failures == 0) + { + std::cout << "Test PASSED\n"; + return EXIT_SUCCESS; + } + + std::cerr << "Test FAILED with " << failures << " mismatches\n"; + return EXIT_FAILURE; +} + +// Comparison runner. Half of the second operands are forced equal to the first so the +// equal/less-equal/greater-equal branches are exercised alongside the strict ones. +template +int run_compare(Pred pred) +{ + sycl::queue q; + std::cout << "SYCL device: " + << q.get_device().get_info() << "\n"; + + InT* a {sycl::malloc_shared(num_elements, q)}; + InT* b {sycl::malloc_shared(num_elements, q)}; + int* out {sycl::malloc_shared(num_elements, q)}; + + std::mt19937_64 rng {42}; + for (int i {0}; i < num_elements; ++i) + { + a[i] = random_value(rng); + b[i] = (i % 2 == 0) ? a[i] : random_value(rng); + } + + q.submit([&](sycl::handler& h) + { + h.parallel_for(sycl::range<1>(num_elements), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + out[i] = static_cast(pred(a[i], b[i])); + }); + }).wait(); + + int failures {0}; + for (int i {0}; i < num_elements; ++i) + { + if (out[i] != static_cast(pred(a[i], b[i]))) + { + if (failures < 5) + { + std::cerr << "Mismatch at element " << i << "\n"; + } + ++failures; + } + } + + sycl::free(a, q); + sycl::free(b, q); + sycl::free(out, q); + + if (failures == 0) + { + std::cout << "Test PASSED\n"; + return EXIT_SUCCESS; + } + + std::cerr << "Test FAILED with " << failures << " mismatches\n"; + return EXIT_FAILURE; +} + +// to_chars runner. Formats each value on the device via boost::charconv::to_chars (now +// SYCL device-capable) and compares the produced text against the host. +template +int run_to_chars(const int base) +{ + using boost::charconv::to_chars; + + sycl::queue q; + std::cout << "SYCL device: " + << q.get_device().get_info() << "\n"; + + InT* in {sycl::malloc_shared(num_elements, q)}; + char* out_str {sycl::malloc_shared(num_elements * buf_size, q)}; + int* out_len {sycl::malloc_shared(num_elements, q)}; + + std::mt19937_64 rng {42}; + for (int i {0}; i < num_elements; ++i) + { + in[i] = random_value(rng); + } + + q.submit([&](sycl::handler& h) + { + h.parallel_for(sycl::range<1>(num_elements), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + char buf[buf_size]; + const auto r {to_chars(buf, buf + buf_size, in[i], base)}; + const int n {static_cast(r.ptr - buf)}; + char* dest {out_str + i * buf_size}; + for (int k {0}; k < n; ++k) + { + dest[k] = buf[k]; + } + out_len[i] = n; + }); + }).wait(); + + int failures {0}; + for (int i {0}; i < num_elements; ++i) + { + char buf[buf_size]; + const auto r {to_chars(buf, buf + buf_size, in[i], base)}; + const int n {static_cast(r.ptr - buf)}; + if (out_len[i] != n || std::memcmp(out_str + i * buf_size, buf, static_cast(n)) != 0) + { + if (failures < 5) + { + std::cerr << "Mismatch at element " << i << "\n"; + } + ++failures; + } + } + + sycl::free(in, q); + sycl::free(out_str, q); + sycl::free(out_len, q); + + if (failures == 0) + { + std::cout << "Test PASSED\n"; + return EXIT_SUCCESS; + } + + std::cerr << "Test FAILED with " << failures << " mismatches\n"; + return EXIT_FAILURE; +} + +// from_chars runner. Formats each value on the host, parses it back on the device via +// boost::charconv::from_chars (now SYCL device-capable), and checks the round trip. +template +int run_from_chars(const int base) +{ + using boost::charconv::to_chars; + using boost::charconv::from_chars; + + sycl::queue q; + std::cout << "SYCL device: " + << q.get_device().get_info() << "\n"; + + char* in_str {sycl::malloc_shared(num_elements * buf_size, q)}; + int* in_len {sycl::malloc_shared(num_elements, q)}; + InT* out {sycl::malloc_shared(num_elements, q)}; + InT* expected {sycl::malloc_shared(num_elements, q)}; + + std::mt19937_64 rng {42}; + for (int i {0}; i < num_elements; ++i) + { + expected[i] = random_value(rng); + char buf[buf_size]; + const auto r {to_chars(buf, buf + buf_size, expected[i], base)}; + const int n {static_cast(r.ptr - buf)}; + std::memcpy(in_str + i * buf_size, buf, static_cast(n)); + in_len[i] = n; + } + + q.submit([&](sycl::handler& h) + { + h.parallel_for(sycl::range<1>(num_elements), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + const char* first {in_str + i * buf_size}; + InT value {}; + from_chars(first, first + in_len[i], value, base); + out[i] = value; + }); + }).wait(); + + int failures {0}; + for (int i {0}; i < num_elements; ++i) + { + if (out[i] != expected[i]) + { + if (failures < 5) + { + std::cerr << "Mismatch at element " << i << "\n"; + } + ++failures; + } + } + + sycl::free(in_str, q); + sycl::free(in_len, q); + sycl::free(out, q); + sycl::free(expected, q); + + if (failures == 0) + { + std::cout << "Test PASSED\n"; + return EXIT_SUCCESS; + } + + std::cerr << "Test FAILED with " << failures << " mismatches\n"; + return EXIT_FAILURE; +} + +} // namespace int128_sycl_test + +#endif // BOOST_INT128_TEST_SYCL_TEST_HPP From 23951c967d8abe74183dccbf2a33958cbce26d82 Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Thu, 9 Jul 2026 10:32:46 -0400 Subject: [PATCH 6/9] Add SYCL testing --- test/sycl_jamfile | 101 +++++++++++++++++ test/test_bit_ceil_sycl.cpp | 14 +++ test/test_bit_floor_sycl.cpp | 13 +++ test/test_bit_width_sycl.cpp | 13 +++ test/test_byteswap_sycl.cpp | 13 +++ test/test_countl_one_sycl.cpp | 13 +++ test/test_countl_zero_sycl.cpp | 13 +++ test/test_countr_one_sycl.cpp | 13 +++ test/test_countr_zero_sycl.cpp | 13 +++ test/test_has_single_bit_sycl.cpp | 13 +++ test/test_popcount_sycl.cpp | 13 +++ test/test_rotl_sycl.cpp | 13 +++ test/test_rotr_sycl.cpp | 13 +++ test/test_signed_add_sat_sycl.cpp | 13 +++ test/test_signed_add_sycl.cpp | 13 +++ test/test_signed_and_sycl.cpp | 13 +++ test/test_signed_cstdlib_div_sycl.cpp | 17 +++ test/test_signed_div_sat_sycl.cpp | 13 +++ test/test_signed_div_sycl.cpp | 13 +++ test/test_signed_eq_sycl.cpp | 13 +++ test/test_signed_from_chars_bases_sycl.cpp | 17 +++ test/test_signed_from_chars_sycl.cpp | 12 ++ test/test_signed_gcd_sycl.cpp | 13 +++ test/test_signed_ge_sycl.cpp | 13 +++ test/test_signed_gt_sycl.cpp | 13 +++ test/test_signed_lcm_sycl.cpp | 14 +++ test/test_signed_le_sycl.cpp | 13 +++ test/test_signed_left_shift_sycl.cpp | 13 +++ test/test_signed_literals_sycl.cpp | 89 +++++++++++++++ test/test_signed_lt_sycl.cpp | 13 +++ test/test_signed_midpoint_sycl.cpp | 13 +++ test/test_signed_mod_sycl.cpp | 13 +++ test/test_signed_mul_sat_sycl.cpp | 13 +++ test/test_signed_mul_sycl.cpp | 13 +++ test/test_signed_ne_sycl.cpp | 13 +++ test/test_signed_not_sycl.cpp | 13 +++ test/test_signed_or_sycl.cpp | 13 +++ test/test_signed_right_shift_sycl.cpp | 13 +++ test/test_signed_sub_sat_sycl.cpp | 13 +++ test/test_signed_sub_sycl.cpp | 13 +++ test/test_signed_to_chars_bases_sycl.cpp | 17 +++ test/test_signed_to_chars_sycl.cpp | 12 ++ ...est_signed_to_unsigned_conversion_sycl.cpp | 14 +++ test/test_signed_xor_sycl.cpp | 13 +++ test/test_unsigned_add_sat_sycl.cpp | 13 +++ test/test_unsigned_add_sycl.cpp | 13 +++ test/test_unsigned_and_sycl.cpp | 13 +++ test/test_unsigned_cstdlib_div_sycl.cpp | 17 +++ test/test_unsigned_div_sat_sycl.cpp | 13 +++ test/test_unsigned_div_sycl.cpp | 13 +++ test/test_unsigned_eq_sycl.cpp | 13 +++ test/test_unsigned_from_chars_bases_sycl.cpp | 17 +++ test/test_unsigned_from_chars_sycl.cpp | 12 ++ test/test_unsigned_gcd_sycl.cpp | 13 +++ test/test_unsigned_ge_sycl.cpp | 13 +++ test/test_unsigned_gt_sycl.cpp | 13 +++ test/test_unsigned_lcm_sycl.cpp | 14 +++ test/test_unsigned_le_sycl.cpp | 13 +++ test/test_unsigned_left_shift_sycl.cpp | 13 +++ test/test_unsigned_literals_sycl.cpp | 105 ++++++++++++++++++ test/test_unsigned_lt_sycl.cpp | 13 +++ test/test_unsigned_midpoint_sycl.cpp | 13 +++ test/test_unsigned_mod_sycl.cpp | 13 +++ test/test_unsigned_mul_sat_sycl.cpp | 13 +++ test/test_unsigned_mul_sycl.cpp | 13 +++ test/test_unsigned_ne_sycl.cpp | 13 +++ test/test_unsigned_not_sycl.cpp | 13 +++ test/test_unsigned_or_sycl.cpp | 13 +++ test/test_unsigned_right_shift_sycl.cpp | 13 +++ test/test_unsigned_sub_sat_sycl.cpp | 13 +++ test/test_unsigned_sub_sycl.cpp | 13 +++ test/test_unsigned_to_chars_bases_sycl.cpp | 17 +++ test/test_unsigned_to_chars_sycl.cpp | 12 ++ ...est_unsigned_to_signed_conversion_sycl.cpp | 14 +++ test/test_unsigned_xor_sycl.cpp | 13 +++ 75 files changed, 1256 insertions(+) create mode 100644 test/sycl_jamfile create mode 100644 test/test_bit_ceil_sycl.cpp create mode 100644 test/test_bit_floor_sycl.cpp create mode 100644 test/test_bit_width_sycl.cpp create mode 100644 test/test_byteswap_sycl.cpp create mode 100644 test/test_countl_one_sycl.cpp create mode 100644 test/test_countl_zero_sycl.cpp create mode 100644 test/test_countr_one_sycl.cpp create mode 100644 test/test_countr_zero_sycl.cpp create mode 100644 test/test_has_single_bit_sycl.cpp create mode 100644 test/test_popcount_sycl.cpp create mode 100644 test/test_rotl_sycl.cpp create mode 100644 test/test_rotr_sycl.cpp create mode 100644 test/test_signed_add_sat_sycl.cpp create mode 100644 test/test_signed_add_sycl.cpp create mode 100644 test/test_signed_and_sycl.cpp create mode 100644 test/test_signed_cstdlib_div_sycl.cpp create mode 100644 test/test_signed_div_sat_sycl.cpp create mode 100644 test/test_signed_div_sycl.cpp create mode 100644 test/test_signed_eq_sycl.cpp create mode 100644 test/test_signed_from_chars_bases_sycl.cpp create mode 100644 test/test_signed_from_chars_sycl.cpp create mode 100644 test/test_signed_gcd_sycl.cpp create mode 100644 test/test_signed_ge_sycl.cpp create mode 100644 test/test_signed_gt_sycl.cpp create mode 100644 test/test_signed_lcm_sycl.cpp create mode 100644 test/test_signed_le_sycl.cpp create mode 100644 test/test_signed_left_shift_sycl.cpp create mode 100644 test/test_signed_literals_sycl.cpp create mode 100644 test/test_signed_lt_sycl.cpp create mode 100644 test/test_signed_midpoint_sycl.cpp create mode 100644 test/test_signed_mod_sycl.cpp create mode 100644 test/test_signed_mul_sat_sycl.cpp create mode 100644 test/test_signed_mul_sycl.cpp create mode 100644 test/test_signed_ne_sycl.cpp create mode 100644 test/test_signed_not_sycl.cpp create mode 100644 test/test_signed_or_sycl.cpp create mode 100644 test/test_signed_right_shift_sycl.cpp create mode 100644 test/test_signed_sub_sat_sycl.cpp create mode 100644 test/test_signed_sub_sycl.cpp create mode 100644 test/test_signed_to_chars_bases_sycl.cpp create mode 100644 test/test_signed_to_chars_sycl.cpp create mode 100644 test/test_signed_to_unsigned_conversion_sycl.cpp create mode 100644 test/test_signed_xor_sycl.cpp create mode 100644 test/test_unsigned_add_sat_sycl.cpp create mode 100644 test/test_unsigned_add_sycl.cpp create mode 100644 test/test_unsigned_and_sycl.cpp create mode 100644 test/test_unsigned_cstdlib_div_sycl.cpp create mode 100644 test/test_unsigned_div_sat_sycl.cpp create mode 100644 test/test_unsigned_div_sycl.cpp create mode 100644 test/test_unsigned_eq_sycl.cpp create mode 100644 test/test_unsigned_from_chars_bases_sycl.cpp create mode 100644 test/test_unsigned_from_chars_sycl.cpp create mode 100644 test/test_unsigned_gcd_sycl.cpp create mode 100644 test/test_unsigned_ge_sycl.cpp create mode 100644 test/test_unsigned_gt_sycl.cpp create mode 100644 test/test_unsigned_lcm_sycl.cpp create mode 100644 test/test_unsigned_le_sycl.cpp create mode 100644 test/test_unsigned_left_shift_sycl.cpp create mode 100644 test/test_unsigned_literals_sycl.cpp create mode 100644 test/test_unsigned_lt_sycl.cpp create mode 100644 test/test_unsigned_midpoint_sycl.cpp create mode 100644 test/test_unsigned_mod_sycl.cpp create mode 100644 test/test_unsigned_mul_sat_sycl.cpp create mode 100644 test/test_unsigned_mul_sycl.cpp create mode 100644 test/test_unsigned_ne_sycl.cpp create mode 100644 test/test_unsigned_not_sycl.cpp create mode 100644 test/test_unsigned_or_sycl.cpp create mode 100644 test/test_unsigned_right_shift_sycl.cpp create mode 100644 test/test_unsigned_sub_sat_sycl.cpp create mode 100644 test/test_unsigned_sub_sycl.cpp create mode 100644 test/test_unsigned_to_chars_bases_sycl.cpp create mode 100644 test/test_unsigned_to_chars_sycl.cpp create mode 100644 test/test_unsigned_to_signed_conversion_sycl.cpp create mode 100644 test/test_unsigned_xor_sycl.cpp diff --git a/test/sycl_jamfile b/test/sycl_jamfile new file mode 100644 index 00000000..ae480b5d --- /dev/null +++ b/test/sycl_jamfile @@ -0,0 +1,101 @@ +# Copyright 2026 Matt Borland +# Distributed under the Boost Software License, Version 1.0. +# https://www.boost.org/LICENSE_1_0.txt + +import testing ; +import ../../config/checks/config : requires ; + +project : requirements + [ requires cxx14_decltype_auto cxx14_generic_lambdas cxx14_return_type_deduction cxx14_variable_templates cxx14_constexpr ] + ; + +run test_unsigned_add_sycl.cpp ; +run test_signed_add_sycl.cpp ; +run test_unsigned_sub_sycl.cpp ; +run test_signed_sub_sycl.cpp ; +run test_unsigned_mul_sycl.cpp ; +run test_signed_mul_sycl.cpp ; +run test_unsigned_div_sycl.cpp ; +run test_signed_div_sycl.cpp ; +run test_unsigned_mod_sycl.cpp ; +run test_signed_mod_sycl.cpp ; + +run test_unsigned_or_sycl.cpp ; +run test_signed_or_sycl.cpp ; +run test_unsigned_and_sycl.cpp ; +run test_signed_and_sycl.cpp ; +run test_unsigned_xor_sycl.cpp ; +run test_signed_xor_sycl.cpp ; +run test_unsigned_not_sycl.cpp ; +run test_signed_not_sycl.cpp ; +run test_unsigned_left_shift_sycl.cpp ; +run test_signed_left_shift_sycl.cpp ; +run test_unsigned_right_shift_sycl.cpp ; +run test_signed_right_shift_sycl.cpp ; + +run test_has_single_bit_sycl.cpp ; +run test_countl_zero_sycl.cpp ; +run test_countl_one_sycl.cpp ; +run test_bit_width_sycl.cpp ; +run test_bit_ceil_sycl.cpp ; +run test_bit_floor_sycl.cpp ; +run test_countr_zero_sycl.cpp ; +run test_countr_one_sycl.cpp ; +run test_rotl_sycl.cpp ; +run test_rotr_sycl.cpp ; +run test_popcount_sycl.cpp ; +run test_byteswap_sycl.cpp ; + +run test_unsigned_eq_sycl.cpp ; +run test_signed_eq_sycl.cpp ; +run test_unsigned_ne_sycl.cpp ; +run test_signed_ne_sycl.cpp ; +run test_unsigned_lt_sycl.cpp ; +run test_signed_lt_sycl.cpp ; +run test_unsigned_le_sycl.cpp ; +run test_signed_le_sycl.cpp ; +run test_unsigned_gt_sycl.cpp ; +run test_signed_gt_sycl.cpp ; +run test_unsigned_ge_sycl.cpp ; +run test_signed_ge_sycl.cpp ; + +run test_unsigned_cstdlib_div_sycl.cpp ; +run test_signed_cstdlib_div_sycl.cpp ; + +run test_signed_to_unsigned_conversion_sycl.cpp ; +run test_unsigned_to_signed_conversion_sycl.cpp ; + +run test_unsigned_add_sat_sycl.cpp ; +run test_signed_add_sat_sycl.cpp ; +run test_unsigned_sub_sat_sycl.cpp ; +run test_signed_sub_sat_sycl.cpp ; +run test_unsigned_mul_sat_sycl.cpp ; +run test_signed_mul_sat_sycl.cpp ; +run test_unsigned_div_sat_sycl.cpp ; +run test_signed_div_sat_sycl.cpp ; +run test_unsigned_gcd_sycl.cpp ; +run test_signed_gcd_sycl.cpp ; +run test_unsigned_lcm_sycl.cpp ; +run test_signed_lcm_sycl.cpp ; +run test_unsigned_midpoint_sycl.cpp ; +# test_signed_midpoint_sycl.cpp is disabled: the signed midpoint kernel triggers an internal +# compiler error in the Intel oneAPI DPC++ device compiler (2026.1.0) during the loop-unroll +# pass (and an "Invalid cast" SPIR-V link error with range rounding disabled). The midpoint +# implementation itself is correct and runs on the host, on CUDA, and via the unsigned overload +# on SYCL. Re-enable once the toolchain is fixed. +# run test_signed_midpoint_sycl.cpp ; + +run test_unsigned_to_chars_sycl.cpp ; +run test_signed_to_chars_sycl.cpp ; +run test_unsigned_from_chars_sycl.cpp ; +run test_signed_from_chars_sycl.cpp ; + +run test_unsigned_to_chars_bases_sycl.cpp ; +run test_signed_to_chars_bases_sycl.cpp ; +run test_unsigned_from_chars_bases_sycl.cpp ; +run test_signed_from_chars_bases_sycl.cpp ; + +run test_unsigned_literals_sycl.cpp ; +run test_signed_literals_sycl.cpp ; + +run ../examples/sycl.cpp ; diff --git a/test/test_bit_ceil_sycl.cpp b/test/test_bit_ceil_sycl.cpp new file mode 100644 index 00000000..3ce20af6 --- /dev/null +++ b/test/test_bit_ceil_sycl.cpp @@ -0,0 +1,14 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + // Shift right by one so bit_ceil never has to round past 2^127. + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int) { return boost::int128::bit_ceil(a >> 1); }); +} diff --git a/test/test_bit_floor_sycl.cpp b/test/test_bit_floor_sycl.cpp new file mode 100644 index 00000000..44cc6046 --- /dev/null +++ b/test/test_bit_floor_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int) { return boost::int128::bit_floor(a); }); +} diff --git a/test/test_bit_width_sycl.cpp b/test/test_bit_width_sycl.cpp new file mode 100644 index 00000000..51f864fb --- /dev/null +++ b/test/test_bit_width_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int) { return boost::int128::bit_width(a); }); +} diff --git a/test/test_byteswap_sycl.cpp b/test/test_byteswap_sycl.cpp new file mode 100644 index 00000000..87fb124b --- /dev/null +++ b/test/test_byteswap_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int) { return boost::int128::byteswap(a); }); +} diff --git a/test/test_countl_one_sycl.cpp b/test/test_countl_one_sycl.cpp new file mode 100644 index 00000000..4fc726e8 --- /dev/null +++ b/test/test_countl_one_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int) { return boost::int128::countl_one(a); }); +} diff --git a/test/test_countl_zero_sycl.cpp b/test/test_countl_zero_sycl.cpp new file mode 100644 index 00000000..32b8e969 --- /dev/null +++ b/test/test_countl_zero_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int) { return boost::int128::countl_zero(a); }); +} diff --git a/test/test_countr_one_sycl.cpp b/test/test_countr_one_sycl.cpp new file mode 100644 index 00000000..cb337e3f --- /dev/null +++ b/test/test_countr_one_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int) { return boost::int128::countr_one(a); }); +} diff --git a/test/test_countr_zero_sycl.cpp b/test/test_countr_zero_sycl.cpp new file mode 100644 index 00000000..543d2f88 --- /dev/null +++ b/test/test_countr_zero_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int) { return boost::int128::countr_zero(a); }); +} diff --git a/test/test_has_single_bit_sycl.cpp b/test/test_has_single_bit_sycl.cpp new file mode 100644 index 00000000..53089127 --- /dev/null +++ b/test/test_has_single_bit_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int) { return static_cast(boost::int128::has_single_bit(a)); }); +} diff --git a/test/test_popcount_sycl.cpp b/test/test_popcount_sycl.cpp new file mode 100644 index 00000000..dcddf7b3 --- /dev/null +++ b/test/test_popcount_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int) { return boost::int128::popcount(a); }); +} diff --git a/test/test_rotl_sycl.cpp b/test/test_rotl_sycl.cpp new file mode 100644 index 00000000..455937e2 --- /dev/null +++ b/test/test_rotl_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int i) { return boost::int128::rotl(a, i % 128); }); +} diff --git a/test/test_rotr_sycl.cpp b/test/test_rotr_sycl.cpp new file mode 100644 index 00000000..3b84d634 --- /dev/null +++ b/test/test_rotr_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int i) { return boost::int128::rotr(a, i % 128); }); +} diff --git a/test/test_signed_add_sat_sycl.cpp b/test/test_signed_add_sat_sycl.cpp new file mode 100644 index 00000000..c60c4599 --- /dev/null +++ b/test/test_signed_add_sat_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return boost::int128::add_sat(a, b); }); +} diff --git a/test/test_signed_add_sycl.cpp b/test/test_signed_add_sycl.cpp new file mode 100644 index 00000000..503607e9 --- /dev/null +++ b/test/test_signed_add_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return a + b; }); +} diff --git a/test/test_signed_and_sycl.cpp b/test/test_signed_and_sycl.cpp new file mode 100644 index 00000000..b07cb135 --- /dev/null +++ b/test/test_signed_and_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return a & b; }); +} diff --git a/test/test_signed_cstdlib_div_sycl.cpp b/test/test_signed_cstdlib_div_sycl.cpp new file mode 100644 index 00000000..d29d1e19 --- /dev/null +++ b/test/test_signed_cstdlib_div_sycl.cpp @@ -0,0 +1,17 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) + { + const auto d {boost::int128::div(a, int128_sycl_test::safe_divisor(a, b))}; + return d.quot ^ d.rem; + }); +} diff --git a/test/test_signed_div_sat_sycl.cpp b/test/test_signed_div_sat_sycl.cpp new file mode 100644 index 00000000..3eddd6c2 --- /dev/null +++ b/test/test_signed_div_sat_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return boost::int128::div_sat(a, int128_sycl_test::safe_divisor(a, b)); }); +} diff --git a/test/test_signed_div_sycl.cpp b/test/test_signed_div_sycl.cpp new file mode 100644 index 00000000..c1e79735 --- /dev/null +++ b/test/test_signed_div_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return int128_sycl_test::safe_div(a, b); }); +} diff --git a/test/test_signed_eq_sycl.cpp b/test/test_signed_eq_sycl.cpp new file mode 100644 index 00000000..a189c583 --- /dev/null +++ b/test/test_signed_eq_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run_compare( + [](int128_t a, int128_t b) { return a == b; }); +} diff --git a/test/test_signed_from_chars_bases_sycl.cpp b/test/test_signed_from_chars_bases_sycl.cpp new file mode 100644 index 00000000..4eff473f --- /dev/null +++ b/test/test_signed_from_chars_bases_sycl.cpp @@ -0,0 +1,17 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + int result {0}; + result |= int128_sycl_test::run_from_chars(2); + result |= int128_sycl_test::run_from_chars(8); + result |= int128_sycl_test::run_from_chars(16); + result |= int128_sycl_test::run_from_chars(36); + return result; +} diff --git a/test/test_signed_from_chars_sycl.cpp b/test/test_signed_from_chars_sycl.cpp new file mode 100644 index 00000000..0c79d30a --- /dev/null +++ b/test/test_signed_from_chars_sycl.cpp @@ -0,0 +1,12 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run_from_chars(10); +} diff --git a/test/test_signed_gcd_sycl.cpp b/test/test_signed_gcd_sycl.cpp new file mode 100644 index 00000000..ee10860c --- /dev/null +++ b/test/test_signed_gcd_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return boost::int128::gcd(a, b); }); +} diff --git a/test/test_signed_ge_sycl.cpp b/test/test_signed_ge_sycl.cpp new file mode 100644 index 00000000..25d69c00 --- /dev/null +++ b/test/test_signed_ge_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run_compare( + [](int128_t a, int128_t b) { return a >= b; }); +} diff --git a/test/test_signed_gt_sycl.cpp b/test/test_signed_gt_sycl.cpp new file mode 100644 index 00000000..d5298102 --- /dev/null +++ b/test/test_signed_gt_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run_compare( + [](int128_t a, int128_t b) { return a > b; }); +} diff --git a/test/test_signed_lcm_sycl.cpp b/test/test_signed_lcm_sycl.cpp new file mode 100644 index 00000000..11962a42 --- /dev/null +++ b/test/test_signed_lcm_sycl.cpp @@ -0,0 +1,14 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + // Restrict inputs so the true LCM stays within 128 bits. + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return boost::int128::lcm(a >> 65, b >> 65); }); +} diff --git a/test/test_signed_le_sycl.cpp b/test/test_signed_le_sycl.cpp new file mode 100644 index 00000000..75aa2607 --- /dev/null +++ b/test/test_signed_le_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run_compare( + [](int128_t a, int128_t b) { return a <= b; }); +} diff --git a/test/test_signed_left_shift_sycl.cpp b/test/test_signed_left_shift_sycl.cpp new file mode 100644 index 00000000..0a1e1d03 --- /dev/null +++ b/test/test_signed_left_shift_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t, int i) { return a << (i % 128); }); +} diff --git a/test/test_signed_literals_sycl.cpp b/test/test_signed_literals_sycl.cpp new file mode 100644 index 00000000..9c48fe2c --- /dev/null +++ b/test/test_signed_literals_sycl.cpp @@ -0,0 +1,89 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" +#include +#include +#include + +using boost::int128::int128_t; +using namespace boost::int128::literals; + +// Exercises the signed user-defined literal operators inside a SYCL kernel and checks the +// device-produced values against the same literals evaluated on the host. +int main() +{ + constexpr int N {16}; + + sycl::queue q; + std::cout << "SYCL device: " << q.get_device().get_info() << "\n"; + + int128_t* out {sycl::malloc_shared(N, q)}; + + q.submit([&](sycl::handler& h) + { + h.parallel_for(sycl::range<1>(N), [=](sycl::id<1> idx) + { + switch (idx[0]) + { + case 0: out[0] = 0_i128; break; + case 1: out[1] = 1_i128; break; + case 2: out[2] = -1_i128; break; + case 3: out[3] = 9223372036854775807_i128; break; + case 4: out[4] = 0_I128; break; + case 5: out[5] = 1_I128; break; + case 6: out[6] = -42_I128; break; + case 7: out[7] = 42_I128; break; + case 8: out[8] = "0"_i128; break; + case 9: out[9] = "170141183460469231731687303715884105727"_i128; break; + case 10: out[10] = "-170141183460469231731687303715884105728"_i128; break; + case 11: out[11] = "-999999999999999999"_i128; break; + case 12: out[12] = "0"_I128; break; + case 13: out[13] = "12345678901234567890"_I128; break; + case 14: out[14] = "-12345678901234567890"_I128; break; + case 15: out[15] = "-1"_I128; break; + default: break; + } + }); + }).wait(); + + int128_t expected[N]; + expected[0] = 0_i128; + expected[1] = 1_i128; + expected[2] = -1_i128; + expected[3] = 9223372036854775807_i128; + expected[4] = 0_I128; + expected[5] = 1_I128; + expected[6] = -42_I128; + expected[7] = 42_I128; + expected[8] = "0"_i128; + expected[9] = "170141183460469231731687303715884105727"_i128; + expected[10] = "-170141183460469231731687303715884105728"_i128; + expected[11] = "-999999999999999999"_i128; + expected[12] = "0"_I128; + expected[13] = "12345678901234567890"_I128; + expected[14] = "-12345678901234567890"_I128; + expected[15] = "-1"_I128; + + int failures {0}; + for (int i {0}; i < N; ++i) + { + if (out[i] != expected[i]) + { + std::cerr << "Mismatch at literal case " << i << "\n"; + ++failures; + } + } + + sycl::free(out, q); + + if (failures == 0) + { + std::cout << "Test PASSED\n"; + return EXIT_SUCCESS; + } + + std::cerr << "Test FAILED with " << failures << " mismatches\n"; + return EXIT_FAILURE; +} diff --git a/test/test_signed_lt_sycl.cpp b/test/test_signed_lt_sycl.cpp new file mode 100644 index 00000000..0c0d5a75 --- /dev/null +++ b/test/test_signed_lt_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run_compare( + [](int128_t a, int128_t b) { return a < b; }); +} diff --git a/test/test_signed_midpoint_sycl.cpp b/test/test_signed_midpoint_sycl.cpp new file mode 100644 index 00000000..d350bfe5 --- /dev/null +++ b/test/test_signed_midpoint_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return boost::int128::midpoint(a, b); }); +} diff --git a/test/test_signed_mod_sycl.cpp b/test/test_signed_mod_sycl.cpp new file mode 100644 index 00000000..1b022a29 --- /dev/null +++ b/test/test_signed_mod_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return int128_sycl_test::safe_mod(a, b); }); +} diff --git a/test/test_signed_mul_sat_sycl.cpp b/test/test_signed_mul_sat_sycl.cpp new file mode 100644 index 00000000..920e07ca --- /dev/null +++ b/test/test_signed_mul_sat_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return boost::int128::mul_sat(a, b); }); +} diff --git a/test/test_signed_mul_sycl.cpp b/test/test_signed_mul_sycl.cpp new file mode 100644 index 00000000..d04d886c --- /dev/null +++ b/test/test_signed_mul_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return a * b; }); +} diff --git a/test/test_signed_ne_sycl.cpp b/test/test_signed_ne_sycl.cpp new file mode 100644 index 00000000..5fce2bf3 --- /dev/null +++ b/test/test_signed_ne_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run_compare( + [](int128_t a, int128_t b) { return a != b; }); +} diff --git a/test/test_signed_not_sycl.cpp b/test/test_signed_not_sycl.cpp new file mode 100644 index 00000000..c846f1cf --- /dev/null +++ b/test/test_signed_not_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t, int) { return ~a; }); +} diff --git a/test/test_signed_or_sycl.cpp b/test/test_signed_or_sycl.cpp new file mode 100644 index 00000000..e15cb187 --- /dev/null +++ b/test/test_signed_or_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return a | b; }); +} diff --git a/test/test_signed_right_shift_sycl.cpp b/test/test_signed_right_shift_sycl.cpp new file mode 100644 index 00000000..1ebf1b56 --- /dev/null +++ b/test/test_signed_right_shift_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t, int i) { return a >> (i % 128); }); +} diff --git a/test/test_signed_sub_sat_sycl.cpp b/test/test_signed_sub_sat_sycl.cpp new file mode 100644 index 00000000..c4f90682 --- /dev/null +++ b/test/test_signed_sub_sat_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return boost::int128::sub_sat(a, b); }); +} diff --git a/test/test_signed_sub_sycl.cpp b/test/test_signed_sub_sycl.cpp new file mode 100644 index 00000000..f0b18f31 --- /dev/null +++ b/test/test_signed_sub_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return a - b; }); +} diff --git a/test/test_signed_to_chars_bases_sycl.cpp b/test/test_signed_to_chars_bases_sycl.cpp new file mode 100644 index 00000000..b64c9114 --- /dev/null +++ b/test/test_signed_to_chars_bases_sycl.cpp @@ -0,0 +1,17 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + int result {0}; + result |= int128_sycl_test::run_to_chars(2); + result |= int128_sycl_test::run_to_chars(8); + result |= int128_sycl_test::run_to_chars(16); + result |= int128_sycl_test::run_to_chars(36); + return result; +} diff --git a/test/test_signed_to_chars_sycl.cpp b/test/test_signed_to_chars_sycl.cpp new file mode 100644 index 00000000..176b5e95 --- /dev/null +++ b/test/test_signed_to_chars_sycl.cpp @@ -0,0 +1,12 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run_to_chars(10); +} diff --git a/test/test_signed_to_unsigned_conversion_sycl.cpp b/test/test_signed_to_unsigned_conversion_sycl.cpp new file mode 100644 index 00000000..093d06b1 --- /dev/null +++ b/test/test_signed_to_unsigned_conversion_sycl.cpp @@ -0,0 +1,14 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t, int) { return static_cast(a); }); +} diff --git a/test/test_signed_xor_sycl.cpp b/test/test_signed_xor_sycl.cpp new file mode 100644 index 00000000..8687865d --- /dev/null +++ b/test/test_signed_xor_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; + +int main() +{ + return int128_sycl_test::run( + [](int128_t a, int128_t b, int) { return a ^ b; }); +} diff --git a/test/test_unsigned_add_sat_sycl.cpp b/test/test_unsigned_add_sat_sycl.cpp new file mode 100644 index 00000000..8157fd84 --- /dev/null +++ b/test/test_unsigned_add_sat_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return boost::int128::add_sat(a, b); }); +} diff --git a/test/test_unsigned_add_sycl.cpp b/test/test_unsigned_add_sycl.cpp new file mode 100644 index 00000000..85e35937 --- /dev/null +++ b/test/test_unsigned_add_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return a + b; }); +} diff --git a/test/test_unsigned_and_sycl.cpp b/test/test_unsigned_and_sycl.cpp new file mode 100644 index 00000000..0f04a181 --- /dev/null +++ b/test/test_unsigned_and_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return a & b; }); +} diff --git a/test/test_unsigned_cstdlib_div_sycl.cpp b/test/test_unsigned_cstdlib_div_sycl.cpp new file mode 100644 index 00000000..706261ab --- /dev/null +++ b/test/test_unsigned_cstdlib_div_sycl.cpp @@ -0,0 +1,17 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) + { + const auto d {boost::int128::div(a, int128_sycl_test::safe_divisor(a, b))}; + return d.quot ^ d.rem; + }); +} diff --git a/test/test_unsigned_div_sat_sycl.cpp b/test/test_unsigned_div_sat_sycl.cpp new file mode 100644 index 00000000..b7fdd633 --- /dev/null +++ b/test/test_unsigned_div_sat_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return boost::int128::div_sat(a, int128_sycl_test::safe_divisor(a, b)); }); +} diff --git a/test/test_unsigned_div_sycl.cpp b/test/test_unsigned_div_sycl.cpp new file mode 100644 index 00000000..54b993dd --- /dev/null +++ b/test/test_unsigned_div_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return int128_sycl_test::safe_div(a, b); }); +} diff --git a/test/test_unsigned_eq_sycl.cpp b/test/test_unsigned_eq_sycl.cpp new file mode 100644 index 00000000..41dc174c --- /dev/null +++ b/test/test_unsigned_eq_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run_compare( + [](uint128_t a, uint128_t b) { return a == b; }); +} diff --git a/test/test_unsigned_from_chars_bases_sycl.cpp b/test/test_unsigned_from_chars_bases_sycl.cpp new file mode 100644 index 00000000..6b16e642 --- /dev/null +++ b/test/test_unsigned_from_chars_bases_sycl.cpp @@ -0,0 +1,17 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + int result {0}; + result |= int128_sycl_test::run_from_chars(2); + result |= int128_sycl_test::run_from_chars(8); + result |= int128_sycl_test::run_from_chars(16); + result |= int128_sycl_test::run_from_chars(36); + return result; +} diff --git a/test/test_unsigned_from_chars_sycl.cpp b/test/test_unsigned_from_chars_sycl.cpp new file mode 100644 index 00000000..d37f1af5 --- /dev/null +++ b/test/test_unsigned_from_chars_sycl.cpp @@ -0,0 +1,12 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run_from_chars(10); +} diff --git a/test/test_unsigned_gcd_sycl.cpp b/test/test_unsigned_gcd_sycl.cpp new file mode 100644 index 00000000..58b821b0 --- /dev/null +++ b/test/test_unsigned_gcd_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return boost::int128::gcd(a, b); }); +} diff --git a/test/test_unsigned_ge_sycl.cpp b/test/test_unsigned_ge_sycl.cpp new file mode 100644 index 00000000..45a96fe3 --- /dev/null +++ b/test/test_unsigned_ge_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run_compare( + [](uint128_t a, uint128_t b) { return a >= b; }); +} diff --git a/test/test_unsigned_gt_sycl.cpp b/test/test_unsigned_gt_sycl.cpp new file mode 100644 index 00000000..1525b9c5 --- /dev/null +++ b/test/test_unsigned_gt_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run_compare( + [](uint128_t a, uint128_t b) { return a > b; }); +} diff --git a/test/test_unsigned_lcm_sycl.cpp b/test/test_unsigned_lcm_sycl.cpp new file mode 100644 index 00000000..ca8e958c --- /dev/null +++ b/test/test_unsigned_lcm_sycl.cpp @@ -0,0 +1,14 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + // Restrict inputs so the true LCM stays within 128 bits. + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return boost::int128::lcm(a >> 64, b >> 64); }); +} diff --git a/test/test_unsigned_le_sycl.cpp b/test/test_unsigned_le_sycl.cpp new file mode 100644 index 00000000..4f094f07 --- /dev/null +++ b/test/test_unsigned_le_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run_compare( + [](uint128_t a, uint128_t b) { return a <= b; }); +} diff --git a/test/test_unsigned_left_shift_sycl.cpp b/test/test_unsigned_left_shift_sycl.cpp new file mode 100644 index 00000000..7580f4e1 --- /dev/null +++ b/test/test_unsigned_left_shift_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int i) { return a << (i % 128); }); +} diff --git a/test/test_unsigned_literals_sycl.cpp b/test/test_unsigned_literals_sycl.cpp new file mode 100644 index 00000000..dcbf6809 --- /dev/null +++ b/test/test_unsigned_literals_sycl.cpp @@ -0,0 +1,105 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" +#include +#include +#include + +using boost::int128::uint128_t; +using namespace boost::int128::literals; + +// Exercises the user-defined literal operators inside a SYCL kernel and checks the +// device-produced values against the same literals evaluated on the host. +int main() +{ + constexpr int N {24}; + + sycl::queue q; + std::cout << "SYCL device: " << q.get_device().get_info() << "\n"; + + uint128_t* out {sycl::malloc_shared(N, q)}; + + q.submit([&](sycl::handler& h) + { + h.parallel_for(sycl::range<1>(N), [=](sycl::id<1> idx) + { + switch (idx[0]) + { + case 0: out[0] = 0_u128; break; + case 1: out[1] = 1_u128; break; + case 2: out[2] = 18446744073709551615_u128; break; + case 3: out[3] = 999999999999999999_u128; break; + case 4: out[4] = 0_U128; break; + case 5: out[5] = 1_U128; break; + case 6: out[6] = 18446744073709551615_U128; break; + case 7: out[7] = 999999999999999999_U128; break; + case 8: out[8] = "0"_u128; break; + case 9: out[9] = "1"_u128; break; + case 10: out[10] = "340282366920938463463374607431768211455"_u128; break; + case 11: out[11] = "999999999999999999"_u128; break; + case 12: out[12] = "0"_U128; break; + case 13: out[13] = "1"_U128; break; + case 14: out[14] = "340282366920938463463374607431768211455"_U128; break; + case 15: out[15] = "999999999999999999"_U128; break; + case 16: out[16] = 0_u128; break; + case 17: out[17] = 1_u128; break; + case 18: out[18] = 18446744073709551615_u128; break; + case 19: out[19] = 42_u128; break; + case 20: out[20] = 0_U128; break; + case 21: out[21] = 1_U128; break; + case 22: out[22] = 18446744073709551615_U128; break; + case 23: out[23] = 42_U128; break; + default: break; + } + }); + }).wait(); + + uint128_t expected[N]; + expected[0] = 0_u128; + expected[1] = 1_u128; + expected[2] = 18446744073709551615_u128; + expected[3] = 999999999999999999_u128; + expected[4] = 0_U128; + expected[5] = 1_U128; + expected[6] = 18446744073709551615_U128; + expected[7] = 999999999999999999_U128; + expected[8] = "0"_u128; + expected[9] = "1"_u128; + expected[10] = "340282366920938463463374607431768211455"_u128; + expected[11] = "999999999999999999"_u128; + expected[12] = "0"_U128; + expected[13] = "1"_U128; + expected[14] = "340282366920938463463374607431768211455"_U128; + expected[15] = "999999999999999999"_U128; + expected[16] = 0_u128; + expected[17] = 1_u128; + expected[18] = 18446744073709551615_u128; + expected[19] = 42_u128; + expected[20] = 0_U128; + expected[21] = 1_U128; + expected[22] = 18446744073709551615_U128; + expected[23] = 42_U128; + + int failures {0}; + for (int i {0}; i < N; ++i) + { + if (out[i] != expected[i]) + { + std::cerr << "Mismatch at literal case " << i << "\n"; + ++failures; + } + } + + sycl::free(out, q); + + if (failures == 0) + { + std::cout << "Test PASSED\n"; + return EXIT_SUCCESS; + } + + std::cerr << "Test FAILED with " << failures << " mismatches\n"; + return EXIT_FAILURE; +} diff --git a/test/test_unsigned_lt_sycl.cpp b/test/test_unsigned_lt_sycl.cpp new file mode 100644 index 00000000..daaa9a91 --- /dev/null +++ b/test/test_unsigned_lt_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run_compare( + [](uint128_t a, uint128_t b) { return a < b; }); +} diff --git a/test/test_unsigned_midpoint_sycl.cpp b/test/test_unsigned_midpoint_sycl.cpp new file mode 100644 index 00000000..7d9d4c4e --- /dev/null +++ b/test/test_unsigned_midpoint_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return boost::int128::midpoint(a, b); }); +} diff --git a/test/test_unsigned_mod_sycl.cpp b/test/test_unsigned_mod_sycl.cpp new file mode 100644 index 00000000..0d032003 --- /dev/null +++ b/test/test_unsigned_mod_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return int128_sycl_test::safe_mod(a, b); }); +} diff --git a/test/test_unsigned_mul_sat_sycl.cpp b/test/test_unsigned_mul_sat_sycl.cpp new file mode 100644 index 00000000..86505059 --- /dev/null +++ b/test/test_unsigned_mul_sat_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return boost::int128::mul_sat(a, b); }); +} diff --git a/test/test_unsigned_mul_sycl.cpp b/test/test_unsigned_mul_sycl.cpp new file mode 100644 index 00000000..7b08ecf3 --- /dev/null +++ b/test/test_unsigned_mul_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return a * b; }); +} diff --git a/test/test_unsigned_ne_sycl.cpp b/test/test_unsigned_ne_sycl.cpp new file mode 100644 index 00000000..18af865b --- /dev/null +++ b/test/test_unsigned_ne_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run_compare( + [](uint128_t a, uint128_t b) { return a != b; }); +} diff --git a/test/test_unsigned_not_sycl.cpp b/test/test_unsigned_not_sycl.cpp new file mode 100644 index 00000000..e359a9e2 --- /dev/null +++ b/test/test_unsigned_not_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int) { return ~a; }); +} diff --git a/test/test_unsigned_or_sycl.cpp b/test/test_unsigned_or_sycl.cpp new file mode 100644 index 00000000..a19267a7 --- /dev/null +++ b/test/test_unsigned_or_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return a | b; }); +} diff --git a/test/test_unsigned_right_shift_sycl.cpp b/test/test_unsigned_right_shift_sycl.cpp new file mode 100644 index 00000000..300fdaae --- /dev/null +++ b/test/test_unsigned_right_shift_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int i) { return a >> (i % 128); }); +} diff --git a/test/test_unsigned_sub_sat_sycl.cpp b/test/test_unsigned_sub_sat_sycl.cpp new file mode 100644 index 00000000..1e731c77 --- /dev/null +++ b/test/test_unsigned_sub_sat_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return boost::int128::sub_sat(a, b); }); +} diff --git a/test/test_unsigned_sub_sycl.cpp b/test/test_unsigned_sub_sycl.cpp new file mode 100644 index 00000000..503d285c --- /dev/null +++ b/test/test_unsigned_sub_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return a - b; }); +} diff --git a/test/test_unsigned_to_chars_bases_sycl.cpp b/test/test_unsigned_to_chars_bases_sycl.cpp new file mode 100644 index 00000000..01066aca --- /dev/null +++ b/test/test_unsigned_to_chars_bases_sycl.cpp @@ -0,0 +1,17 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + int result {0}; + result |= int128_sycl_test::run_to_chars(2); + result |= int128_sycl_test::run_to_chars(8); + result |= int128_sycl_test::run_to_chars(16); + result |= int128_sycl_test::run_to_chars(36); + return result; +} diff --git a/test/test_unsigned_to_chars_sycl.cpp b/test/test_unsigned_to_chars_sycl.cpp new file mode 100644 index 00000000..d4dd01f7 --- /dev/null +++ b/test/test_unsigned_to_chars_sycl.cpp @@ -0,0 +1,12 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run_to_chars(10); +} diff --git a/test/test_unsigned_to_signed_conversion_sycl.cpp b/test/test_unsigned_to_signed_conversion_sycl.cpp new file mode 100644 index 00000000..e7c5aac2 --- /dev/null +++ b/test/test_unsigned_to_signed_conversion_sycl.cpp @@ -0,0 +1,14 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::int128_t; +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t, int) { return static_cast(a); }); +} diff --git a/test/test_unsigned_xor_sycl.cpp b/test/test_unsigned_xor_sycl.cpp new file mode 100644 index 00000000..30f9ea1d --- /dev/null +++ b/test/test_unsigned_xor_sycl.cpp @@ -0,0 +1,13 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +using boost::int128::uint128_t; + +int main() +{ + return int128_sycl_test::run( + [](uint128_t a, uint128_t b, int) { return a ^ b; }); +} From ae18b131fde26916db99edf6fe9c714d070a71b2 Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Thu, 9 Jul 2026 10:33:02 -0400 Subject: [PATCH 7/9] Add SYCL testing path to test CML --- test/CMakeLists.txt | 17 ++++++++++++++++- 1 file changed, 16 insertions(+), 1 deletion(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index e601babb..effb4ae4 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -20,7 +20,22 @@ if(HAVE_BOOST_TEST) enable_testing() boost_test_jamfile(FILE cuda_jamfile LINK_LIBRARIES Boost::int128 Boost::random Boost::charconv CUDA::cudart COMPILE_DEFINITIONS BOOST_INT128_ENABLE_CUDA=1 ) - + + elseif(BOOST_INT128_ENABLE_SYCL) + + message(STATUS "Building Boost.int128 with SYCL") + + set(CMAKE_CXX_COMPILER "icpx") + set(CMAKE_C_COMPILER "icx") + + # -fsycl must be on the link line too (it bundles the device image); COMPILE_OPTIONS + # only reaches the compile step, so add it to the linker flags as well. + set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} -fsycl") + + enable_testing() + + boost_test_jamfile(FILE sycl_jamfile LINK_LIBRARIES Boost::int128 Boost::charconv sycl COMPILE_DEFINITIONS BOOST_INT128_ENABLE_SYCL=1 COMPILE_OPTIONS -fsycl) + else() boost_test_jamfile(FILE Jamfile LINK_LIBRARIES Boost::int128 Boost::core Boost::random Boost::multiprecision Boost::mp11 Boost::charconv) From 864570f5d576d4c36237553000d8a66642242ca4 Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Thu, 9 Jul 2026 10:33:18 -0400 Subject: [PATCH 8/9] Add SCYL testing path in CI like Math and Charconv --- .github/workflows/ci.yml | 62 ++++++++++++++++++++++++++++++++++++++++ 1 file changed, 62 insertions(+) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index b0889369..9ee23c78 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -1211,3 +1211,65 @@ jobs: ctest --output-on-failure --no-tests=error cd $GITHUB_WORKSPACE done + + sycl-cmake-test: + strategy: + fail-fast: false + + runs-on: ubuntu-latest + + steps: + - name: Intel Apt repository + timeout-minutes: 1 + run: | + wget https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS-2023.PUB + sudo apt-key add GPG-PUB-KEY-INTEL-SW-PRODUCTS-2023.PUB + rm GPG-PUB-KEY-INTEL-SW-PRODUCTS-2023.PUB + echo "deb https://apt.repos.intel.com/oneapi all main" | sudo tee /etc/apt/sources.list.d/oneAPI.list + sudo apt-get update + - name: Install Intel oneAPI compilers + timeout-minutes: 5 + run: sudo apt-get install -y intel-oneapi-compiler-fortran intel-oneapi-compiler-dpcpp-cpp + + - name: Setup Intel oneAPI environment + run: | + source /opt/intel/oneapi/setvars.sh + printenv >> $GITHUB_ENV + + - uses: actions/checkout@v5 + + - name: Install Packages + run: | + sudo apt-get install -y cmake make + - name: Setup Boost + run: | + echo GITHUB_REPOSITORY: $GITHUB_REPOSITORY + LIBRARY=${GITHUB_REPOSITORY#*/} + echo LIBRARY: $LIBRARY + echo "LIBRARY=$LIBRARY" >> $GITHUB_ENV + echo GITHUB_BASE_REF: $GITHUB_BASE_REF + echo GITHUB_REF: $GITHUB_REF + REF=${GITHUB_BASE_REF:-$GITHUB_REF} + REF=${REF#refs/heads/} + echo REF: $REF + BOOST_BRANCH=develop && [ "$REF" == "master" ] && BOOST_BRANCH=master || true + echo BOOST_BRANCH: $BOOST_BRANCH + cd .. + git clone -b $BOOST_BRANCH --depth 1 https://github.com/boostorg/boost.git boost-root + cd boost-root + mkdir -p libs/$LIBRARY + cp -r $GITHUB_WORKSPACE/* libs/$LIBRARY + git submodule update --init tools/boostdep + python3 tools/boostdep/depinst/depinst.py --git_args "--jobs 3" $LIBRARY + - name: Test C++17/20/23 + run: | + for std in 17 20 23; do + echo "======== Testing C++${std} ========" + cd ../boost-root + rm -rf __build__ + mkdir __build__ && cd __build__ + cmake -DBOOST_INCLUDE_LIBRARIES=$LIBRARY -DBUILD_TESTING=ON -DBOOST_INT128_ENABLE_SYCL=ON -DCMAKE_CXX_COMPILER=icpx -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_STANDARD=${std} .. + cmake --build . --target tests -j $(nproc) + ctest --output-on-failure --no-tests=error + cd $GITHUB_WORKSPACE + done From 9782646996ccad60187b1003bad225e02d8fb7b7 Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Thu, 9 Jul 2026 10:33:37 -0400 Subject: [PATCH 9/9] Add SYCL support documentation and example --- doc/modules/ROOT/nav.adoc | 1 + doc/modules/ROOT/pages/config.adoc | 12 +++- doc/modules/ROOT/pages/examples.adoc | 19 +++++ doc/modules/ROOT/pages/getting_started.adoc | 4 ++ doc/modules/ROOT/pages/overview.adoc | 2 +- examples/sycl.cpp | 78 +++++++++++++++++++++ 6 files changed, 113 insertions(+), 3 deletions(-) create mode 100644 examples/sycl.cpp diff --git a/doc/modules/ROOT/nav.adoc b/doc/modules/ROOT/nav.adoc index 5934b755..3779768a 100644 --- a/doc/modules/ROOT/nav.adoc +++ b/doc/modules/ROOT/nav.adoc @@ -17,6 +17,7 @@ ** xref:examples.adoc#examples_fmt_format[pass:[{fmt}] Integration] ** xref:examples.adoc#examples_cstdlib[`` support (Combined div and mod)] ** xref:examples.adoc#examples_cuda[Use of the library in a CUDA kernel] +** xref:examples.adoc#examples_sycl[Use of the library in a SYCL kernel] * xref:api_reference.adoc[] ** xref:api_reference.adoc#api_namespaces[Namespaces] ** xref:api_reference.adoc#api_types[Types] diff --git a/doc/modules/ROOT/pages/config.adoc b/doc/modules/ROOT/pages/config.adoc index bb36ed59..007c42a5 100644 --- a/doc/modules/ROOT/pages/config.adoc +++ b/doc/modules/ROOT/pages/config.adoc @@ -18,6 +18,12 @@ These macros allow customization of library behavior. User-configurable macros s - `BOOST_INT128_ENABLE_CUDA`: Defining this macro allows both types and selected functions to be run on both host and device when compiling with NVCC. Allowed functions have `BOOST_INT128_HOST_DEVICE` as part of their function signature in their documentation. +[#enable_sycl] +- `BOOST_INT128_ENABLE_SYCL`: Defining this macro allows both types and selected functions to be run on the device when compiling with a SYCL compiler (for example Intel oneAPI `icpx -fsycl`). +Unlike CUDA, SYCL device support is opt-in and is not detected automatically, so this macro must be defined and `` must be included before any Boost.int128 header (so that `SYCL_EXTERNAL` is available). +Device-enabled functions have `BOOST_INT128_HOST_DEVICE` as part of their function signature in their documentation. +The `boost::charconv::to_chars` / `from_chars` overloads in `` also run on the SYCL device: defining `BOOST_INT128_ENABLE_SYCL` forwards to `BOOST_CHARCONV_ENABLE_SYCL`, and Boost.Charconv provides matching SYCL device support (as it does for CUDA). + [#no_int128] - `BOOST_INT128_NO_BUILTIN_INT128`: The user may define this when they do not want the internal implementations to rely on builtin `pass:[__int128]` or `pass:[unsigned __int128]` types. @@ -41,6 +47,8 @@ See xref:getting_started.adoc[Getting Started]. - `BOOST_INT128_HAS_MSVC_INT128`: This is defined when the MSVC standard library provides the simulated 128-bit integer types in `pass:[<__msvc_int128.hpp>]` (MSVC in pass:[C++20] mode). When defined, the library uses those types as its native 128-bit backing. +- `BOOST_INT128_HAS_GPU_SUPPORT`: This is defined when the library is compiled with device support enabled, that is under either `BOOST_INT128_ENABLE_CUDA` (with NVCC) or `BOOST_INT128_ENABLE_SYCL`. It is the single check for whether a GPU device target is in play. Note that on the SYCL device pass the library falls back to its portable path because the spir64 target has no native `pass:[__int128]`, so `BOOST_INT128_HAS_INT128` is not defined there. + - `BOOST_INT128_HAS_SPACESHIP_OPERATOR`: This is defined when the compiler and standard library provide pass:[C++20] three-way comparison (`operator<=>` and ``). When defined, both types provide `operator<=>`. - `BOOST_INT128_HAS_FORMAT`: This is defined by `` when pass:[C++20] `` is available. When defined, that header provides `std::formatter` specializations for both types. @@ -48,5 +56,5 @@ See xref:getting_started.adoc[Getting Started]. - `BOOST_INT128_HAS_FMT_FORMAT`: This is defined by `` when the pass:[{fmt}] library is available. When defined, that header provides `fmt::formatter` specializations for both types. [#host_device] -- `BOOST_INT128_HOST_DEVICE`: This is defined to `pass:[__host__ __device__]` when compiling with NVCC (`pass:[__NVCC__]` is defined), and to nothing otherwise. -The core type operations (constructors, conversion operators, and the comparison, bitwise, arithmetic, and shift operators) and most free functions are annotated with this macro, allowing `int128_t` and `uint128_t` to be used in CUDA device code. Host-only integrations that depend on host facilities are not annotated: `std::hash`, the iostream operators, the `std::formatter` / pass:[{fmt}] support, and the `core::string_view` overloads of `from_chars`. +- `BOOST_INT128_HOST_DEVICE`: This is defined to `pass:[__host__ __device__]` when compiling for CUDA (`pass:[__CUDACC__]` with `BOOST_INT128_ENABLE_CUDA`), to `SYCL_EXTERNAL` when compiling for SYCL (`BOOST_INT128_ENABLE_SYCL`), and to nothing otherwise. +The core type operations (constructors, conversion operators, and the comparison, bitwise, arithmetic, and shift operators) and most free functions are annotated with this macro, allowing `int128_t` and `uint128_t` to be used in CUDA and SYCL device code. Host-only integrations that depend on host facilities are not annotated: `std::hash`, the iostream operators, the `std::formatter` / pass:[{fmt}] support, and the `core::string_view` overloads of `from_chars`. Under SYCL the `operator long double` conversion is additionally unavailable because the spir64 device target has no `long double`. diff --git a/doc/modules/ROOT/pages/examples.adoc b/doc/modules/ROOT/pages/examples.adoc index ab880458..17abcdbd 100644 --- a/doc/modules/ROOT/pages/examples.adoc +++ b/doc/modules/ROOT/pages/examples.adoc @@ -562,3 +562,22 @@ CUDA kernel launch with 196 blocks of 256 threads All CPU and GPU computed elements match! ---- ==== + +[#examples_sycl] +== SYCL Usage + +.This https://github.com/cppalliance/int128/blob/develop/examples/sycl.cpp[example] demonstrates how to use library types and functions inside a SYCL kernel. Define `BOOST_INT128_ENABLE_SYCL` and include `` before any Boost.int128 header, then compile with `icpx -fsycl`. +==== +[source, c++] +---- +include::example$sycl.cpp[] +---- + +.Expected Output +[listing] +---- +[Vector operation on 50000 elements] +SYCL device: Intel(R) ... +All CPU and GPU computed elements match! +---- +==== diff --git a/doc/modules/ROOT/pages/getting_started.adoc b/doc/modules/ROOT/pages/getting_started.adoc index 4fd9cf69..392c2ca6 100644 --- a/doc/modules/ROOT/pages/getting_started.adoc +++ b/doc/modules/ROOT/pages/getting_started.adoc @@ -83,6 +83,10 @@ Building the module requires defining `BOOST_INT128_BUILD_MODULE` and compiling The library types and many of the functions can run on both host and device under CUDA. Compile with a CUDA-aware toolchain and define `BOOST_INT128_ENABLE_CUDA`. Functions carrying `BOOST_INT128_HOST_DEVICE` in their signature run on both host and device; all others run on host only. See xref:config.adoc[Configuration Macros]. +== SYCL Support + +The library types and many of the functions can also run on the device under SYCL (for example Intel oneAPI `icpx -fsycl`). SYCL device support is opt-in: include `` before any Boost.int128 header and define `BOOST_INT128_ENABLE_SYCL`. Functions carrying `BOOST_INT128_HOST_DEVICE` in their signature run on both host and device; all others run on host only. The `boost::charconv::to_chars` / `from_chars` overloads from `` also run on the SYCL device (Boost.Charconv provides SYCL support, and defining `BOOST_INT128_ENABLE_SYCL` enables it automatically). See xref:config.adoc[Configuration Macros]. + == Dependencies The core library (everything included by ``) has no required dependencies and needs only a conforming pass:[C++14] compiler. The optional integration headers depend on the following, and are only usable when those components are available: diff --git a/doc/modules/ROOT/pages/overview.adoc b/doc/modules/ROOT/pages/overview.adoc index eb785908..64aaa172 100644 --- a/doc/modules/ROOT/pages/overview.adoc +++ b/doc/modules/ROOT/pages/overview.adoc @@ -25,7 +25,7 @@ GCC and Clang offer `__int128` as a non-standard extension on 64-bit targets, bu Multiprecision libraries can fill the gap, but typically at the cost of a larger `sizeof` and additional overhead (e.g., Boost.Multiprecision always has an extra word). Boost.Int128 solves this by providing types that are exactly 128-bits on every platform. Operation implementations rely on compiler intrinsics where available for native performance, and optimized software implementations elsewhere. -The types provided by the library also natively support running on GPUs using CUDA, along with many of the functions. +The types provided by the library also natively support running on GPUs using CUDA and SYCL, along with many of the functions. == Use Cases diff --git a/examples/sycl.cpp b/examples/sycl.cpp new file mode 100644 index 00000000..5b6f5a58 --- /dev/null +++ b/examples/sycl.cpp @@ -0,0 +1,78 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include +#include +#include +#include + +using test_type = boost::int128::uint128_t; + +// Calculates the GCD of two values on the SYCL device and verifies against the host +int main() +{ + std::mt19937_64 rng {42}; + + const int numElements {50000}; + std::cout << "[Vector operation on " << numElements << " elements]" << std::endl; + + sycl::queue q; + std::cout << "SYCL device: " << q.get_device().get_info() << std::endl; + + // Allocate shared (USM) memory so the arrays are usable on both host and device + test_type* in1 {sycl::malloc_shared(numElements, q)}; + test_type* in2 {sycl::malloc_shared(numElements, q)}; + test_type* out {sycl::malloc_shared(numElements, q)}; + + for (int i {0}; i < numElements; ++i) + { + in1[i] = test_type{rng(), rng()}; + in2[i] = test_type{rng(), rng()}; + } + + // Launch the SYCL kernel: each work item computes one gcd + q.submit([&](sycl::handler& h) + { + h.parallel_for(sycl::range<1>(numElements), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + out[i] = boost::int128::gcd(in1[i], in2[i]); + }); + }).wait(); + + // Perform the same operation on the host and compare + std::vector results; + results.reserve(numElements); + for (int i {0}; i < numElements; ++i) + { + results.emplace_back(boost::int128::gcd(in1[i], in2[i])); + } + + int ret {EXIT_SUCCESS}; + for (int i {0}; i < numElements; ++i) + { + if (out[i] != results[i]) + { + std::cerr << "Result verification failed at element: " << i << "!" << std::endl; + ret = EXIT_FAILURE; + break; + } + } + + if (ret == EXIT_SUCCESS) + { + std::cout << "All CPU and GPU computed elements match!" << std::endl; + } + + sycl::free(in1, q); + sycl::free(in2, q); + sycl::free(out, q); + + return ret; +}