Add GPU support for SYCL platforms

.github/workflows/ci.yml

@@ -740,3 +740,66 @@ jobs:
       #  run: |
       #    cd ../boost-root/__build__
       #    ctest --output-on-failure --no-tests=error
+
+  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@v6
+
+      - 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: Configure
+        run: |
+          cd ../boost-root
+          mkdir __build__ && cd __build__
+          cmake -DBOOST_INCLUDE_LIBRARIES=$LIBRARY -DBUILD_TESTING=ON -DBOOST_CHARCONV_ENABLE_SYCL=ON -DCMAKE_CXX_COMPILER=icpx -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_STANDARD=17 ..
+      - name: Build tests
+        run: |
+          cd ../boost-root/__build__
+          cmake --build . --target tests -j $(nproc)
+      - name: Run tests
+        run: |
+          cd ../boost-root/__build__
+          ctest --output-on-failure --no-tests=error

doc/charconv/api_reference.adoc

@@ -31,6 +31,7 @@ https://www.boost.org/LICENSE_1_0.txt
 == Macros
 
 - <<enable_cuda_, `BOOST_CHARCONV_ENABLE_CUDA`>>
+- <<enable_sycl_, `BOOST_CHARCONV_ENABLE_SYCL`>>
 - <<enable_cuda_, `BOOST_CHARCONV_HOST_DEVICE`>>
 - <<integral_usage_notes_, `BOOST_CHARCONV_CONSTEXPR`>>
 - <<run_benchmarks_, `BOOST_CHARCONV_RUN_BENCHMARKS`>>

doc/charconv/build.adoc

