Add SYCL support

.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

doc/modules/ROOT/nav.adoc

@@ -17,6 +17,7 @@
 ** xref:examples.adoc#examples_fmt_format[pass:[{fmt}] Integration]
 ** xref:examples.adoc#examples_cstdlib[`<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]

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 `<sycl/sycl.hpp>` 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 `<boost/int128/charconv.hpp>` 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,12 +47,14 @@ 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 `<compare>`). When defined, both types provide `operator<=>`.
 
 - `BOOST_INT128_HAS_FORMAT`: This is defined by `<boost/int128/format.hpp>` when pass:[C++20] `<format>` is available. When defined, that header provides `std::formatter` specializations for both types.
 
 - `BOOST_INT128_HAS_FMT_FORMAT`: This is defined by `<boost/int128/fmt_format.hpp>` 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`.

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 `<sycl/sycl.hpp>` 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!
+----
+====

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 `<sycl/sycl.hpp>` 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 `<boost/int128/charconv.hpp>` 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 `<boost/int128.hpp>`) 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:

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
 

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 <sycl/sycl.hpp>
+#include <boost/int128.hpp>
+#include <boost/int128/numeric.hpp>
+#include <iostream>
+#include <vector>
+#include <random>
+#include <cstdint>
+#include <cstdlib>
+
+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<sycl::info::device::name>() << std::endl;
+
+    // Allocate shared (USM) memory so the arrays are usable on both host and device
+    test_type* in1 {sycl::malloc_shared<test_type>(numElements, q)};
+    test_type* in2 {sycl::malloc_shared<test_type>(numElements, q)};
+    test_type* out {sycl::malloc_shared<test_type>(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<int>(idx[0])};
+            out[i] = boost::int128::gcd(in1[i], in2[i]);
+        });
+    }).wait();
+
+    // Perform the same operation on the host and compare
+    std::vector<test_type> 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;
+}

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 <boost/int128/int128.hpp>
 #include <boost/int128/literals.hpp>
 #include <boost/charconv.hpp>
@@ -44,23 +48,23 @@ struct make_signed<int128::uint128_t> { using type = int128::int128_t; };
 template <>
 struct make_signed<int128::int128_t> { 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<int128::uint128_t>()
+BOOST_INT128_HOST_DEVICE constexpr int128::uint128_t get_max_value<int128::uint128_t>()
 {
     return std::numeric_limits<int128::uint128_t>::max();
 }
 
 template <>
-__host__ __device__ constexpr int128::int128_t get_max_value<int128::int128_t>()
+BOOST_INT128_HOST_DEVICE constexpr int128::int128_t get_max_value<int128::int128_t>()
 {
     return std::numeric_limits<int128::int128_t>::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)
     {

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<std::uint64_t>(product >> 64U);
         return static_cast<std::uint64_t>(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;

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;
+// <sycl/sycl.hpp> must be included before <boost/int128.hpp> 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

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<double>(high) * detail::offset_value_v<double> + static_cast<double>(low);
 }
 
-#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
+#if !defined(BOOST_INT128_HAS_GPU_SUPPORT)
 
 constexpr int128_t::operator long double() const noexcept
 {