[libcu++] Implement P3793R2 Better Shifting (without SIMD)

libcudacxx/include/cuda/std/__bit/shl.h

@@ -0,0 +1,91 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_STD___BIT_SHL_H
+#define _CUDA_STD___BIT_SHL_H
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/__cmath/uabs.h>
+#include <cuda/std/__concepts/concept_macros.h>
+#include <cuda/std/__type_traits/is_integer.h>
+#include <cuda/std/__type_traits/is_signed.h>
+#include <cuda/std/__type_traits/make_nbit_int.h>
+#include <cuda/std/__type_traits/make_unsigned.h>
+#include <cuda/std/__type_traits/num_bits.h>
+#include <cuda/std/__utility/cmp.h>
+#include <cuda/std/cstdint>
+
+#if _CCCL_CUDA_COMPILATION() && !_CCCL_TILE_COMPILATION()
+#  include <cuda/__ptx/instructions/shl.h>
+#  include <cuda/__ptx/instructions/shr.h>
+#endif // _CCCL_CUDA_COMPILATION() && !_CCCL_TILE_COMPILATION()
+
+#include <cuda/std/__cccl/prologue.h>
+
+_CCCL_BEGIN_NAMESPACE_CUDA_STD
+
+_CCCL_TEMPLATE(class _Tp, class _Shift)
+_CCCL_REQUIRES(__cccl_is_integer_v<_Tp> _CCCL_AND __cccl_is_integer_v<_Shift>)
+[[nodiscard]] _CCCL_API constexpr _Tp shl(const _Tp __v, const _Shift __shift) noexcept
+{
+  constexpr auto __width = uint32_t{__num_bits_v<_Tp>};
+  const auto __ushift    = ::cuda::uabs(__shift);
+
+  if constexpr (is_signed_v<_Shift>)
+  {
+    if (__shift < 0)
+    {
+#if !_CCCL_TILE_COMPILATION()
+      _CCCL_IF_NOT_CONSTEVAL_DEFAULT
+      {
+        // On device, shr PTX instruction clamps the shift to width, however only 32-bit shifts are supported.
+        NV_IF_TARGET(NV_IS_DEVICE, ({
+                       if constexpr (sizeof(_Shift) <= sizeof(uint32_t) && sizeof(_Tp) <= sizeof(int64_t))
+                       {
+                         using _Up = __make_nbit_int_t<sizeof(_Tp) < sizeof(int64_t) ? 32 : 64, is_signed_v<_Tp>>;
+                         return static_cast<_Tp>(::cuda::ptx::shr(_Up{__v}, static_cast<uint32_t>(__ushift)));
+                       }
+                     }))
+      }
+#endif // !_CCCL_TILE_COMPILATION()
+      return (__ushift < __width) ? (__v >> __ushift) : static_cast<_Tp>(::cuda::std::cmp_less(__v, 0) ? -1 : 0);
+    }
+  }
+
+#if !_CCCL_TILE_COMPILATION() // error: asm statement is unsupported in tile code
+  _CCCL_IF_NOT_CONSTEVAL_DEFAULT
+  {
+    // On device, shl PTX instruction clamps the shift to width, however only 32-bit shifts are supported.
+    NV_IF_TARGET(NV_IS_DEVICE, ({
+                   if constexpr (sizeof(_Shift) <= sizeof(uint32_t) && sizeof(_Tp) <= sizeof(int64_t))
+                   {
+                     using _Up = __make_nbit_int_t<sizeof(_Tp) < sizeof(int64_t) ? 32 : 64, is_signed_v<_Tp>>;
+                     return static_cast<_Tp>(::cuda::ptx::shl(_Up{__v}, static_cast<uint32_t>(__ushift)));
+                   }
+                 }))
+  }
+#endif // !_CCCL_TILE_COMPILATION()
+  return (__ushift < __width) ? static_cast<_Tp>(::cuda::std::__to_unsigned_like(__v) << __ushift) : _Tp{0};
+}
+
+_CCCL_END_NAMESPACE_CUDA_STD
+
+#include <cuda/std/__cccl/epilogue.h>
+
+#endif // _CUDA_STD___BIT_SHL_H

libcudacxx/include/cuda/std/__bit/shr.h

