cuda::access_property

Defined in header <cuda/annotated_ptr>.

The class cuda::access_property provides an opaque encoding for L2 cache memory residence control and memory space properties. It is used in combination with cuda::annotated_ptr, cuda::associate_access_property and cuda::apply_access_property to request the application of properties to memory operations.

namespace cuda {

class access_property {
  public:
    // Static memory space property:
    struct shared {};
    struct global {};

    // Static global memory residence control property:
    struct normal {
        __host__ __device__ constexpr operator cudaAccessProperty() const noexcept;
    };
    struct persisting {
        __host__ __device__ constexpr operator cudaAccessProperty() const noexcept;
    };
    struct streaming {
        __host__ __device__ constexpr operator cudaAccessProperty() const noexcept;
    };

    access_property() noexcept = default;

    // Constructors from static global memory residence control properties:
    __host__ __device__ constexpr access_property(global)     noexcept;
    __host__ __device__ constexpr access_property(normal)     noexcept;
    __host__ __device__ constexpr access_property(streaming)  noexcept;
    __host__ __device__ constexpr access_property(persisting) noexcept;

    // Dynamic interleaved global memory residence control property constructors:
    __host__ __device__ constexpr access_property(normal,     float probability) noexcept;
    __host__ __device__ constexpr access_property(streaming,  float probability) noexcept;
    __host__ __device__ constexpr access_property(persisting, float probability) noexcept;
    __host__ __device__ constexpr access_property(normal,     float probability, streaming) noexcept;
    __host__ __device__ constexpr access_property(persisting, float probability, streaming) noexcept;

    // Dynamic range global memory residence control property constructors:
    __host__ __device__ access_property(void* ptr, size_t primary_bytes, size_t total_bytes, normal) noexcept;
    __host__ __device__ access_property(void* ptr, size_t primary_bytes, size_t total_bytes, streaming) noexcept;
    __host__ __device__ access_property(void* ptr, size_t primary_bytes, size_t total_bytes, persisting) noexcept;
    __host__ __device__ access_property(void* ptr, size_t primary_bytes, size_t total_bytes, global,     streaming) noexcept;
    __host__ __device__ access_property(void* ptr, size_t primary_bytes, size_t total_bytes, normal,     streaming) noexcept;
    __host__ __device__ access_property(void* ptr, size_t primary_bytes, size_t total_bytes, persisting, streaming) noexcept;
    __host__ __device__ access_property(void* ptr, size_t primary_bytes, size_t total_bytes, streaming,  streaming) noexcept;
};

} // namespace cuda

Kinds of Access Properties

Access properties are either static compile-time values or dynamic runtime values. The following properties of a memory access are provided:

Shared Memory property:

  • cuda::access_property::shared: memory access to the shared memory space.

Global Memory properties:

  • cuda::access_property::global: memory access to the global memory space without indicating an expected frequency of access to that memory, namely the access behavior is not modified.

  • cuda::access_property::normal: memory access to the global memory space expecting the memory to be accessed as frequent as other memory.

  • cuda::access_property::persisting: memory access to the global memory space expecting the memory to be accessed more frequently than other memory; this priority is suitable for data that should remain persistent in cache.

  • cuda::access_property::streaming: memory access to the global memory space expecting the memory to be accessed infrequently; this priority is suitable for streaming data.

Note: The difference between cuda::access_property::global and cuda::access_property::normal is subtle. The cuda::access_property::normal hints that the pointer points to the global address space and the memory will be accessed with “normal frequency”, while cuda::access_property::global only hints that the pointer points to the global address-space, it does not hint about how frequent the accesses will be.

Warning

The behavior of requesting the application of cuda::access_property to memory accesses, or their association with memory addresses, outside of the corresponding address space is undefined (note: even if that address is not used). The correctness of the input pointer and memory properties are verified in debug mode.

Global Memory Property Definition

The L2 residence control can be specified in two ways:

  • Interleaved: A memory address is accessed with a property with a given probability, while the remaining 1 - probability accesses are performed with a second one.

  • Range: The first primary_bytes of a memory address is accessed with one property and the remaining total_bytes - primary_bytes addresses with a second one.

Default constructor

access_property() noexcept = default;

Effects: as if access_property(global) (unchanged).

Static global memory residence control property constructors

__host__ __device__ constexpr access_property::access_property(global) noexcept;
__host__ __device__ constexpr access_property::access_property(normal) noexcept;
__host__ __device__ constexpr access_property::access_property(streaming) noexcept;
__host__ __device__ constexpr access_property::access_property(persisting) noexcept;

Effects: as if access_property(PROPERTY, 1.0) where PROPERTY is one of global (unchanged), normal, streaming, or persisting.

