diff --git a/include/boost/int128/detail/mini_from_chars.hpp b/include/boost/int128/detail/mini_from_chars.hpp
index dfdf9b4f..b5023105 100644
--- a/include/boost/int128/detail/mini_from_chars.hpp
+++ b/include/boost/int128/detail/mini_from_chars.hpp
@@ -22,6 +22,9 @@ namespace int128 {
 namespace detail {
 
 namespace impl {
+
+#ifndef __NVCC__
+
 BOOST_INT128_INLINE_CONSTEXPR unsigned char uchar_values[] =
      {255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255,
       255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255,
@@ -42,9 +45,35 @@ BOOST_INT128_INLINE_CONSTEXPR unsigned char uchar_values[] =
 
 static_assert(sizeof(uchar_values) == 256, "uchar_values should represent all 256 values of unsigned char");
 
+#endif // __NVCC__
+
 // 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
 {
+    #ifdef __NVCC__
+
+    constexpr unsigned char uchar_values[] =
+    {255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255,
+     255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255,
+     255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255,
+       0,   1,   2,   3,   4,   5,   6,   7,   8,   9, 255, 255, 255, 255, 255, 255,
+     255,  10,  11,  12,  13,  14,  15,  16,  17,  18,  19,  20,  21,  22,  23,  24,
+      25,  26,  27,  28,  29,  30,  31,  32,  33,  34,  35, 255, 255, 255, 255, 255,
+     255,  10,  11,  12,  13,  14,  15,  16,  17,  18,  19,  20,  21,  22,  23,  24,
+      25,  26,  27,  28,  29,  30,  31,  32,  33,  34,  35, 255, 255, 255, 255, 255,
+     255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255,
+     255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255,
+     255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255,
+     255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255,
+     255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255,
+     255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255,
+     255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255,
+     255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255, 255};
+
+    static_assert(sizeof(uchar_values) == 256, "uchar_values should represent all 256 values of unsigned char");
+
+    #endif // __NVCC__
+
     return uchar_values[static_cast<unsigned char>(val)];
 }
 
diff --git a/test/cuda_jamfile b/test/cuda_jamfile
index dab3ff9d..1a67d48d 100644
--- a/test/cuda_jamfile
+++ b/test/cuda_jamfile
@@ -89,3 +89,6 @@ run test_unsigned_to_chars_bases.cu ;
 run test_signed_to_chars_bases.cu ;
 run test_unsigned_from_chars_bases.cu ;
 run test_signed_from_chars_bases.cu ;
+
+run test_unsigned_literals.cu ;
+run test_signed_literals.cu ;
diff --git a/test/test_signed_literals.cu b/test/test_signed_literals.cu
new file mode 100644
index 00000000..8723b06c
--- /dev/null
+++ b/test/test_signed_literals.cu
@@ -0,0 +1,161 @@
+//  Copyright Matt Borland 2024 - 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)
+
+#define BOOST_INT128_ALLOW_SIGN_CONVERSION
+
+#include <iostream>
+#include <iomanip>
+#include <boost/int128.hpp>
+#include <boost/int128/literals.hpp>
+#include "cuda_managed_ptr.hpp"
+#include "stopwatch.hpp"
+
+// For the CUDA runtime routines (prefixed with "cuda_")
+#include <cuda_runtime.h>
+
+using boost::int128::int128_t;
+using namespace boost::int128::literals;
+
+// Number of test cases: we test each literal operator with several values
+// Operators:
+//   1. operator""_i128(const char*)              - raw literal, decimal
+//   2. operator""_I128(const char*)              - raw literal, decimal
+//   3. operator""_i128(const char*, size_t)      - cooked string literal
+//   4. operator""_I128(const char*, size_t)      - cooked string literal
+//   5. operator""_i128(unsigned long long)        - integer literal
+//   6. operator""_I128(unsigned long long)        - integer literal
+
+constexpr int NUM_TESTS = 30;
+
+__global__ void cuda_test(int128_t *out)
+{
+    int i = threadIdx.x;
+
+    // operator""_i128(const char*) - raw literal (values must fit unsigned long long to avoid NVCC warnings)
+    if (i == 0) { out[i] = 0_i128; }
+    if (i == 1) { out[i] = 1_i128; }
+    if (i == 2) { out[i] = 18446744073709551615_i128; }
+    if (i == 3) { out[i] = 999999999999999999_i128; }
+    if (i == 4) { out[i] = 42_i128; }
+
+    // operator""_I128(const char*) - raw literal (values must fit unsigned long long to avoid NVCC warnings)
+    if (i == 5) { out[i] = 0_I128; }
+    if (i == 6) { out[i] = 1_I128; }
+    if (i == 7) { out[i] = 18446744073709551615_I128; }
+    if (i == 8) { out[i] = 999999999999999999_I128; }
+    if (i == 9) { out[i] = 42_I128; }
+
+    // operator""_i128(const char*, size_t) - string literal (supports negative)
+    if (i == 10) { out[i] = "0"_i128; }
+    if (i == 11) { out[i] = "1"_i128; }
+    if (i == 12) { out[i] = "170141183460469231731687303715884105727"_i128; }
+    if (i == 13) { out[i] = "-1"_i128; }
+    if (i == 14) { out[i] = "-170141183460469231731687303715884105727"_i128; }
+
+    // operator""_I128(const char*, size_t) - string literal (supports negative)
+    if (i == 15) { out[i] = "0"_I128; }
+    if (i == 16) { out[i] = "1"_I128; }
+    if (i == 17) { out[i] = "170141183460469231731687303715884105727"_I128; }
+    if (i == 18) { out[i] = "-1"_I128; }
+    if (i == 19) { out[i] = "-170141183460469231731687303715884105727"_I128; }
+
+    // operator""_i128(unsigned long long) - integer literal
+    if (i == 20) { out[i] = 0_i128; }
+    if (i == 21) { out[i] = 1_i128; }
+    if (i == 22) { out[i] = 18446744073709551615_i128; }
+    if (i == 23) { out[i] = 42_i128; }
+    if (i == 24) { out[i] = 100_i128; }
+
+    // operator""_I128(unsigned long long) - integer literal
+    if (i == 25) { out[i] = 0_I128; }
+    if (i == 26) { out[i] = 1_I128; }
+    if (i == 27) { out[i] = 18446744073709551615_I128; }
+    if (i == 28) { out[i] = 42_I128; }
+    if (i == 29) { out[i] = 100_I128; }
+}
+
+int main(void)
+{
+    cudaError_t err = cudaSuccess;
+
+    std::cout << "[Signed literal tests: " << NUM_TESTS << " cases]" << std::endl;
+
+    cuda_managed_ptr<int128_t> output(NUM_TESTS);
+
+    // Launch with 1 block of NUM_TESTS threads
+    watch w;
+
+    cuda_test<<<1, NUM_TESTS>>>(output.get());
+    cudaDeviceSynchronize();
+
+    std::cout << "CUDA kernal done in: " << w.elapsed() << "s" << std::endl;
+
+    err = cudaGetLastError();
+    if (err != cudaSuccess)
+    {
+        std::cerr << "Failed to launch kernel (error code " << cudaGetErrorString(err) << ")!" << std::endl;
+        return EXIT_FAILURE;
+    }
+
+    // Build expected values on host using the same literals
+    int128_t expected[NUM_TESTS];
+
+    // operator""_i128(const char*) - raw literal (values must fit unsigned long long to avoid NVCC warnings)
+    expected[0]  = 0_i128;
+    expected[1]  = 1_i128;
+    expected[2]  = 18446744073709551615_i128;
+    expected[3]  = 999999999999999999_i128;
+    expected[4]  = 42_i128;
+
+    // operator""_I128(const char*) - raw literal (values must fit unsigned long long to avoid NVCC warnings)
+    expected[5]  = 0_I128;
+    expected[6]  = 1_I128;
+    expected[7]  = 18446744073709551615_I128;
+    expected[8]  = 999999999999999999_I128;
+    expected[9]  = 42_I128;
+
+    // operator""_i128(const char*, size_t) - string literal (supports negative)
+    expected[10] = "0"_i128;
+    expected[11] = "1"_i128;
+    expected[12] = "170141183460469231731687303715884105727"_i128;
+    expected[13] = "-1"_i128;
+    expected[14] = "-170141183460469231731687303715884105727"_i128;
+
+    // operator""_I128(const char*, size_t) - string literal (supports negative)
+    expected[15] = "0"_I128;
+    expected[16] = "1"_I128;
+    expected[17] = "170141183460469231731687303715884105727"_I128;
+    expected[18] = "-1"_I128;
+    expected[19] = "-170141183460469231731687303715884105727"_I128;
+
+    // operator""_i128(unsigned long long) - integer literal
+    expected[20] = 0_i128;
+    expected[21] = 1_i128;
+    expected[22] = 18446744073709551615_i128;
+    expected[23] = 42_i128;
+    expected[24] = 100_i128;
+
+    // operator""_I128(unsigned long long) - integer literal
+    expected[25] = 0_I128;
+    expected[26] = 1_I128;
+    expected[27] = 18446744073709551615_I128;
+    expected[28] = 42_I128;
+    expected[29] = 100_I128;
+
+    // Verify
+    for (int i = 0; i < NUM_TESTS; ++i)
+    {
+        if (output[i] != expected[i])
+        {
+            std::cerr << "Result verification failed at test case " << i << "!" << std::endl;
+            return EXIT_FAILURE;
+        }
+    }
+
+    std::cout << "Test PASSED" << std::endl;
+    std::cout << "Done\n";
+
+    return 0;
+}
diff --git a/test/test_unsigned_literals.cu b/test/test_unsigned_literals.cu
new file mode 100644
index 00000000..d3fad8c3
--- /dev/null
+++ b/test/test_unsigned_literals.cu
@@ -0,0 +1,149 @@
+//  Copyright Matt Borland 2024 - 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)
+
+#define BOOST_INT128_ALLOW_SIGN_CONVERSION
+
+#include <iostream>
+#include <iomanip>
+#include <boost/int128.hpp>
+#include <boost/int128/literals.hpp>
+#include "cuda_managed_ptr.hpp"
+#include "stopwatch.hpp"
+
+// For the CUDA runtime routines (prefixed with "cuda_")
+#include <cuda_runtime.h>
+
+using boost::int128::uint128_t;
+using namespace boost::int128::literals;
+
+// Number of test cases: we test each literal operator with several values
+// Operators:
+//   1. operator""_u128(const char*)              - raw literal, decimal
+//   2. operator""_U128(const char*)              - raw literal, decimal
+//   3. operator""_u128(const char*, size_t)      - cooked string literal
+//   4. operator""_U128(const char*, size_t)      - cooked string literal
+//   5. operator""_u128(unsigned long long)        - integer literal
+//   6. operator""_U128(unsigned long long)        - integer literal
+
+constexpr int NUM_TESTS = 24;
+
+__global__ void cuda_test(uint128_t *out)
+{
+    int i = threadIdx.x;
+
+    // operator""_u128(const char*) - raw literal
+    if (i == 0) { out[i] = 0_u128; }
+    if (i == 1) { out[i] = 1_u128; }
+    if (i == 2) { out[i] = 18446744073709551615_u128; }
+    if (i == 3) { out[i] = 999999999999999999_u128; }
+
+    // operator""_U128(const char*) - raw literal
+    if (i == 4) { out[i] = 0_U128; }
+    if (i == 5) { out[i] = 1_U128; }
+    if (i == 6) { out[i] = 18446744073709551615_U128; }
+    if (i == 7) { out[i] = 999999999999999999_U128; }
+
+    // operator""_u128(const char*, size_t) - string literal
+    if (i == 8)  { out[i] = "0"_u128; }
+    if (i == 9)  { out[i] = "1"_u128; }
+    if (i == 10) { out[i] = "340282366920938463463374607431768211455"_u128; }
+    if (i == 11) { out[i] = "999999999999999999"_u128; }
+
+    // operator""_U128(const char*, size_t) - string literal
+    if (i == 12) { out[i] = "0"_U128; }
+    if (i == 13) { out[i] = "1"_U128; }
+    if (i == 14) { out[i] = "340282366920938463463374607431768211455"_U128; }
+    if (i == 15) { out[i] = "999999999999999999"_U128; }
+
+    // operator""_u128(unsigned long long) - integer literal
+    if (i == 16) { out[i] = 0_u128; }
+    if (i == 17) { out[i] = 1_u128; }
+    if (i == 18) { out[i] = 18446744073709551615_u128; }
+    if (i == 19) { out[i] = 42_u128; }
+
+    // operator""_U128(unsigned long long) - integer literal
+    if (i == 20) { out[i] = 0_U128; }
+    if (i == 21) { out[i] = 1_U128; }
+    if (i == 22) { out[i] = 18446744073709551615_U128; }
+    if (i == 23) { out[i] = 42_U128; }
+}
+
+int main(void)
+{
+    cudaError_t err = cudaSuccess;
+
+    std::cout << "[Unsigned literal tests: " << NUM_TESTS << " cases]" << std::endl;
+
+    cuda_managed_ptr<uint128_t> output(NUM_TESTS);
+
+    // Launch with 1 block of NUM_TESTS threads
+    watch w;
+
+    cuda_test<<<1, NUM_TESTS>>>(output.get());
+    cudaDeviceSynchronize();
+
+    std::cout << "CUDA kernal done in: " << w.elapsed() << "s" << std::endl;
+
+    err = cudaGetLastError();
+    if (err != cudaSuccess)
+    {
+        std::cerr << "Failed to launch kernel (error code " << cudaGetErrorString(err) << ")!" << std::endl;
+        return EXIT_FAILURE;
+    }
+
+    // Build expected values on host using the same literals
+    uint128_t expected[NUM_TESTS];
+
+    // operator""_u128(const char*) - raw literal
+    expected[0]  = 0_u128;
+    expected[1]  = 1_u128;
+    expected[2]  = 18446744073709551615_u128;
+    expected[3]  = 999999999999999999_u128;
+
+    // operator""_U128(const char*) - raw literal
+    expected[4]  = 0_U128;
+    expected[5]  = 1_U128;
+    expected[6]  = 18446744073709551615_U128;
+    expected[7]  = 999999999999999999_U128;
+
+    // operator""_u128(const char*, size_t) - string literal
+    expected[8]  = "0"_u128;
+    expected[9]  = "1"_u128;
+    expected[10] = "340282366920938463463374607431768211455"_u128;
+    expected[11] = "999999999999999999"_u128;
+
+    // operator""_U128(const char*, size_t) - string literal
+    expected[12] = "0"_U128;
+    expected[13] = "1"_U128;
+    expected[14] = "340282366920938463463374607431768211455"_U128;
+    expected[15] = "999999999999999999"_U128;
+
+    // operator""_u128(unsigned long long) - integer literal
+    expected[16] = 0_u128;
+    expected[17] = 1_u128;
+    expected[18] = 18446744073709551615_u128;
+    expected[19] = 42_u128;
+
+    // operator""_U128(unsigned long long) - integer literal
+    expected[20] = 0_U128;
+    expected[21] = 1_U128;
+    expected[22] = 18446744073709551615_U128;
+    expected[23] = 42_U128;
+
+    // Verify
+    for (int i = 0; i < NUM_TESTS; ++i)
+    {
+        if (output[i] != expected[i])
+        {
+            std::cerr << "Result verification failed at test case " << i << "!" << std::endl;
+            return EXIT_FAILURE;
+        }
+    }
+
+    std::cout << "Test PASSED" << std::endl;
+    std::cout << "Done\n";
+
+    return 0;
+}