Skip to content

Commit 89bad60

Browse files
mborlandmborland
authored
Merge pull request #364 from cppalliance/cudacc
2 parents 2ea99ac + bd1ccfe commit 89bad60

12 files changed

Lines changed: 45 additions & 32 deletions

File tree

‎doc/modules/ROOT/pages/api_reference.adoc‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -307,6 +307,9 @@ Listed by analogous STL header.
307307

308308
| xref:config.adoc#disable_exceptions[`BOOST_INT128_DISABLE_EXCEPTIONS`]
309309
| Disables exception throwing
310+
311+
| xref:config.adoc#enable_cuda[`BOOST_INT128_ENABLE_CUDA`]
312+
| Enables CUDA support allowing the library types and functions to be run on both host and device
310313
|===
311314

312315
==== Automatic Configuration

‎doc/modules/ROOT/pages/config.adoc‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -17,6 +17,10 @@ https://www.boost.org/LICENSE_1_0.txt
1717

1818
These macros allow customization of library behavior. User-configurable macros should be defined before including any library headers.
1919

20+
[#enable_cuda]
21+
- `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.
22+
Allowed functions have `BOOST_IN128_HOST_DEVICE` as part of their function signature in their documentation.
23+
2024
[#no_int128]
2125
- `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.
2226

‎include/boost/int128/bit.hpp‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -65,7 +65,7 @@ BOOST_INT128_EXPORT BOOST_INT128_HOST_DEVICE constexpr uint128_t rotr(const uint
6565
return x >> (static_cast<unsigned>(s) & mask) | x << (static_cast<unsigned>(-s) & mask);
6666
}
6767

68-
#if BOOST_INT128_HAS_BUILTIN(__builtin_popcountll) && !defined(__NVCC__)
68+
#if BOOST_INT128_HAS_BUILTIN(__builtin_popcountll) && !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
6969

7070
BOOST_INT128_EXPORT BOOST_INT128_HOST_DEVICE constexpr int popcount(const uint128_t x) noexcept
7171
{
@@ -139,7 +139,7 @@ BOOST_INT128_EXPORT BOOST_INT128_HOST_DEVICE constexpr int popcount(const uint12
139139
}
140140
}
141141

142-
#elif !BOOST_INT128_HAS_BUILTIN(__builtin_popcountll) || defined(__NVCC__)
142+
#elif !BOOST_INT128_HAS_BUILTIN(__builtin_popcountll) || (defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
143143

144144
BOOST_INT128_EXPORT BOOST_INT128_HOST_DEVICE constexpr int popcount(const uint128_t x) noexcept
145145
{
@@ -148,7 +148,7 @@ BOOST_INT128_EXPORT BOOST_INT128_HOST_DEVICE constexpr int popcount(const uint12
148148

149149
#endif
150150

151-
#if BOOST_INT128_HAS_BUILTIN(__builtin_bswap64) && !defined(__NVCC__)
151+
#if BOOST_INT128_HAS_BUILTIN(__builtin_bswap64) && !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
152152

153153
BOOST_INT128_EXPORT BOOST_INT128_HOST_DEVICE constexpr uint128_t byteswap(const uint128_t x) noexcept
154154
{
@@ -187,7 +187,7 @@ BOOST_INT128_EXPORT BOOST_INT128_HOST_DEVICE constexpr uint128_t byteswap(const
187187
}
188188
}
189189

190-
#elif !BOOST_INT128_HAS_BUILTIN(__builtin_bswap64) || defined(__NVCC__)
190+
#elif !BOOST_INT128_HAS_BUILTIN(__builtin_bswap64) || (defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
191191

192192
BOOST_INT128_EXPORT BOOST_INT128_HOST_DEVICE constexpr uint128_t byteswap(const uint128_t x) noexcept
193193
{

‎include/boost/int128/charconv.hpp‎

Lines changed: 11 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,12 @@
1010

1111
#if __has_include(<boost/charconv.hpp>)
1212

13+
// Define for the user automatically,
14+
// otherwise we'll have an ever-increasing number of these required as we go down the dependency chain
15+
#if defined(BOOST_INT128_ENABLE_CUDA) && !defined(BOOST_CHARCONV_ENABLE_CUDA)
16+
# define BOOST_CHARCONV_ENABLE_CUDA
17+
#endif
18+
1319
#include <boost/int128/int128.hpp>
1420
#include <boost/int128/literals.hpp>
1521
#include <boost/charconv.hpp>
@@ -38,7 +44,7 @@ struct make_signed<int128::uint128_t> { using type = int128::int128_t; };
3844
template <>
3945
struct make_signed<int128::int128_t> { using type = int128::int128_t; };
4046

41-
#ifdef __NVCC__
47+
#if defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)
4248

4349
template <>
4450
__host__ __device__ constexpr int128::uint128_t get_max_value<int128::uint128_t>()
@@ -54,7 +60,7 @@ __host__ __device__ constexpr int128::int128_t get_max_value<int128::int128_t>()
5460

5561
#endif // __NVCC__
5662

57-
#ifndef __NVCC__
63+
#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
5864

5965
BOOST_INT128_INLINE_CONSTEXPR int128::uint128_t int128_pow10[39] =
6066
{
@@ -103,7 +109,7 @@ BOOST_INT128_INLINE_CONSTEXPR int128::uint128_t int128_pow10[39] =
103109

104110
BOOST_INT128_HOST_DEVICE constexpr int num_digits(const int128::uint128_t& x) noexcept
105111
{
106-
#ifdef __NVCC__
112+
#if defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)
107113

108114
constexpr int128::uint128_t int128_pow10[39] =
109115
{
@@ -181,7 +187,7 @@ BOOST_INT128_HOST_DEVICE constexpr int num_digits(const int128::uint128_t& x) no
181187

182188
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
183189
{
184-
#ifndef __NVCC__
190+
#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
185191

186192
if (base == 10)
187193
{
@@ -195,7 +201,7 @@ BOOST_INT128_HOST_DEVICE BOOST_CHARCONV_CONSTEXPR to_chars_result to_chars(char*
195201

196202
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
197203
{
198-
#ifndef __NVCC__
204+
#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
199205

200206
if (base == 10)
201207
{

‎include/boost/int128/detail/clz.hpp‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@ namespace detail {
2020

2121
namespace impl {
2222

23-
#ifndef __NVCC__
23+
#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
2424

2525
// See: http://graphics.stanford.edu/~seander/bithacks.html#IntegerLogDeBruijn
2626
BOOST_INT128_INLINE_CONSTEXPR int index64[64] = {
@@ -38,7 +38,7 @@ BOOST_INT128_INLINE_CONSTEXPR int index64[64] = {
3838

3939
BOOST_INT128_HOST_DEVICE constexpr int bit_scan_reverse(std::uint64_t bb) noexcept
4040
{
41-
#ifdef __NVCC__
41+
#if defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)
4242

4343
constexpr int index64[64] = {
4444
0, 47, 1, 56, 48, 27, 2, 60,
@@ -67,7 +67,7 @@ BOOST_INT128_HOST_DEVICE constexpr int bit_scan_reverse(std::uint64_t bb) noexce
6767
return index64[(bb * debruijn64) >> 58];
6868
}
6969

70-
#ifndef __NVCC__
70+
#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
7171

7272
BOOST_INT128_INLINE_CONSTEXPR int countl_mod37[37] = {
7373
32, 31, 6, 30, 9, 5, 0, 29,
@@ -81,7 +81,7 @@ BOOST_INT128_INLINE_CONSTEXPR int countl_mod37[37] = {
8181

8282
BOOST_INT128_HOST_DEVICE constexpr int backup_countl_impl(std::uint32_t x) noexcept
8383
{
84-
#ifdef __NVCC__
84+
#if defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)
8585

8686
constexpr int countl_mod37[37] = {
8787
32, 31, 6, 30, 9, 5, 0, 29,
@@ -102,7 +102,7 @@ BOOST_INT128_HOST_DEVICE constexpr int backup_countl_impl(std::uint32_t x) noexc
102102
return countl_mod37[x % 37];
103103
}
104104

105-
#if BOOST_INT128_HAS_BUILTIN(__builtin_clz) && !defined(__NVCC__)
105+
#if BOOST_INT128_HAS_BUILTIN(__builtin_clz) && !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
106106

107107
constexpr int countl_impl(unsigned int x) noexcept
108108
{
@@ -119,7 +119,7 @@ constexpr int countl_impl(unsigned long long x) noexcept
119119
return x ? __builtin_clzll(x) : std::numeric_limits<unsigned long long>::digits;
120120
}
121121

122-
#elif (defined(_M_AMD64) || defined(_M_ARM64)) && !defined(BOOST_INT128_NO_CONSTEVAL_DETECTION) && !defined(__NVCC__)
122+
#elif (defined(_M_AMD64) || defined(_M_ARM64)) && !defined(BOOST_INT128_NO_CONSTEVAL_DETECTION) && !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
123123

124124
constexpr int countl_impl(std::uint32_t x) noexcept
125125
{

‎include/boost/int128/detail/config.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -285,7 +285,7 @@ using builtin_u128 = std::_Unsigned128;
285285
# endif
286286
#endif
287287

288-
#ifdef __NVCC__
288+
#if defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)
289289
# define BOOST_INT128_HOST_DEVICE __host__ __device__
290290
#else
291291
# define BOOST_INT128_HOST_DEVICE

‎include/boost/int128/detail/ctz.hpp‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@ namespace detail {
2020

2121
namespace impl {
2222

23-
#if BOOST_INT128_HAS_BUILTIN(__builtin_ctz) && !defined(__NVCC__)
23+
#if BOOST_INT128_HAS_BUILTIN(__builtin_ctz) && !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
2424

2525
constexpr int countr_impl(unsigned int x) noexcept
2626
{
@@ -39,7 +39,7 @@ constexpr int countr_impl(unsigned long long x) noexcept
3939

4040
#endif
4141

42-
#ifndef __NVCC__
42+
#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
4343

4444
BOOST_INT128_INLINE_CONSTEXPR int countr_mod37[37] = {
4545
32, 0, 1, 26, 2, 23, 27, 0,
@@ -79,7 +79,7 @@ constexpr int countr_impl(std::uint32_t x) noexcept
7979

8080
#pragma warning(pop)
8181

82-
#elif !BOOST_INT128_HAS_BUILTIN(__builtin_ctz) || defined(__NVCC__)
82+
#elif !BOOST_INT128_HAS_BUILTIN(__builtin_ctz) || (defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
8383

8484
#ifdef _MSC_VER
8585
#pragma warning(push)
@@ -88,7 +88,7 @@ constexpr int countr_impl(std::uint32_t x) noexcept
8888

8989
BOOST_INT128_HOST_DEVICE constexpr int countr_impl(std::uint32_t x) noexcept
9090
{
91-
#ifdef __NVCC__
91+
#if defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)
9292

9393
constexpr int countr_mod37[37] = {
9494
32, 0, 1, 26, 2, 23, 27, 0,
@@ -109,7 +109,7 @@ BOOST_INT128_HOST_DEVICE constexpr int countr_impl(std::uint32_t x) noexcept
109109

110110
#endif
111111

112-
#if (defined(_M_AMD64) || defined(_M_ARM64)) && !defined(BOOST_INT128_NO_CONSTEVAL_DETECTION) && !BOOST_INT128_HAS_BUILTIN(__builtin_ctz) && !defined(__NVCC__)
112+
#if (defined(_M_AMD64) || defined(_M_ARM64)) && !defined(BOOST_INT128_NO_CONSTEVAL_DETECTION) && !BOOST_INT128_HAS_BUILTIN(__builtin_ctz) && !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
113113

114114
constexpr int countr_impl(std::uint64_t x) noexcept
115115
{
@@ -132,7 +132,7 @@ constexpr int countr_impl(std::uint64_t x) noexcept
132132
}
133133
}
134134

135-
#elif !BOOST_INT128_HAS_BUILTIN(__builtin_ctz) || defined(__NVCC__)
135+
#elif !BOOST_INT128_HAS_BUILTIN(__builtin_ctz) || (defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
136136

137137
BOOST_INT128_HOST_DEVICE constexpr int countr_impl(std::uint64_t x) noexcept
138138
{

‎include/boost/int128/detail/int128_imp.hpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -120,7 +120,7 @@ int128_t
120120
BOOST_INT128_HOST_DEVICE explicit constexpr operator double() const noexcept;
121121

122122
// Long double does not exist on device
123-
#ifndef __NVCC__
123+
#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
124124
explicit constexpr operator long double() const noexcept;
125125
#endif
126126

@@ -297,7 +297,7 @@ BOOST_INT128_HOST_DEVICE constexpr int128_t::operator double() const noexcept
297297
return static_cast<double>(high) * detail::offset_value_v<double> + static_cast<double>(low);
298298
}
299299

300-
#ifndef __NVCC__
300+
#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
301301

302302
constexpr int128_t::operator long double() const noexcept
303303
{
@@ -2266,7 +2266,7 @@ BOOST_INT128_HOST_DEVICE BOOST_INT128_FORCE_INLINE constexpr int128_t library_su
22662266

22672267
BOOST_INT128_HOST_DEVICE BOOST_INT128_FORCE_INLINE constexpr int128_t default_sub(const int128_t lhs, const int128_t rhs) noexcept
22682268
{
2269-
#if defined(BOOST_INT128_HAS_BUILTIN_SUB_OVERFLOW) && (!defined(__aarch64__) || defined(__APPLE__) || !defined(BOOST_INT128_HAS_INT128)) && !defined(__NVCC__)
2269+
#if defined(BOOST_INT128_HAS_BUILTIN_SUB_OVERFLOW) && (!defined(__aarch64__) || defined(__APPLE__) || !defined(BOOST_INT128_HAS_INT128)) && !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
22702270

22712271
// __builtin_sub_overflow is marked constexpr so we don't need if consteval handling
22722272
std::uint64_t result_low {};

‎include/boost/int128/detail/mini_from_chars.hpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -23,7 +23,7 @@ namespace detail {
2323

2424
namespace impl {
2525

26-
#ifndef __NVCC__
26+
#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
2727

2828
BOOST_INT128_INLINE_CONSTEXPR unsigned char uchar_values[] =
2929
{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
5050
// Convert characters for 0-9, A-Z, a-z to 0-35. Anything else is 255
5151
BOOST_INT128_HOST_DEVICE BOOST_INT128_FORCE_INLINE constexpr auto digit_from_char(char val) noexcept -> unsigned char
5252
{
53-
#ifdef __NVCC__
53+
#if defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)
5454

5555
constexpr unsigned char uchar_values[] =
5656
{255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255,

‎include/boost/int128/detail/mini_to_chars.hpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -12,7 +12,7 @@ namespace boost {
1212
namespace int128 {
1313
namespace detail {
1414

15-
#ifndef __NVCC__
15+
#if !(defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA))
1616

1717
BOOST_INT128_INLINE_CONSTEXPR char lower_case_digit_table[] = {
1818
'0', '1', '2', '3', '4', '5', '6', '7', '8', '9',
@@ -32,7 +32,7 @@ static_assert(sizeof(upper_case_digit_table) == sizeof(char) * 16, "10 numbers,
3232

3333
BOOST_INT128_HOST_DEVICE constexpr char* mini_to_chars(char (&buffer)[64], uint128_t v, const int base, const bool uppercase) noexcept
3434
{
35-
#ifdef __NVCC__
35+
#if defined(__CUDACC__) && defined(BOOST_INT128_ENABLE_CUDA)
3636
constexpr char lower_case_digit_table[] = {
3737
'0', '1', '2', '3', '4', '5', '6', '7', '8', '9',
3838
'a', 'b', 'c', 'd', 'e', 'f'

0 commit comments

Comments
 (0)