Replace `thrust::raw_pointer_cast` with `cuda::std::to_address` by miscco · Pull Request #7439 · NVIDIA/cccl

miscco · 2026-01-30T10:55:13Z

This replaces thrust::raw_pointer_cast with the standard conforming cuda::std::to_address

We enable usage of that by specializing cuda::std::pointer_traits for the thrust smart pointers.

NOTE: I did not yet deprecate thrust::raw_pointer_cast because we should give users, e.g rapids a way to update and its not yet available. I plan on deprecating it in the 3.4 time frame

thrust/thrust/detail/pointer.h

bernhardmgruber

LGTM, just some questions.

thrust/thrust/detail/pointer.h

Jacobfaib · 2026-03-17T16:19:11Z

thrust/thrust/detail/pointer.h

+
+// Specialize pointer traits for everything that has the raw_pointer alias
+template <typename Pointer>
+struct ::cuda::std::pointer_traits<Pointer, ::cuda::std::void_t<typename Pointer::raw_pointer>>


Consider specializing std::pointer_traits as well here?

I am a bit on the fence, because the standard does not know about execution spaces, but with that I am fine

I think it would be useful for downstream users. In my experience it is common for downstream projects to use thrust fancy pointers, but not yet cuda::std:: utilities. I myself have (carefully) rolled my own specializations for std::pointer_traits for thrust fancy pointers, simply to have generic code that needs to introspect pointers Just Work(TM).

Jacobfaib · 2026-03-17T16:21:28Z

thrust/thrust/device_ptr.h

+
+// Specialize pointer traits for everything that has the raw_pointer alias
+template <class T>
+struct ::cuda::std::pointer_traits<THRUST_NS_QUALIFIER::device_ptr<T>>


Same here for std::pointer_traits?

I definitely do not want to do that because that is a device pointer and the std traits have no dealing with that

I'm not sure I follow, what would be the danger here?

device only sequences should work with cuda::std:: types and not std::types

bernhardmgruber · 2026-03-19T14:14:59Z

I have the feeling this PR massively grew in size after I approved. Why was this necessary and why can't we ship the changes added later in a separate PR?

miscco · 2026-03-19T15:00:23Z

Actually it is much smaller.

It introduces pointer_traits specializations for the contiguous thrust iterators.

It then replaces raw_pointer_cast with cuda::std::to_address because that is the standard interface to retrieve a pointer from a smart pointer

The issue is that we use raw_pointer_cast a lot

pauleonix

Everything in cub/ looks fine to me.

…hrust

miscco · 2026-03-20T10:28:14Z

Not to reviewers. all relevant product code changes are in f87d697

The rest is only search & replace

github-actions · 2026-03-20T13:48:10Z

🥳 CI Workflow Results

🟩 Finished in 3h 14m: Pass: 100%/398 | Total: 7d 06h | Max: 3h 03m | Hits: 90%/511025

See results here.

bernhardmgruber

If there are no uses of thrust::raw_pointer_cast remaining, can we deprecate it?

bernhardmgruber · 2026-03-20T16:27:37Z

thrust/testing/cuda/partition.cu

 #include <thrust/partition.h>

-#include "thrust/detail/raw_pointer_cast.h"
+#include "cuda/std/memory"


Suggested change

#include "cuda/std/memory"

#include <cuda/std/memory>

miscco requested review from a team as code owners January 30, 2026 10:55

miscco requested a review from wmaxey January 30, 2026 10:55

github-project-automation bot added this to CCCL Jan 30, 2026

miscco requested a review from davebayer January 30, 2026 10:55

github-project-automation bot moved this to Todo in CCCL Jan 30, 2026

cccl-authenticator-app bot moved this from Todo to In Review in CCCL Jan 30, 2026

miscco force-pushed the pointer_traits_address_of branch from c7fab2a to 0403ea8 Compare January 30, 2026 11:08

davebayer reviewed Jan 30, 2026

View reviewed changes

thrust/thrust/detail/pointer.h Show resolved Hide resolved

bernhardmgruber approved these changes Jan 30, 2026

View reviewed changes

thrust/thrust/detail/pointer.h Outdated Show resolved Hide resolved

thrust/thrust/detail/pointer.h Outdated Show resolved Hide resolved

This comment has been minimized.

Sign in to view

miscco force-pushed the pointer_traits_address_of branch 2 times, most recently from 138f219 to 1a5c52d Compare March 6, 2026 12:49

This comment has been minimized.

Sign in to view

Jacobfaib approved these changes Mar 17, 2026

View reviewed changes

miscco force-pushed the pointer_traits_address_of branch from 1a5c52d to 6597b5f Compare March 19, 2026 13:33

miscco requested review from a team as code owners March 19, 2026 13:33

miscco requested review from caugonnet, gonidelis and pauleonix March 19, 2026 13:33

miscco force-pushed the pointer_traits_address_of branch from 6597b5f to b9668c9 Compare March 19, 2026 13:46

This comment has been minimized.

Sign in to view

pauleonix approved these changes Mar 20, 2026

View reviewed changes

miscco force-pushed the pointer_traits_address_of branch 3 times, most recently from 5e21068 to da7abeb Compare March 20, 2026 09:13

miscco changed the title ~~Cleanup some of thrust pointer traits~~ Replace thrust::raw_pointer_cast with cuda::std::to_address Mar 20, 2026

miscco requested review from Jacobfaib, bernhardmgruber and pauleonix March 20, 2026 10:15

miscco added 5 commits March 20, 2026 11:27

Specialize pointer_traits for thrust smart pointers

f87d697

Replace thrust::raw_pointer_cast with cuda::std::to_address

bfeb38a

Replace `thrust::raw_pointer_cast with cuda::std::to_address in t…

e5b50bf

…hrust

No need to wrap a standard facility

04dda11

More occurences

0575f94

miscco force-pushed the pointer_traits_address_of branch from da7abeb to eeda680 Compare March 20, 2026 10:27

Drop all the includes too

fc0439a

miscco force-pushed the pointer_traits_address_of branch from eeda680 to fc0439a Compare March 20, 2026 10:30

oleksandr-pavlyk approved these changes Mar 20, 2026

View reviewed changes

Jacobfaib approved these changes Mar 20, 2026

View reviewed changes

bernhardmgruber reviewed Mar 20, 2026

View reviewed changes

Conversation

miscco commented Jan 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

bernhardmgruber left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

This comment has been minimized.

This comment has been minimized.

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!

bernhardmgruber commented Mar 19, 2026

Uh oh!

miscco commented Mar 19, 2026

Uh oh!

This comment has been minimized.

This comment has been minimized.

pauleonix left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

miscco commented Mar 20, 2026

Uh oh!

github-actions bot commented Mar 20, 2026

🥳 CI Workflow Results

🟩 Finished in 3h 14m: Pass: 100%/398 | Total: 7d 06h | Max: 3h 03m | Hits: 90%/511025

Uh oh!

bernhardmgruber left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

miscco commented Jan 30, 2026 •

edited

Loading

pauleonix left a comment •

edited

Loading