Dynamic interleaved global memory residence control property constructors

__host__ __device__ constexpr access_property::access_property(normal,     float probability) noexcept;
__host__ __device__ constexpr access_property::access_property(streaming,  float probability) noexcept;
__host__ __device__ constexpr access_property::access_property(persisting, float probability) noexcept;
__host__ __device__ constexpr access_property::access_property(normal,     float probability, streaming) noexcept;
__host__ __device__ constexpr access_property::access_property(persisting, float probability, streaming) noexcept;

Preconditions: 0 < probability <= 1.0.

Effects: constructs an interleaved access property that requests the first and third arguments - access properties - to be applied with probability and 1 - probability to memory accesses. The overloads without a third argument request applying global (unchanged) with 1 - probability.

Dynamic range global memory residence control property constructors

__host__ __device__ access_property::access_property(void* ptr, size_t leading_bytes, size_t total_bytes, normal) noexcept;
__host__ __device__ access_property::access_property(void* ptr, size_t leading_bytes, size_t total_bytes, streaming) noexcept;
__host__ __device__ access_property::access_property(void* ptr, size_t leading_bytes, size_t total_bytes, persisting) noexcept;
__host__ __device__ access_property::access_property(void* ptr, size_t leading_bytes, size_t total_bytes, normal,     streaming) noexcept;
__host__ __device__ access_property::access_property(void* ptr, size_t leading_bytes, size_t total_bytes, persisting, streaming) noexcept;

note: pointer arithmetic below performed char* ptr instead of void* ptr

Preconditions:

  • ptr is a generic pointer that is valid to cast to a pointer to the global memory address space.

  • 0 < leading_bytes <= total_bytes <= 4GB.

Postconditions: memory accesses requesting the application of this property must be in range [ptr, ptr + total_bytes).

Effects: the fourth and fifth arguments, access properties, are called primary and secondary properties. The overloads without a fifth argument use global as the secondary property. Constructs a range access property requesting the properties to be approximately applied to memory accesses as follows:

  • primary property to accesses in address-range: [ptr, ptr + leading_bytes)

  • secondary property to accesses in address-range: [ptr + leading_bytes, ptr + total_bytes)

Note: This property enables two main use cases:

  1. Unary range [ptr, ptr + total_bytes) with primary property by using leading_bytes == total_bytes.

  2. Binary range [ptr, ptr + leading_bytes) and [ptr + leading_bytes, ptr + total_bytes) with primary and secondary properties respectively.

Conversion operators

__host__ __device__ constexpr access_property::normal::operator     cudaAccessProperty() const noexcept;
__host__ __device__ constexpr access_property::streaming::operator  cudaAccessProperty() const noexcept;
__host__ __device__ constexpr access_property::persisting::operator cudaAccessProperty() const noexcept;

Allows constexpr cuda::access_property::normal{}, cuda::access_property::streaming{}, and cuda::access_property::persisting{} to be used in lieu of the corresponding CUDA Runtime cudaAccessProperty. See also L2 Policy for Persisting Accesses.

Example

#include <cuda/access_property>

__global__ void kernel(int* global_ptr, size_t num_bytes) {
    __shared__ int smem;
    cuda::access_property shared_prop{&smem, cuda::access_property::shared{}};
    cuda::access_property streaming_prop{global_ptr, sizeof(int), sizeof(int), cuda::access_property::streaming{}};
    cuda::access_property streaming_interleaved_prop{cuda::access_property::streaming{}, 1.0};
    cuda::access_property persisting_prop{global_ptr, num_bytes, num_bytes, cuda::access_property::persisting{});
}

__global__ void undefined_behavior(int* global_ptr) { // verified in debug mode
    __shared__ int smem;
    // Associating pointers with mismatching address spaces is undefined:
    cuda::access_property{global_ptr, cuda::access_property::shared{}}; // undefined behavior
    cuda::access_property{&smem, cuda::access_property::normal{}};      // undefined behavior
    cuda::access_property{&smem, cuda::access_property::streaming{}};   // undefined behavior
    cuda::access_property{&smem, cuda::access_property::persisting{}};  // undefined behavior

    // Using a zero probability or probability out-of-range (0, 1] is undefined:
    cuda::access_property{cuda::access_property::streaming{}, 0.0f};    // undefined behavior
    cuda::access_property{cuda::access_property::streaming{}, 2.0f};    // undefined behavior

    // Providing size values out-of-range is undefined:
    cuda::access_property{global_ptr, 0, 0, cuda::access_property::streaming{}, 0.0f}; // undefined behavior
    cuda::access_property{global_ptr, 8, 4, cuda::access_property::streaming{}, 2.0f}; // undefined behavior
}