[libcu++] Implement P3793R2 Better Shifting (without SIMD) by davebayer · Pull Request #9993 · NVIDIA/cccl

davebayer · 2026-07-17T12:56:47Z

This PR implement P3793R2 Better Shifting. I've omitted the simd support, should be done in a separate PR.

coderabbitai · 2026-07-17T13:02:32Z

📝 Walkthrough

Summary by CodeRabbit

New Features
- Added cuda::std::shl and cuda::std::shr constexpr utilities for integer left/right shifts.
- Works with signed/unsigned operands, including negative shift counts and well-defined out-of-range saturation behavior.
- Available through the standard cuda/std/bit umbrella header.
Tests
- Added new pass tests covering many integer types and shift edge cases for both shl and shr.
- Refined the test “do not optimize” barrier to use separate device vs host handling.

Walkthrough

Changes

Adds cuda::std::shl and cuda::std::shr with constexpr integer shifting, signed-shift normalization, width clamping, and CUDA PTX dispatch. Exposes both APIs through cuda/std/bit and adds host/device tests.

Integer shift utilities

Layer / File(s)	Summary
Shift API implementations `libcudacxx/include/cuda/std/__bit/*`, `libcudacxx/include/cuda/std/bit`	Adds constrained `shl` and `shr` functions with constexpr/noexcept behavior, signed-shift handling, out-of-range results, and device PTX paths; exports them from the umbrella header.
shl validation `libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shl.pass.cpp`	Tests positive and negative shifts, width boundaries, constant evaluation, noexcept/type properties, and signed, unsigned, and optional 128-bit integer combinations.
shr validation and device barriers `libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shr.pass.cpp`, `libcudacxx/test/support/test_macros.h`	Tests positive and negative `shr` behavior across integer types and architectures, while separating device and host `DoNotOptimize` assembly constraints.

Suggested reviewers: wmaxey, bernhardmgruber

_{Comment @coderabbitai help to get the list of available commands.}

coderabbitai

Actionable comments posted: 1

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 487bc115-c3d6-4a80-845a-898abd27190e

📥 Commits

Reviewing files that changed from the base of the PR and between 81279b3 and 247d04d.

📒 Files selected for processing (6)

libcudacxx/include/cuda/std/__bit/shl.h
libcudacxx/include/cuda/std/__bit/shr.h
libcudacxx/include/cuda/std/bit
libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shl.pass.cpp
libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shr.pass.cpp
libcudacxx/test/support/test_macros.h

coderabbitai

Actionable comments posted: 1

🧹 Nitpick comments (1)

libcudacxx/test/support/test_macros.h (1)

112-112: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: Make the pointer object const. ptr is not reassigned; use a top-level const while retaining its volatile qualification for the assembly operand.

Source: Coding guidelines

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: edbeeff7-f119-4e50-a24a-968539124e5c

📥 Commits

Reviewing files that changed from the base of the PR and between d62708c and bd57e62.

📒 Files selected for processing (1)

libcudacxx/test/support/test_macros.h

fbusato · 2026-07-17T18:56:34Z

+[[nodiscard]] _CCCL_API constexpr _Tp shl(_Tp __v, _Shift __shift) noexcept
+{
+  constexpr auto __width = static_cast<uint32_t>(__num_bits_v<_Tp>);
+  const auto __ushift    = ::cuda::uabs(__shift);


what about we add an assertion for -min value instead?

fbusato · 2026-07-17T18:56:45Z

+
+_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(_Tp __v, _Shift __shift) noexcept


Suggested change

[[nodiscard]] _CCCL_API constexpr _Tp shl(_Tp __v, _Shift __shift) noexcept

[[nodiscard]] _CCCL_API constexpr _Tp shl(const _Tp __v, const _Shift __shift) noexcept

fbusato · 2026-07-17T18:59:31Z

+      }
+      return (__ushift < __width) ? (__v >> __ushift) : static_cast<_Tp>(::cuda::std::cmp_less(__v, 0) ? -1 : 0);
+    }
+  }


Suggested change

}

}

