Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Compile-time Error

Component

Thrust

Describe the bug

The device_reference<T> type acts as a transparent proxy for an object that resides in device memory.

It has an assignment operator declared as follows:

_CCCL_HOST_DEVICE device_reference& operator=(const value_type& x);

I believe the signature of this operator should be relaxed as marked as const. A constant memory location can still be represented as a device_reference<const T>.

The motivation for this comes from the indirectly_writable concept. This concept requires that the iterator's "reference" type be assignable to, even when const-qualified. See the cppreference page for more details and rationale on this requirement.

This in turn prevents std::output_iterator from being satisfied, which is needed by many STL ranges algorithms.

How to Reproduce

static_assert(std::indirectly_writable<uint8_t*, uint8_t>); // This succeeds
static_assert(std::indirectly_writable<thrust::device_ptr<uint8_t>, uint8_t>); // This fails

https://godbolt.org/z/bMd9Yaa5h

Expected behavior

std::indirectly_writable<thrust::device_ptr<uint8_t>, uint8_t> should be satisfied.

Reproduction link

https://godbolt.org/z/bMd9Yaa5h

Operating System

No response

nvidia-smi output

No response

NVCC version

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2024 NVIDIA Corporation
Built on Thu_Mar_28_02:18:24_PDT_2024
Cuda compilation tools, release 12.4, V12.4.131
Build cuda_12.4.r12.4/compiler.34097967_0

Same results with 12.5.1 on Godbolt.