@@ -42,6 +42,14 @@ IMPORTANT: libquadmath is only available on supported platforms (e.g. Linux with
 This library has partial support for CUDA which can be enabled during compilation with `BOOST_CHARCONV_ENABLE_CUDA`.
 Functions with `BOOST_CHARCONV_HOST_DEVICE` in their signature can be run on both host and device, all others are strictly run on host.
 
+[#enable_sycl_]
+== SYCL Support
+
+This library also has partial support for SYCL, which can be enabled during compilation with `BOOST_CHARCONV_ENABLE_SYCL`.
+Unlike CUDA it is fully opt-in: define the macro and include `<sycl/sycl.hpp>` before any Boost.Charconv header so that `SYCL_EXTERNAL` is available.
+The same functions carrying `BOOST_CHARCONV_HOST_DEVICE` can then be run on both host and device.
+On the SYCL device the builtin 128-bit integer type is not used (the spir64 target has none), so the emulated 128-bit implementation is used instead.
+
 == Dependencies
 
 This library depends on: Boost.Assert, Boost.Config, Boost.Core, and optionally libquadmath (see above).

include/boost/charconv/detail/config.hpp

@@ -19,8 +19,10 @@
 #  define BOOST_CHARCONV_DEBUG_ASSERT(expr)
 #endif
 
-// Use 128-bit integers and suppress warnings for using extensions
-#if defined(BOOST_HAS_INT128) && !(defined(BOOST_CHARCONV_ENABLE_CUDA) && defined(__CUDACC__))
+// Use 128-bit integers and suppress warnings for using extensions.
+// The CUDA and SYCL device paths do not use the builtin 128-bit integer (the SYCL spir64
+// target has none, and the CUDA path is kept portable), so the emulated uint128 is used there.
+#if defined(BOOST_HAS_INT128) && !(defined(BOOST_CHARCONV_ENABLE_CUDA) && defined(__CUDACC__)) && !defined(BOOST_CHARCONV_ENABLE_SYCL)
 #  define BOOST_CHARCONV_HAS_INT128
 #  define BOOST_CHARCONV_INT128_MAX  static_cast<boost::int128_type>((static_cast<boost::uint128_type>(1) << 127) - 1)
 #  define BOOST_CHARCONV_INT128_MIN  (-BOOST_CHARCONV_INT128_MAX - 1)
@@ -205,8 +207,12 @@ static_assert(std::is_same<long double, __float128>::value, "__float128 should b
 
 #endif
 
+// Device function markers. CUDA uses __host__ __device__; SYCL uses SYCL_EXTERNAL and is
+// opt-in via BOOST_CHARCONV_ENABLE_SYCL (define it, and include <sycl/sycl.hpp> beforehand).
 #if defined(BOOST_CHARCONV_ENABLE_CUDA) && defined(__CUDACC__)
 #  define BOOST_CHARCONV_HOST_DEVICE __host__ __device__
+#elif defined(BOOST_CHARCONV_ENABLE_SYCL)
+#  define BOOST_CHARCONV_HOST_DEVICE SYCL_EXTERNAL
 #else
 #  define BOOST_CHARCONV_HOST_DEVICE
 #endif

include/boost/charconv/detail/to_chars_integer_impl.hpp

@@ -343,8 +343,8 @@ BOOST_CHARCONV_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars_int
 
     Unsigned_Integer unsigned_value {};
     const auto unsigned_base = static_cast<Unsigned_Integer>(base);
-
-    BOOST_IF_CONSTEXPR (std::is_signed<Integer>::value)
+    
+    BOOST_IF_CONSTEXPR (std::numeric_limits<Integer>::is_signed)
     {
         if (value < static_cast<Integer>(0))
         {

test/CMakeLists.txt

@@ -18,6 +18,19 @@ if(HAVE_BOOST_TEST)
 
         boost_test_jamfile(FILE cuda_jamfile LINK_LIBRARIES Boost::charconv Boost::core Boost::assert ${CUDA_LIBRARIES} COMPILE_DEFINITIONS BOOST_CHARCONV_ENABLE_CUDA=1 INCLUDE_DIRECTORIES ${CUDA_INCLUDE_DIRS} )
 
+    elseif(BOOST_CHARCONV_ENABLE_SYCL)
+
+        message(STATUS "Building Boost.charconv 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")
+
+        boost_test_jamfile(FILE sycl_jamfile LINK_LIBRARIES Boost::charconv Boost::core Boost::assert sycl COMPILE_DEFINITIONS BOOST_CHARCONV_ENABLE_SYCL=1 COMPILE_OPTIONS -fsycl )
+
     else()
 
         # https://crascit.com/2015/03/28/enabling-cxx11-in-cmake/

test/sycl_charconv_test.hpp

@@ -0,0 +1,202 @@
+//  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.Charconv SYCL tests. Each test runs to_chars/from_chars
+//  element-wise on the SYCL device over random inputs and verifies the device results
+//  against a host recomputation of the same operation.
+
+#ifndef BOOST_CHARCONV_TEST_SYCL_CHARCONV_TEST_HPP
+#define BOOST_CHARCONV_TEST_SYCL_CHARCONV_TEST_HPP
+
+#include <sycl/sycl.hpp>
+#include <boost/charconv.hpp>
+#include <boost/charconv/detail/integer_search_trees.hpp>
+#include <iostream>
+#include <random>
+#include <cstddef>
+#include <cstdint>
+#include <cstdlib>
+#include <cstring>
+
+namespace charconv_sycl_test {
+
+constexpr int num_elements {20000};
+// Large enough for the widest built-in integer in base 2 (64 digits) plus a sign.
+constexpr int buf_size {128};
+
+// to_chars: format each value on the device and compare against the host.
+template <typename T>
+int test_to_chars(const int base = 10)
+{
+    sycl::queue q;
+    std::cout << "SYCL device: " << q.get_device().get_info<sycl::info::device::name>() << "\n";
+
+    T* in {sycl::malloc_shared<T>(num_elements, q)};
+    char* out_str {sycl::malloc_shared<char>(num_elements * buf_size, q)};
+    int* out_len {sycl::malloc_shared<int>(num_elements, q)};
+
+    std::mt19937_64 rng {42};
+    for (int i {0}; i < num_elements; ++i)
+    {
+        in[i] = static_cast<T>(rng());
+    }
+
+    q.submit([&](sycl::handler& h)
+    {
+        h.parallel_for(sycl::range<1>(num_elements), [=](sycl::id<1> idx)
+        {
+            const int i {static_cast<int>(idx[0])};
+            char* buf {out_str + i * buf_size};
+            const auto r {boost::charconv::to_chars(buf, buf + buf_size, in[i], base)};
+            out_len[i] = static_cast<int>(r.ptr - buf);
+        });
+    }).wait();
+
+    int failures {0};
+    for (int i {0}; i < num_elements; ++i)
+    {
+        char buf[buf_size];
+        const auto r {boost::charconv::to_chars(buf, buf + buf_size, in[i], base)};
+        const int n {static_cast<int>(r.ptr - buf)};
+        if (out_len[i] != n || std::memcmp(out_str + i * buf_size, buf, static_cast<std::size_t>(n)) != 0)
+        {
+            if (failures < 5)
+            {
+                std::cerr << "Mismatch at element " << i << " (base " << base << ")\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: format each value on the host, parse it back on the device, verify round trip.
+template <typename T>
+int test_from_chars(const int base = 10)
+{
+    sycl::queue q;
+    std::cout << "SYCL device: " << q.get_device().get_info<sycl::info::device::name>() << "\n";
+
+    char* in_str {sycl::malloc_shared<char>(num_elements * buf_size, q)};
+    int* in_len {sycl::malloc_shared<int>(num_elements, q)};
+    T* out {sycl::malloc_shared<T>(num_elements, q)};
+    T* expected {sycl::malloc_shared<T>(num_elements, q)};
+
+    std::mt19937_64 rng {42};
+    for (int i {0}; i < num_elements; ++i)
+    {
+        expected[i] = static_cast<T>(rng());
+        char buf[buf_size];
+        const auto r {boost::charconv::to_chars(buf, buf + buf_size, expected[i], base)};
+        const int n {static_cast<int>(r.ptr - buf)};
+        std::memcpy(in_str + i * buf_size, buf, static_cast<std::size_t>(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<int>(idx[0])};
+            const char* buf {in_str + i * buf_size};
+            T value {};
+            boost::charconv::from_chars(buf, buf + 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 << " (base " << base << ")\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;
+}
+
+// num_digits: count digits on the device and compare against the host.
+template <typename T>
+int test_num_digits()
+{
+    sycl::queue q;
+    std::cout << "SYCL device: " << q.get_device().get_info<sycl::info::device::name>() << "\n";
+
+    T* in {sycl::malloc_shared<T>(num_elements, q)};
+    int* out {sycl::malloc_shared<int>(num_elements, q)};
+
+    std::mt19937_64 rng {42};
+    for (int i {0}; i < num_elements; ++i)
+    {
+        in[i] = static_cast<T>(rng());
+    }
+
+    q.submit([&](sycl::handler& h)
+    {
+        h.parallel_for(sycl::range<1>(num_elements), [=](sycl::id<1> idx)
+        {
+            const int i {static_cast<int>(idx[0])};
+            out[i] = boost::charconv::detail::num_digits(in[i]);
+        });
+    }).wait();
+
+    int failures {0};
+    for (int i {0}; i < num_elements; ++i)
+    {
+        if (out[i] != boost::charconv::detail::num_digits(in[i]))
+        {
+            if (failures < 5)
+            {
+                std::cerr << "Mismatch at element " << i << "\n";
+            }
+            ++failures;
+        }
+    }
+
+    sycl::free(in, 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;
+}
+
+} // namespace charconv_sycl_test
+
+#endif // BOOST_CHARCONV_TEST_SYCL_CHARCONV_TEST_HPP