else {

fbusato · 2026-07-17T19:02:03Z

+_CCCL_REQUIRES(__cccl_is_integer_v<_Tp> _CCCL_AND __cccl_is_integer_v<_Shift>)
+[[nodiscard]] _CCCL_API constexpr _Tp shr(_Tp __v, _Shift __shift) noexcept
+{
+  constexpr auto __width = static_cast<uint32_t>(__num_bits_v<_Tp>);


Suggested change

constexpr auto __width = static_cast<uint32_t>(__num_bits_v<_Tp>);

constexpr auto __width = uint32_t{__num_bits_v<_Tp>};

fbusato · 2026-07-17T19:02:34Z

+
+_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(_Tp __v, _Shift __shift) noexcept


it uses PTX so Tile is not supported

Suggested change

[[nodiscard]] _CCCL_API constexpr _Tp shr(_Tp __v, _Shift __shift) noexcept

[[nodiscard]] _CCCL_HOST_DEVICE_API constexpr _Tp shr(_Tp __v, _Shift __shift) noexcept

fbusato · 2026-07-17T19:03:04Z

+#include <cuda/std/__utility/cmp.h>
+#include <cuda/std/cstdint>
+
+#if _CCCL_CUDA_COMPILATION()


Suggested change

#if _CCCL_CUDA_COMPILATION()

#if _CCCL_CUDA_COMPILATION() && !_CCCL_TILE_COMPILATION()

(or similar)

fbusato · 2026-07-17T19:04:54Z

 // 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.


should we also update __cpp_lib_bitops?

fbusato · 2026-07-17T19:05:27Z

@@ -106,31 +106,23 @@
 #  define TEST_NVRTC_VIRTUAL_DEFAULT_DTOR_ANNOTATION


question. why we need to modify this file?

fbusato · 2026-07-17T19:06:02Z

+
+    const S neg_shifts[] = {smin, S{-65}, S{-33}, S{-32}, S{-15}, S{-7}, S{-1}};
+
+    _CCCL_PRAGMA_NOUNROLL()


I don't think make sense to use pragma unroll in tests

github-actions · 2026-07-17T22:13:04Z

😬 CI Workflow Results

🟥 Finished in 4h 21m: Pass: 89%/120 | Total: 5d 11h | Max: 3h 37m | Hits: 45%/2480803

See results here.

davebayer requested a review from a team as a code owner July 17, 2026 12:56

davebayer requested a review from wmaxey July 17, 2026 12:56

github-project-automation Bot added this to CCCL Jul 17, 2026

github-project-automation Bot moved this to Todo in CCCL Jul 17, 2026

cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Jul 17, 2026

coderabbitai Bot reviewed Jul 17, 2026

View reviewed changes

Comment thread libcudacxx/include/cuda/std/__bit/shl.h Outdated

[libcu++] Implement P3793R2 Better Shifting

dfc3703

davebayer force-pushed the better_shifting branch from 247d04d to dfc3703 Compare July 17, 2026 13:04

davebayer changed the title ~~[libcu++] Implement P3793R2 Better Shifting~~ [libcu++] Implement P3793R2 Better Shifting (without SIMD) Jul 17, 2026

davebayer added 3 commits July 17, 2026 19:04

fixes

d62708c

fix donotoptimize

0aed510

fix typo

bd57e62

coderabbitai Bot reviewed Jul 17, 2026

View reviewed changes

Comment thread libcudacxx/test/support/test_macros.h

coderabbitai Bot mentioned this pull request Jul 17, 2026

Fix cuda::ptx shr and bfind long, long long handling #10000

Open

fbusato reviewed Jul 17, 2026

View reviewed changes

	[[nodiscard]] _CCCL_API constexpr _Tp shl(_Tp __v, _Shift __shift) noexcept
	[[nodiscard]] _CCCL_API constexpr _Tp shl(const _Tp __v, const _Shift __shift) noexcept

	constexpr auto __width = static_cast<uint32_t>(__num_bits_v<_Tp>);
	constexpr auto __width = uint32_t{__num_bits_v<_Tp>};

	[[nodiscard]] _CCCL_API constexpr _Tp shr(_Tp __v, _Shift __shift) noexcept
	[[nodiscard]] _CCCL_HOST_DEVICE_API constexpr _Tp shr(_Tp __v, _Shift __shift) noexcept

	#if _CCCL_CUDA_COMPILATION()
	#if _CCCL_CUDA_COMPILATION() && !_CCCL_TILE_COMPILATION()

		@@ -106,31 +106,23 @@
		# define TEST_NVRTC_VIRTUAL_DEFAULT_DTOR_ANNOTATION


		const S neg_shifts[] = {smin, S{-65}, S{-33}, S{-32}, S{-15}, S{-7}, S{-1}};

		_CCCL_PRAGMA_NOUNROLL()

Uh oh!

Conversation

davebayer commented Jul 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

coderabbitai Bot commented Jul 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary by CodeRabbit

Walkthrough

Changes

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

github-actions Bot commented Jul 17, 2026

😬 CI Workflow Results

🟥 Finished in 4h 21m: Pass: 89%/120 | Total: 5d 11h | Max: 3h 37m | Hits: 45%/2480803

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

davebayer commented Jul 17, 2026 •

edited

Loading

coderabbitai Bot commented Jul 17, 2026 •

edited

Loading