Unified shared memory, owned properly, and the evidence that it is unified. More...

#include <cstddef>
#include <cstdint>
#include <new>
#include <span>
#include <utility>
#include <sycl/sycl.hpp>

Classes
class	orrery::backend::UsmArray< T >
	An owning, movable, cache-line-aligned shared USM allocation. More...

Enumerations
enum class	orrery::backend::UsmKind : std::uint8_t { kUnknown , kHost , kDevice , kShared }
	Which kind of USM allocation a pointer is, as the runtime sees it. More...

Functions
constexpr const char *	orrery::backend::to_string (UsmKind kind) noexcept
UsmKind	orrery::backend::allocation_kind (const void *pointer, const sycl::queue &queue) noexcept
	What the runtime reports pointer to be, in queue's context.

Detailed Description

Unified shared memory, owned properly, and the evidence that it is unified.

On the target hardware the GPU is on the same package as the CPU and reads the same physical memory (section 2 of the implementation plan). SYCL's unified shared memory is how a program says so: an allocation made this way has one address that is valid in host code and in device code, and passing it to a kernel passes a pointer rather than scheduling a transfer.

Which of the three kinds, and why it matters

USM offers three allocation kinds and only one of them is right here.

Device allocations live in device memory and are not addressable from the host, so filling one means an explicit copy. On a discrete card that is the correct choice and the copy is the price of the fastest memory on the board. Here there is no separate memory for it to live in, so the copy would be from system memory to system memory: pure loss.

Host allocations live in host memory and are addressable from the device, which sounds right but is not. On a discrete system they are read across the bus on demand. The kind is meant for data touched once, and the runtime is entitled to treat it as uncached from the device's point of view.

Shared allocations are addressable from both and the runtime is free to place the pages wherever it likes, including nowhere in particular when host and device share a memory controller. That is the kind this backend uses, and on a unified part it is the kind for which no copy exists to elide.

Why the kind is queried rather than assumed

Phase 9's definition of done asks for the absence of a host-to-device copy to be demonstrated rather than asserted, and "I called the shared allocator" is an assertion. allocation_kind asks the runtime what a pointer actually is, through the same query any tool would use, and tests/backend/sycl_usm_test.cpp checks that a pointer the host wrote is reported as shared, is readable from a kernel, and comes back with the kernel's changes visible on the host without any copy having been requested. That is a test rather than a paragraph, which is the distinction the phase is asking for.

Enumeration Type Documentation

◆ UsmKind

enum class orrery::backend::UsmKind : std::uint8_t

strong

Which kind of USM allocation a pointer is, as the runtime sees it.

Mirrors sycl::usm::alloc in a type that does not require the caller to include a device runtime header, so that a test message or a report can name the kind without the file that prints it becoming a SYCL translation unit.

Enumerator
kUnknown	Not a USM pointer at all. What an ordinary new returns, and what a mistake looks like.
kHost	Host memory, device-addressable.
kDevice	Device memory, not host-addressable.
kShared	Addressable from both. What this backend allocates.

Function Documentation

◆ allocation_kind()

UsmKind orrery::backend::allocation_kind	(	const void *	pointer,
		const sycl::queue &	queue )

nodiscardnoexcept

What the runtime reports pointer to be, in queue's context.

The demonstration described above. A pointer that this backend allocated answers kShared; one that came from new answers kUnknown.

Classes

Enumerations