cuda::associate_access_property

Defined in header <cuda/annotated_ptr>.

template <typename T, typename Property>
[[nodiscard]] __host__ __device__
T* associate_access_property(T* ptr, Property prop) noexcept;

Associate an cuda::access_property to the input 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.

Constraints

• Property is convertible to cuda::access_property.

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::global, cuda::access_property::persisting, cuda::access_property::normal, or cuda::access_property::streaming, then it must be valid to cast the generic pointer ptr to a pointer to the global memory address space.

Note: currently associate_access_property is ignored by nvcc and nvc++ on the host.

Example

#include <cuda/cooperative_groups.h>

__global__ void memcpy_kernel(const int* in, int* out) {
    __shared__ int smem[N];
    auto in1          = cuda::associate_access_property(in, cuda::access_property::streaming{});
    auto idx          = cooperative_groups::this_grid().thread_rank();
    smem[threadIdx.x] = in1[idx]; // streaming access
    // compute...
}