Skip to content

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

Open
fbusato wants to merge 1 commit into
NVIDIA:mainfrom
fbusato:fix-ptx-shr-bfind
Open

Fix cuda::ptx shr and bfind long, long long handling#10000
fbusato wants to merge 1 commit into
NVIDIA:mainfrom
fbusato:fix-ptx-shr-bfind

Conversation

@fbusato

@fbusato fbusato commented Jul 17, 2026

Copy link
Copy Markdown
Contributor

Description

Fix cuda::ptx shr and bfind long, long long (+ unsigned) handling across Linux and Windows

Fixes #9992

@fbusato fbusato self-assigned this Jul 17, 2026
@fbusato
fbusato requested a review from a team as a code owner July 17, 2026 18:52
@fbusato fbusato added the libcu++ For all items related to libcu++ label Jul 17, 2026
@fbusato
fbusato requested a review from a team as a code owner July 17, 2026 18:52
@fbusato fbusato added this to CCCL Jul 17, 2026
@fbusato
fbusato requested a review from ericniebler July 17, 2026 18:52
@github-project-automation github-project-automation Bot moved this to Todo in CCCL Jul 17, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Jul 17, 2026
@coderabbitai

coderabbitai Bot commented Jul 17, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 600c17f8-19fc-430a-8fb0-64b4b052b2d4

📥 Commits

Reviewing files that changed from the base of the PR and between b757d02 and dba1bf7.

⛔ Files ignored due to path filters (6)
  • docs/libcudacxx/ptx/instructions/generated/bfind.rst is excluded by !**/generated/**
  • docs/libcudacxx/ptx/instructions/generated/shr.rst is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/bfind.h is excluded by !**/generated/**
  • libcudacxx/include/cuda/__ptx/instructions/generated/shr.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/bfind.h is excluded by !**/generated/**
  • libcudacxx/test/libcudacxx/cuda/ptx/generated/shr.h is excluded by !**/generated/**
📒 Files selected for processing (2)
  • libcudacxx/include/cuda/__ptx/instructions/bfind.h
  • libcudacxx/include/cuda/__ptx/instructions/shr.h

📝 Walkthrough

Summary by CodeRabbit

  • Bug Fixes
    • Improved compatibility and reliability for CUDA PTX bit-manipulation instruction support across integral and signed data types.

Walkthrough

Changes

PTX instruction type-trait dependencies

Layer / File(s) Summary
Add type-trait includes
libcudacxx/include/cuda/__ptx/instructions/bfind.h, libcudacxx/include/cuda/__ptx/instructions/shr.h
The PTX instruction headers now include CUDA standard-library traits for integral and signedness checks.

Assessment against linked issues

Objective Addressed Explanation
Ensure all signed integer types use arithmetic shift in cuda::ptx::shr [#9992] The required traits are included, but no overload or dispatch logic changes are shown to establish arithmetic shifting for long long or other signed integer types.

Out-of-scope changes

Code Change Explanation
Add is_integral, is_signed, and is_unsigned includes to bfind.h (libcudacxx/include/cuda/__ptx/instructions/bfind.h:27-29) The linked issue concerns cuda::ptx::shr; no stated objective covers bfind dependencies.

Possibly related PRs

  • NVIDIA/cccl#9993: Updates the PTX shr header with type-trait includes related to signed-shift dispatch.

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

@github-actions

Copy link
Copy Markdown
Contributor

😬 CI Workflow Results

🟥 Finished in 2h 42m: Pass: 99%/120 | Total: 4d 06h | Max: 2h 35m | Hits: 44%/1345151

See results here.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

libcu++ For all items related to libcu++

Projects

Status: In Review

Development

Successfully merging this pull request may close these issues.

[BUG]: Invalid instruction generated by cuda::ptx::shr for non-stdint type

1 participant