cuda::associate_access_property

template <class T, class Property>
__host__ __device__
T* associate_access_property(T* ptr, Property prop);

Preconditions:

if Property is cuda::access_property::shared then it must be valid to cast the generic pointer ptr to a pointer to the shared memory address space.
if Property is one of cuda::access_property::shared, cuda::access_property::shared, cuda::access_property::shared, or cuda::access_property::shared then it must be valid to cast the generic pointer ptr to a pointer to the global memory address space.
if Property is a cuda::access_property of “range” kind, then ptr must be in the valid range.

Mandates: Property is convertible to cuda::access_property.

Effects: no effects.

Hint: to associate an access property with the returned pointer, such that subsequent memory operations with the returned pointer or pointers derived from it may apply the access property.

The “association” is not part of the value representation of the pointer.
The compiler is allowed to drop the association; it does not have a functional consequence.
The association may hold through simple expressions, sequence of simple statements, or fully inlined function calls where the pointer value or C++ reference is provably unchanged; this includes offset pointers used for array access.
The association is not expected to hold through the ABI of an unknown function call, e.g., when the pointer is passed through a separately-compiled function interface, unless link-time optimizations are used.

Note: currently associate_access_property is ignored by nvcc and nvc++ on the host; but this might change any time.

Example

#include <cuda/cooperative_groups.h>
__global__ void memcpy(int const* in_, int* out) {
    int const* in = cuda::associate_access_property(in_, cuda::access_property::streaming{});
    auto idx = cooperative_groups::this_grid().thread_rank();

    __shared__ int shmem[N];
    shmem[threadIdx.x] = in[idx]; // streaming access

    // compute...
}