@@ -0,0 +1,91 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_STD___BIT_SHR_H
+#define _CUDA_STD___BIT_SHR_H
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/__cmath/uabs.h>
+#include <cuda/std/__concepts/concept_macros.h>
+#include <cuda/std/__type_traits/is_integer.h>
+#include <cuda/std/__type_traits/is_signed.h>
+#include <cuda/std/__type_traits/make_nbit_int.h>
+#include <cuda/std/__type_traits/make_unsigned.h>
+#include <cuda/std/__type_traits/num_bits.h>
+#include <cuda/std/__utility/cmp.h>
+#include <cuda/std/cstdint>
+
+#if _CCCL_CUDA_COMPILATION() && !_CCCL_TILE_COMPILATION()
+#  include <cuda/__ptx/instructions/shl.h>
+#  include <cuda/__ptx/instructions/shr.h>
+#endif // _CCCL_CUDA_COMPILATION() && !_CCCL_TILE_COMPILATION()
+
+#include <cuda/std/__cccl/prologue.h>
+
+_CCCL_BEGIN_NAMESPACE_CUDA_STD
+
+_CCCL_TEMPLATE(class _Tp, class _Shift)
+_CCCL_REQUIRES(__cccl_is_integer_v<_Tp> _CCCL_AND __cccl_is_integer_v<_Shift>)
+[[nodiscard]] _CCCL_API constexpr _Tp shr(const _Tp __v, const _Shift __shift) noexcept
+{
+  constexpr auto __width = uint32_t{__num_bits_v<_Tp>};
+  const auto __ushift    = ::cuda::uabs(__shift);
+
+  if constexpr (is_signed_v<_Shift>)
+  {
+    if (__shift < 0)
+    {
+#if !_CCCL_TILE_COMPILATION() // error: asm statement is unsupported in tile code
+      _CCCL_IF_NOT_CONSTEVAL_DEFAULT
+      {
+        // On device, shl PTX instruction clamps the shift to width, however only 32-bit shifts are supported.
+        NV_IF_TARGET(NV_IS_DEVICE, ({
+                       if constexpr (sizeof(_Shift) <= sizeof(uint32_t) && sizeof(_Tp) <= sizeof(int64_t))
+                       {
+                         using _Up = __make_nbit_int_t<sizeof(_Tp) < sizeof(int64_t) ? 32 : 64, is_signed_v<_Tp>>;
+                         return static_cast<_Tp>(::cuda::ptx::shl(_Up{__v}, static_cast<uint32_t>(__ushift)));
+                       }
+                     }))
+      }
+#endif // !_CCCL_TILE_COMPILATION()
+      return (__ushift < __width) ? static_cast<_Tp>(::cuda::std::__to_unsigned_like(__v) << __ushift) : _Tp{0};
+    }
+  }
+
+#if !_CCCL_TILE_COMPILATION() // error: asm statement is unsupported in tile code
+  _CCCL_IF_NOT_CONSTEVAL_DEFAULT
+  {
+    // On device, shr PTX instruction clamps the shift to width, however only 32-bit shifts are supported.
+    NV_IF_TARGET(NV_IS_DEVICE, ({
+                   if constexpr (sizeof(_Shift) <= sizeof(uint32_t) && sizeof(_Tp) <= sizeof(int64_t))
+                   {
+                     using _Up = __make_nbit_int_t<sizeof(_Tp) < sizeof(int64_t) ? 32 : 64, is_signed_v<_Tp>>;
+                     return static_cast<_Tp>(::cuda::ptx::shr(_Up{__v}, static_cast<uint32_t>(__ushift)));
+                   }
+                 }))
+  }
+#endif // !_CCCL_TILE_COMPILATION()
+  return (__ushift < __width) ? (__v >> __ushift) : static_cast<_Tp>(::cuda::std::cmp_less(__v, 0) ? -1 : 0);
+}
+
+_CCCL_END_NAMESPACE_CUDA_STD
+
+#include <cuda/std/__cccl/epilogue.h>
+
+#endif // _CUDA_STD___BIT_SHR_H

libcudacxx/include/cuda/std/bit

@@ -4,7 +4,7 @@
 // under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES.
 //
 //===----------------------------------------------------------------------===//
 
@@ -30,6 +30,8 @@
 #include <cuda/std/__bit/integral.h>
 #include <cuda/std/__bit/popcount.h>
 #include <cuda/std/__bit/rotate.h>
+#include <cuda/std/__bit/shl.h>
+#include <cuda/std/__bit/shr.h>
 #include <cuda/std/version>
 
 #endif // _CUDA_STD_BIT

libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shl.pass.cpp

