Skip to content

Commit af47854

Browse files
mborlandmborland
authored
Merge pull request #2 from cppalliance/to_chars_2
Add CUDA support to `to_chars` for integers
2 parents ffba31a + a8f7c29 commit af47854

28 files changed

+2480
-18
lines changed

‎include/boost/charconv/detail/apply_sign.hpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -26,13 +26,13 @@ namespace boost { namespace charconv { namespace detail {
2626

2727
template <typename Integer, typename Unsigned_Integer = detail::make_unsigned_t<Integer>,
2828
typename std::enable_if<detail::is_signed<Integer>::value, bool>::type = true>
29-
constexpr Unsigned_Integer apply_sign(Integer val) noexcept
29+
BOOST_CHARCONV_HOST_DEVICE constexpr Unsigned_Integer apply_sign(Integer val) noexcept
3030
{
3131
return -(static_cast<Unsigned_Integer>(val));
3232
}
3333

3434
template <typename Unsigned_Integer, typename std::enable_if<!detail::is_signed<Unsigned_Integer>::value, bool>::type = true>
35-
constexpr Unsigned_Integer apply_sign(Unsigned_Integer val) noexcept
35+
BOOST_CHARCONV_HOST_DEVICE constexpr Unsigned_Integer apply_sign(Unsigned_Integer val) noexcept
3636
{
3737
return val;
3838
}

‎include/boost/charconv/detail/memcpy.hpp‎

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,22 @@
2525

2626
namespace boost { namespace charconv { namespace detail {
2727

28+
#ifdef __NVCC__
29+
30+
__host__ __device__ constexpr char* memcpy(char* dest, const char* src, std::size_t count)
31+
{
32+
for (std::size_t i = 0; i < count; ++i)
33+
{
34+
*(dest + i) = *(src + i);
35+
}
36+
37+
return dest;
38+
}
39+
40+
#define BOOST_CHARCONV_CONSTEXPR constexpr
41+
42+
#else
43+
2844
#if !defined(BOOST_CHARCONV_NO_CONSTEXPR_DETECTION) && defined(BOOST_CXX14_CONSTEXPR)
2945

3046
#define BOOST_CHARCONV_CONSTEXPR constexpr
@@ -69,6 +85,8 @@ inline void* memcpy(void* dest, const void* src, std::size_t count)
6985

7086
#endif
7187

88+
#endif // NVCC
89+
7290
}}} // Namespace boost::charconv::detail
7391

7492
#ifdef BOOST_CHARCONV_STRINGOP_OVERFLOW_DISABLED

‎include/boost/charconv/detail/to_chars_integer_impl.hpp‎

Lines changed: 34 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -52,6 +52,8 @@ static constexpr char radix_table[] = {
5252
'9', '5', '9', '6', '9', '7', '9', '8', '9', '9'
5353
};
5454

55+
#ifndef __NVCC__
56+
5557
static constexpr char digit_table[] = {
5658
'0', '1', '2', '3', '4', '5', '6', '7', '8', '9',
5759
'a', 'b', 'c', 'd', 'e', 'f', 'g', 'h', 'i', 'j',
@@ -301,6 +303,8 @@ BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars_128integer_impl(char* first, c
301303
return {first + converted_value_digits, std::errc()};
302304
}
303305

306+
#endif // __NVCC__
307+
304308
// Conversion warning from shift operators with unsigned char
305309
#if defined(__GNUC__) && __GNUC__ >= 5
306310
# pragma GCC diagnostic push
@@ -313,8 +317,19 @@ BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars_128integer_impl(char* first, c
313317
// All other bases
314318
// Use a simple lookup table to put together the Integer in character form
315319
template <typename Integer, typename Unsigned_Integer>
316-
BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars_integer_impl(char* first, char* last, Integer value, int base) noexcept
320+
BOOST_CHARCONV_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars_integer_impl(char* first, char* last, Integer value, int base) noexcept
317321
{
322+
#ifdef __NVCC__
323+
324+
constexpr char digit_table[] = {
325+
'0', '1', '2', '3', '4', '5', '6', '7', '8', '9',
326+
'a', 'b', 'c', 'd', 'e', 'f', 'g', 'h', 'i', 'j',
327+
'k', 'l', 'm', 'n', 'o', 'p', 'q', 'r', 's', 't',
328+
'u', 'v', 'w', 'x', 'y', 'z'
329+
};
330+
331+
#endif
332+
318333
if (!((first <= last) && (base >= 2 && base <= 36)))
319334
{
320335
return {last, std::errc::invalid_argument};
@@ -381,6 +396,18 @@ BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars_integer_impl(char* first, char
381396
}
382397
break;
383398

399+
#ifdef __NVCC__
400+
401+
case 10:
402+
while (unsigned_value != static_cast<Unsigned_Integer>(0))
403+
{
404+
*end-- = static_cast<char>(zero + (unsigned_value % 10U));
405+
unsigned_value /= 10U;
406+
}
407+
break;
408+
409+
#endif
410+
384411
case 16:
385412
while (unsigned_value != static_cast<Unsigned_Integer>(0))
386413
{
@@ -430,13 +457,18 @@ BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars_integer_impl(char* first, char
430457
#endif
431458

432459
template <typename Integer>
433-
BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars_int(char* first, char* last, Integer value, int base = 10) noexcept
460+
BOOST_CHARCONV_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars_int(char* first, char* last, Integer value, int base = 10) noexcept
434461
{
435462
using Unsigned_Integer = typename std::make_unsigned<Integer>::type;
463+
464+
// The specialized base 10 path requires lookup tables and memcpy
465+
// On device, we instead use the trivial divide and mod to avoid these
466+
#ifndef __NVCC__
436467
if (base == 10)
437468
{
438469
return to_chars_integer_impl(first, last, value);
439470
}
471+
#endif
440472

441473
return to_chars_integer_impl<Integer, Unsigned_Integer>(first, last, value, base);
442474
}

‎include/boost/charconv/detail/to_chars_result.hpp‎

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,7 @@
55
#ifndef BOOST_CHARCONV_DETAIL_TO_CHARS_RESULT_HPP
66
#define BOOST_CHARCONV_DETAIL_TO_CHARS_RESULT_HPP
77

8+
#include <boost/charconv/detail/config.hpp>
89
#include <system_error>
910

1011
// 22.13.2, Primitive numerical output conversion
@@ -16,17 +17,17 @@ struct to_chars_result
1617
char *ptr;
1718
std::errc ec;
1819

19-
constexpr friend bool operator==(const to_chars_result &lhs, const to_chars_result &rhs) noexcept
20+
BOOST_CHARCONV_HOST_DEVICE constexpr friend bool operator==(const to_chars_result &lhs, const to_chars_result &rhs) noexcept
2021
{
2122
return lhs.ptr == rhs.ptr && lhs.ec == rhs.ec;
2223
}
2324

24-
constexpr friend bool operator!=(const to_chars_result &lhs, const to_chars_result &rhs) noexcept
25+
BOOST_CHARCONV_HOST_DEVICE constexpr friend bool operator!=(const to_chars_result &lhs, const to_chars_result &rhs) noexcept
2526
{
2627
return !(lhs == rhs);
2728
}
2829

29-
constexpr explicit operator bool() const noexcept { return ec == std::errc{}; }
30+
BOOST_CHARCONV_HOST_DEVICE constexpr explicit operator bool() const noexcept { return ec == std::errc{}; }
3031
};
3132

3233
}} // Namespaces

‎include/boost/charconv/to_chars.hpp‎

Lines changed: 11 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -17,47 +17,47 @@ namespace charconv {
1717

1818
// integer overloads
1919
BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, bool value, int base) noexcept = delete;
20-
BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, char value, int base = 10) noexcept
20+
BOOST_CHARCONV_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, char value, int base = 10) noexcept
2121
{
2222
return detail::to_chars_int(first, last, value, base);
2323
}
24-
BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, signed char value, int base = 10) noexcept
24+
BOOST_CHARCONV_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, signed char value, int base = 10) noexcept
2525
{
2626
return detail::to_chars_int(first, last, value, base);
2727
}
28-
BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, unsigned char value, int base = 10) noexcept
28+
BOOST_CHARCONV_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, unsigned char value, int base = 10) noexcept
2929
{
3030
return detail::to_chars_int(first, last, value, base);
3131
}
32-
BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, short value, int base = 10) noexcept
32+
BOOST_CHARCONV_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, short value, int base = 10) noexcept
3333
{
3434
return detail::to_chars_int(first, last, value, base);
3535
}
36-
BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, unsigned short value, int base = 10) noexcept
36+
BOOST_CHARCONV_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, unsigned short value, int base = 10) noexcept
3737
{
3838
return detail::to_chars_int(first, last, value, base);
3939
}
40-
BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, int value, int base = 10) noexcept
40+
BOOST_CHARCONV_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, int value, int base = 10) noexcept
4141
{
4242
return detail::to_chars_int(first, last, value, base);
4343
}
44-
BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, unsigned int value, int base = 10) noexcept
44+
BOOST_CHARCONV_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, unsigned int value, int base = 10) noexcept
4545
{
4646
return detail::to_chars_int(first, last, value, base);
4747
}
48-
BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, long value, int base = 10) noexcept
48+
BOOST_CHARCONV_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, long value, int base = 10) noexcept
4949
{
5050
return detail::to_chars_int(first, last, value, base);
5151
}
52-
BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, unsigned long value, int base = 10) noexcept
52+
BOOST_CHARCONV_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, unsigned long value, int base = 10) noexcept
5353
{
5454
return detail::to_chars_int(first, last, value, base);
5555
}
56-
BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, long long value, int base = 10) noexcept
56+
BOOST_CHARCONV_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, long long value, int base = 10) noexcept
5757
{
5858
return detail::to_chars_int(first, last, value, base);
5959
}
60-
BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, unsigned long long value, int base = 10) noexcept
60+
BOOST_CHARCONV_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char* first, char* last, unsigned long long value, int base = 10) noexcept
6161
{
6262
return detail::to_chars_int(first, last, value, base);
6363
}

