Skip to content

Provide cuda::ptr_alignment#7701

Open
fbusato wants to merge 4 commits intoNVIDIA:mainfrom
fbusato:pointer-alignment
Open

Provide cuda::ptr_alignment#7701
fbusato wants to merge 4 commits intoNVIDIA:mainfrom
fbusato:pointer-alignment

Conversation

@fbusato
Copy link
Contributor

@fbusato fbusato commented Feb 17, 2026

Description

Retrieving the alignment of a pointer is a common utility in CCCL, as well as in other CUDA libraries.; see for example cudax::hyperloglog

The PR adds cuda::ptr_alignment(ptr, max_alignment) (<cuda/memory>) to get the alignment of a pointer (the largest power of two dividing its address), optionally capped by a maximum alignment.

@fbusato fbusato self-assigned this Feb 17, 2026
@fbusato fbusato requested a review from a team as a code owner February 17, 2026 23:54
@fbusato fbusato added the libcu++ For all items related to libcu++ label Feb 17, 2026
@fbusato fbusato requested a review from a team as a code owner February 17, 2026 23:54
@fbusato fbusato added this to CCCL Feb 17, 2026
@fbusato fbusato requested review from alliepiper and wmaxey February 17, 2026 23:54
@github-project-automation github-project-automation bot moved this to Todo in CCCL Feb 17, 2026
@cccl-authenticator-app cccl-authenticator-app bot moved this from Todo to In Review in CCCL Feb 17, 2026
@github-actions

This comment has been minimized.

Comment on lines 12 to 13
[[nodiscard]] __host__ __device__ inline
size_t ptr_alignment(const void* ptr, size_t max_alignment = 0) noexcept;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm not sure if I understand the motivation of having this function. Why using is_aligned is not enough?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

suppose you have the following situation.

auto align = ptr_alignment(ptr, 8);
switch (align) {
   case 1:
   case 2:
   ....
}

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I see, but is there enough motivation to publicly expose the function?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

that's a good question. We could also use it only internally but, the situation described above is very common in CUDA. @bernhardmgruber, @miscco any thought?

@github-actions
Copy link
Contributor

😬 CI Workflow Results

🟥 Finished in 1h 20m: Pass: 98%/99 | Total: 1d 16h | Max: 1h 00m | Hits: 96%/251275

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.

2 participants

Comments