@@ -0,0 +1,165 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+// template<class T, class S>
+// constexpr T shl(T x, S s) noexcept;
+
+#include <cuda/std/bit>
+#include <cuda/std/cassert>
+#include <cuda/std/cstdint>
+#include <cuda/std/limits>
+#include <cuda/std/numeric>
+#include <cuda/std/type_traits>
+#include <cuda/std/utility>
+
+#include "test_macros.h"
+
+template <class T, class S>
+TEST_FUNC constexpr T invoke_shl(T v, S shift)
+{
+  if (cuda::std::__cccl_default_is_constant_evaluated())
+  {
+    return cuda::std::shl(v, shift);
+  }
+  else
+  {
+    DoNotOptimize(v);
+    DoNotOptimize(shift);
+    return cuda::std::shl(v, shift);
+  }
+}
+
+template <class T, class S>
+TEST_FUNC constexpr void test_pos(T v, S shift)
+{
+  assert(cuda::std::cmp_greater_equal(shift, 0));
+
+  auto result = invoke_shl(v, shift);
+  if (cuda::std::cmp_less(shift, sizeof(T) * CHAR_BIT))
+  {
+    using U = cuda::std::make_unsigned_t<T>;
+    assert(result == static_cast<T>(static_cast<U>(v) << shift));
+  }
+  else
+  {
+    assert(result == T{0});
+  }
+}
+
+template <class T, class S>
+TEST_FUNC constexpr void test_neg(T v, S shift)
+{
+  assert(shift < 0);
+
+  auto result = invoke_shl(v, shift);
+  if (cuda::std::cmp_less(cuda::uabs(shift), (sizeof(T) * CHAR_BIT)))
+  {
+    assert(result == static_cast<T>(v >> cuda::uabs(shift)));
+  }
+  else
+  {
+    assert(result == T(cuda::std::cmp_less(v, 0) ? -1 : 0));
+  }
+}
+
+template <class T, class S>
+TEST_FUNC constexpr void test()
+{
+  constexpr auto tmin = cuda::std::numeric_limits<T>::min();
+  constexpr auto tmax = cuda::std::numeric_limits<T>::max();
+  constexpr auto smax = cuda::std::numeric_limits<S>::max();
+
+  static_assert(cuda::std::is_same_v<T, decltype(cuda::std::shl(T{}, S{}))>);
+  static_assert(noexcept(cuda::std::shl(T{}, S{})));
+
+  const T vs[]         = {tmin, T(-24), T(-1), T{0}, T{1}, T{20}, T{99}, static_cast<T>(1225), tmax};
+  const S pos_shifts[] = {S{0}, S{1}, S{7}, S{17}, S{23}, S{32}, S{33}, S{65}, smax};
+
+  // Disable loop unrolling to reduce ptxas compile times.
+  _CCCL_PRAGMA_NOUNROLL()
+  for (auto v : vs)
+  {
+    _CCCL_PRAGMA_NOUNROLL()
+    for (auto shift : pos_shifts)
+    {
+      test_pos(v, shift);
+    }
+  }
+
+  if constexpr (cuda::std::is_signed_v<S>)
+  {
+    constexpr auto smin = cuda::std::numeric_limits<S>::min();
+
+    const S neg_shifts[] = {smin, S{-65}, S{-33}, S{-32}, S{-15}, S{-7}, S{-1}};
+
+    _CCCL_PRAGMA_NOUNROLL()
+    for (auto v : vs)
+    {
+      _CCCL_PRAGMA_NOUNROLL()
+      for (auto shift : neg_shifts)
+      {
+        test_neg(v, shift);
+      }
+    }
+  }
+}
+
+template <class T>
+TEST_FUNC constexpr void test()
+{
+  test<T, signed char>();
+  test<T, signed short>();
+  test<T, signed>();
+  test<T, signed long>();
+  test<T, signed long long>();
+#if _CCCL_HAS_INT128()
+  test<T, __int128_t>();
+#endif // _CCCL_HAS_INT128()
+
+  test<T, unsigned char>();
+  test<T, unsigned short>();
+  test<T, unsigned>();
+  test<T, unsigned long>();
+  test<T, unsigned long long>();
+#if _CCCL_HAS_INT128()
+  test<T, __uint128_t>();
+#endif // _CCCL_HAS_INT128()
+}
+
+TEST_FUNC constexpr bool test()
+{
+  test<signed char>();
+  test<signed short>();
+  test<signed>();
+  test<signed long>();
+  test<signed long long>();
+#if _CCCL_HAS_INT128()
+  test<__int128_t>();
+#endif // _CCCL_HAS_INT128()
+
+  test<unsigned char>();
+  test<unsigned short>();
+  test<unsigned>();
+  test<unsigned long>();
+  test<unsigned long long>();
+#if _CCCL_HAS_INT128()
+  test<__uint128_t>();
+#endif // _CCCL_HAS_INT128()
+
+  return true;
+}
+
+int main(int, char**)
+{
+  test();
+  static_assert(test());
+
+  return 0;
+}