Link Search Menu Expand Document


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


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.


#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...