‎test/cuda_jamfile‎

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -32,3 +32,27 @@ run test_from_chars_bases_long.cu ;
3232
run test_from_chars_bases_unsigned_long.cu ;
3333
run test_from_chars_bases_long_long.cu ;
3434
run test_from_chars_bases_unsigned_long_long.cu ;
35+
36+
run test_to_chars_char.cu ;
37+
run test_to_chars_signed_char.cu ;
38+
run test_to_chars_unsigned_char.cu ;
39+
run test_to_chars_short.cu ;
40+
run test_to_chars_unsigned_short.cu ;
41+
run test_to_chars_int.cu ;
42+
run test_to_chars_unsigned_int.cu ;
43+
run test_to_chars_long.cu ;
44+
run test_to_chars_unsigned_long.cu ;
45+
run test_to_chars_long_long.cu ;
46+
run test_to_chars_unsigned_long_long.cu ;
47+
48+
run test_to_chars_bases_char.cu ;
49+
run test_to_chars_bases_signed_char.cu ;
50+
run test_to_chars_bases_unsigned_char.cu ;
51+
run test_to_chars_bases_short.cu ;
52+
run test_to_chars_bases_unsigned_short.cu ;
53+
run test_to_chars_bases_int.cu ;
54+
run test_to_chars_bases_unsigned_int.cu ;
55+
run test_to_chars_bases_long.cu ;
56+
run test_to_chars_bases_unsigned_long.cu ;
57+
run test_to_chars_bases_long_long.cu ;
58+
run test_to_chars_bases_unsigned_long_long.cu ;

