Repository navigation
Add cuda::bit_msb to <cuda/bit> - #9624
Conversation
bit_msb returns the zero-based index of the most significant set bit (floor(log2(value))), or -1 when value is zero. It forwards to cuda::std::__bit_log2 so it reuses the optimal bfind/clz lowering, and exposes that internal primitive type safely on all unsigned integers. Addresses the msb request in NVIDIA#6108. Signed-off-by: temujinkz <ttalkenov@gmail.com>
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
No actionable comments were generated in the recent review.
 |
| Layer / File(s) | Summary |
|---|---|
Header implementation and public wiring libcudacxx/include/cuda/__bit/bit_msb.h, libcudacxx/include/cuda/bit |
Defines bit_msb for unsigned integer types as a wrapper over __bit_log2 and includes it from cuda/bit. |
Constexpr and runtime validation libcudacxx/test/libcudacxx/cuda/bit/bit_msb.pass.cpp |
Checks accepted and rejected types, plus zero, single-bit, and multi-bit results across supported unsigned integer types. |
API reference documentation docs/libcudacxx/extended_api/bit.rst, docs/libcudacxx/extended_api/bit/bit_msb.rst |
Adds the reference page, example, toctree entry, capability description, and CCCL/CUDA availability metadata. |
Possibly related issues
- important: NVIDIA/cccl#11989 — The issue describes a test pattern that can remove runtime coverage on Clang; the new test uses the same
assert(test()); static_assert(test());pattern.
Possibly related PRs
- important: NVIDIA/cccl#12220 — Changes the
__bit_log2device lowering and preserves the zero sentinel used bybit_msb.
Suggested reviewers: pciolkosz
Priority:
Change: Feature
Merge Risk: 
d9585
-
The documentation does not introduce the reported broken link. No issue requiring a fix before merge was established.
Comment @coderabbitai help to get the list of available commands.
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 2
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 172f9c7a-18ac-49a7-bd73-6af7368b9912
📒 Files selected for processing (5)
docs/libcudacxx/extended_api/bit.rstdocs/libcudacxx/extended_api/bit/bit_msb.rstlibcudacxx/include/cuda/__bit/bit_msb.hlibcudacxx/include/cuda/bitlibcudacxx/test/libcudacxx/cuda/bit/bit_msb.pass.cpp
| #include <cuda/bit> | ||
| #include <cuda/std/cassert> | ||
| #include <cuda/std/cstdint> | ||
|
|
||
| __global__ void bit_msb_kernel() { | ||
| assert(cuda::bit_msb(uint32_t{0}) == -1); | ||
| assert(cuda::bit_msb(uint32_t{1}) == 0); | ||
| assert(cuda::bit_msb(uint32_t{0b10101000}) == 7); | ||
| assert(cuda::bit_msb(~uint32_t{0}) == 31); | ||
| } | ||
|
|
||
| int main() { | ||
| bit_msb_kernel<<<1, 1>>>(); | ||
| cudaDeviceSynchronize(); | ||
| return 0; |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
important: Make the example self-contained.
The snippet calls cudaDeviceSynchronize() but never includes a CUDA runtime header, so it is not buildable as written. Add the appropriate runtime include (for example, the project-preferred CUDA runtime API header) or drop the synchronization call from the example. As per path instructions, documentation changes should prioritize buildable examples and technical accuracy.
Source: Path instructions
There was a problem hiding this comment.
@fbusato Sounds good — leaving this for temujinkz to update the example with the missing CUDA runtime include (or drop cudaDeviceSynchronize()).
There was a problem hiding this comment.
Added <cuda_runtime_api.h> to the example includes, same header is_pointer_accessible.rst uses. Applied the same fix to the bit_ffs example in #9609 since it had the identical pattern.
There was a problem hiding this comment.
Seems like the humans are having a chat. I'll hop back into my burrow for now. If you need me again, just tag @coderabbitai in a new comment, and I'll come hopping out!
There was a problem hiding this comment.
Please, remove the include, this header is always included during CUDA compilation by the compiler
|
/ok to test 6d37a1f |
This comment has been minimized.
This comment has been minimized.
fbusato
left a comment
fbusato
left a comment
There was a problem hiding this comment.
looks good. A few minor suggestions
- convert the header comment to a doxygen docstring and take the argument by const value - test: trim the type list to the fundamental unsigned types plus __uint128_t, replace the concept detector with a plain SFINAE variable template, and fix a copy-paste comment - docs: use math notation for the floor(log2) formula, say unsigned integer types in the note, and set the CUDA availability to 13.5
|
/ok to test 88bbfe2 |
This comment has been minimized.
This comment has been minimized.
The bit_ffs page lives in a separate PR, so the :ref: label does not exist on this branch and the docs build treats the dangling reference as an error. Use plain text for now; the link can be restored once both pages are in main.
- assume the result range after __bit_log2 so the optimizer can rely on it - docs: mark the signature __tile__ and include cuda_runtime_api.h in the example so it builds as written
|
Went through everything from the last rounds: doxygen docstring, const argument, math notation, the unsigned wording in the note, CUDA 13.5, the trimmed type list, and the SFINAE detector (one tweak: |
|
/ok to test 295df5d |
This comment has been minimized.
This comment has been minimized.
There was a problem hiding this comment.
Actionable comments posted: 1
ℹ️ Review info
⚙️ Run configuration
- Source mode: Auto
- Repository: NVIDIA/cccl
- Branch: main
- File: .coderabbit.yaml
- Commit: baaeed7
- Layers: Repository YAML
- Review profile: CHILL
- Plan: Enterprise
- Run ID:
70c55bbe-20b8-474f-8b32-64552b4eaaea
📒 Files selected for processing: 4
docs/libcudacxx/extended_api/bit.rstdocs/libcudacxx/extended_api/bit/bit_msb.rstlibcudacxx/include/cuda/__bit/bit_msb.hlibcudacxx/include/cuda/bit
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 8 remain after this review.
|
/ok to test 92ab995 |
This comment has been minimized.
This comment has been minimized.
|
/ok to test d95857a |
|
pre-commit.ci autofix |
|
/ok to test 63228c6 |
🥳 CI Workflow Results🟩 Finished in 2h 43m: Pass: 100%/228 | Total: 3d 02h | Max: 1h 47m | Hits: 99%/623210See results here. |
Addresses the msb request from #6108 (sibling of
bit_ffsin #9609).bit_msbreturns the zero-based index of the most significant set bit,floor(log2(value)), and -1 for zero. It forwards tocuda::std::__bit_log2so it reuses the existing bfind/clz lowering rather than adding new codegen.I went with the zero-based convention to match
__bit_log2, since a 1-based version would just bebit_width. Easy to flip to 1-based if you'd prefer. Tests and a docs page included.