‎test/test_to_chars_bases_char.cu‎

Lines changed: 112 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,112 @@
1+
// Copyright Matt Borland 2024 - 2026.
2+
// Use, modification and distribution are subject to the
3+
// Boost Software License, Version 1.0. (See accompanying file
4+
// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
5+
6+
#include <iostream>
7+
#include <iomanip>
8+
#include <vector>
9+
#include <random>
10+
#include <limits>
11+
#include <cstring>
12+
#include <boost/charconv/to_chars.hpp>
13+
#include "cuda_managed_ptr.hpp"
14+
#include "stopwatch.hpp"
15+
16+
// For the CUDA runtime routines (prefixed with "cuda_")
17+
#include <cuda_runtime.h>
18+
19+
using test_type = char;
20+
21+
constexpr int BUF_SIZE = 128;
22+
23+
__global__ void cuda_test(const test_type *in, char *out_strings, int *out_lengths, int numElements, int base)
24+
{
25+
int i = blockDim.x * blockIdx.x + threadIdx.x;
26+
27+
if (i < numElements)
28+
{
29+
char* buf = out_strings + i * BUF_SIZE;
30+
auto res = boost::charconv::to_chars(buf, buf + BUF_SIZE, in[i], base);
31+
out_lengths[i] = static_cast<int>(res.ptr - buf);
32+
}
33+
}
34+
35+
/**
36+
* Host main routine
37+
*/
38+
int main(void)
39+
{
40+
std::mt19937_64 rng {42};
41+
42+
// Error code to check return values for CUDA calls
43+
cudaError_t err = cudaSuccess;
44+
45+
// Print the vector length to be used, and compute its size
46+
int numElements = 50000;
47+
std::cout << "[Vector operation on " << numElements << " elements]" << std::endl;
48+
49+
// Allocate the managed input vector
50+
cuda_managed_ptr<test_type> input_vector(numElements);
51+
52+
// Allocate the managed output vectors
53+
cuda_managed_ptr<char> output_strings(numElements * BUF_SIZE);
54+
cuda_managed_ptr<int> output_lengths(numElements);
55+
56+
int threadsPerBlock = 256;
57+
int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;
58+
59+
std::uniform_int_distribution<short> dist {(std::numeric_limits<test_type>::min)(), (std::numeric_limits<test_type>::max)()};
60+
61+
for (int base = 2; base <= 36; ++base)
62+
{
63+
// Initialize the input vectors
64+
for (std::size_t i = 0; i < numElements; ++i)
65+
{
66+
input_vector[i] = static_cast<test_type>(dist(rng));
67+
}
68+
69+
// Launch the CUDA Kernel
70+
std::cout << "CUDA kernel launch with " << blocksPerGrid << " blocks of " << threadsPerBlock << " threads, base " << base << std::endl;
71+
72+
watch w;
73+
74+
cuda_test<<<blocksPerGrid, threadsPerBlock>>>(input_vector.get(), output_strings.get(), output_lengths.get(), numElements, base);
75+
cudaDeviceSynchronize();
76+
77+
std::cout << "CUDA kernal done in: " << w.elapsed() << "s" << std::endl;
78+
79+
err = cudaGetLastError();
80+
81+
if (err != cudaSuccess)
82+
{
83+
std::cerr << "Failed to launch kernel (error code " << cudaGetErrorString(err) << ")!" << std::endl;
84+
return EXIT_FAILURE;
85+
}
86+
87+
// Verify that the result vector is correct
88+
w.reset();
89+
for(int i = 0; i < numElements; ++i)
90+
{
91+
char cpu_buf[BUF_SIZE];
92+
auto cpu_res = boost::charconv::to_chars(cpu_buf, cpu_buf + BUF_SIZE, input_vector[i], base);
93+
int cpu_len = static_cast<int>(cpu_res.ptr - cpu_buf);
94+
int gpu_len = output_lengths[i];
95+
const char* gpu_buf = &output_strings[i * BUF_SIZE];
96+
97+
if (cpu_len != gpu_len || std::memcmp(cpu_buf, gpu_buf, static_cast<std::size_t>(cpu_len)) != 0)
98+
{
99+
std::cerr << "Result verification failed at element " << i << " base " << base << "!" << std::endl;
100+
return EXIT_FAILURE;
101+
}
102+
}
103+
double t = w.elapsed();
104+
105+
std::cout << "Test base " << base << " PASSED, normal calculation time: " << t << "s" << std::endl;
106+
}
107+
108+
std::cout << "All bases PASSED" << std::endl;
109+
std::cout << "Done\n";
110+
111+
return 0;
112+
}

0 commit comments

Comments
 (0)