Memory expansion API reference.

Shared model

• Page p of a vector is the blob named p in decimal in the vector's tag, page_bytes long, written at score 0.5 (kVectorBlobScore).
• Ranges are (offset, count) pairs in elements. One Fetch or Flush takes up to kMaxFetchRanges (8) pairs.
• The first argument of Fetch and Flush is a generation. 0 is an ordinary transfer. Otherwise a flush stamps that generation on the extents it wrote, and a fetch waits until at least that generation has been published and refetches a resident frame older than it.
• A page with no blob yet reads as zeros on an ordinary fetch.
• PageOf(off), PageLo(off) and PageSpan(off, count) convert an element range to whole pages: Fetch(0, v.PageLo(off), v.PageSpan(off, count)) fetches every page the range touches.

cpu_vector

Header clio_cte/cpu_vector/cpu_vector.h, namespace clio::cte::cpu_vector, header-only. The verbs block, every verb is thread-safe, and a caller error throws std::runtime_error.

Constructor

template <typename T>
class Vector {
 public:
  Vector(const std::string &tag_name, u64 page_bytes, u64 num_elems,
         u32 nsets, u32 set_size,
         clio::run::PoolId storage_pool_id = clio::run::PoolId::GetNull(),
         int compress_lib = 0, int compress_preset = 1,
         u32 capacity_pages = 0);
  Vector(const Vector &) = delete;
};

• tag_name: the CTE tag, created if it does not exist.
• page_bytes: page size; must be non-zero and a multiple of sizeof(T).
• num_elems: logical length. The last page may be partial.
• nsets, set_size: cache geometry, both non-zero. The cache holds up to nsets * set_size pages.
• storage_pool_id: the pool puts and gets go to. Null means the CTE core pool (kCtePoolId, 512.0).
• compress_lib, compress_preset: the CTE compression context used for every page; 0 means none.
• capacity_pages: page-sized regions backing the frames. 0 means one per frame. Smaller than nsets * set_size makes the frames tags over a tighter memory budget.

Verbs

template <typename... Args> void Fetch(u64 generation, Args... off_count_pairs);
Held<T> HoldPage(u64 off, u64 count, bool write = false);
Held<T> UpdateRange(u64 off, u64 count, bool write = false, u64 generation = 0);
template <typename... Args> void Flush(u64 generation, Args... off_count_pairs);
void UnpinRange(u64 off, u64 count);
void UnpinRange2(u64 off0, u64 count0, u64 off1, u64 count1);

• Fetch makes every page of the ranges resident and pins it. It throws if a get fails or, after kAllocWaitMs (30000 ms) of waiting for a free frame, if every frame of a set is pinned or in flight.
• HoldPage returns a view on the part of [off, off+count) that lies on off's page. A read hold requires the elements to have been fetched; a write hold marks them valid. It never fetches.
• UpdateRange is Fetch of the whole pages under the range, then HoldPage. The pins stay until UnpinRange(PageLo(off), PageSpan(off, count)).
• Flush writes the named ranges back and waits for the puts. Only elements a resident frame holds are sent; a range whose page is not resident is dropped and counted in flush_skipped.
• UnpinRange releases the pins Fetch took on every page of the range. UnpinRange2 does two ranges.

Held<T>

template <typename T> class Held {
 public:
  T *ptr() const;                 // first element of the run
  u64 run() const;                // elements on this page
  u64 begin_off() const;          // absolute offset of ptr()[0]
  explicit operator bool() const; // false when empty
  T &operator[](u64 off) const;   // absolute element offset
  void Rescore(float rank) const; // eviction rank of the frame; lowest is evicted first
};

Move-only. A hold returned by HoldPage does not own the pin; UnpinRange releases it.

Host data movement and cache control

template <typename FillFn> void PreloadPages(u64 pg_lo, u64 pg_hi, FillFn fill);   // fill(pg, char *buf)
void Preload(const T *src, u64 n);
template <typename SinkFn> u64 DownloadPages(u64 pg_lo, u64 pg_hi, SinkFn sink);  // sink(pg, const char *bytes); returns pages read
u64 Download(T *dst, u64 n);                                                        // returns elements read
void Prefetch(u64 pg_lo, u64 pg_hi);        // best effort, leaves pages unpinned
void FlushResidentToCte();                  // every resident frame's valid elements
void ClearCache();                          // drops resident pages, discards unflushed writes

PreloadPages and Preload write straight to the CTE and throw if a put fails; resident copies are not updated, so call ClearCache first if those pages may be resident. DownloadPages and Download read straight from the CTE; flush first.

Copies

std::unique_ptr<Vector<T>> Copy(const std::string &new_name, bool sync = false);
u64 MaterializeAll(u32 inflight = 32);      // pages materialized; 0 on the first failure or when not a copy
void ScoreAllPages(float score);

Copy flushes resident state, creates the checkpoint pool if needed, creates new_name with the checkpoint module as its fault handler and this tag as the source, and returns a vector on the copy with the same geometry. It throws if the source tag name is too long for the fault parameters or any step fails. With sync true every page is materialized before it returns.

Tier hints and diagnostics

struct PrefetchHint { u64 page_; float score_; };
void PrefetchNow(const std::vector<PrefetchHint> &hints);   // non-blocking, deduplicated
void DrainPrefetch();                                       // wait for outstanding rescores
void ResetPrefetchHistory();                                // per run, not per phase
PrefetchStats ReadPrefetchStats() const;  void ResetPrefetchStats();

void EnableStats();  void ResetStats();  Stats ReadStats() const;  Stats ReadStats(int gpu_id) const;
std::vector<Page> ReadTable() const;  u32 Occupied() const;  u32 MaxPageCopies() const;

Stats counts faults, puts, evicts, get_errors, gen_ok, gen_stale, gen_busy, flush_skipped, put_errors and alloc_waits, and is populated only after EnableStats(). Page is a snapshot of one frame: page_num, data, last_access, generation, score, pins, flushing, fetching, valid_lo, valid_hi.

Geometry

size(), TagId(), PageBytes(), ElemsPerPage(), NumPages(), NumSets(), SetSize(), NumFrames(), Regions(), PageOf(off), PageLo(off), PageSpan(off, count), HomeSet(pn).

gpu_vector: host side

Header clio_cte/gpu_vector/gpu_vector.h, namespace clio::cte::gpu_vector, header-only. The host Vector<T> owns the device allocations and never pages itself.

Constructor

template <typename T>
class Vector {
 public:
  Vector(const std::string &tag_name, const std::vector<int> &gpu_ids,
         u64 page_bytes, u32 nblocks, u32 set_size, u64 num_elems,
         clio::run::PoolId storage_pool_id = clio::run::PoolId::GetNull(),
         int compress_lib = 0, int compress_preset = 1,
         u32 nsets = 0, u32 capacity_pages = 0);
  Vector(const Vector &) = delete;
  DeviceVector<T> GetDevice(int gpu_id) const;
  static clio::run::PoolId DefaultStoragePool();
};

• gpu_ids: the GPUs to build device views for. GetDevice(gpu_id) returns the value a kernel takes.
• nblocks: the number of blocks the kernel will launch. Each block owns its own fetch and flush task set, so a kernel that calls Init with a block number at or beyond this traps.
• set_size: frames per associative set. nsets defaults to nblocks. There is one cache for the whole grid and a page is stored once; its home set is a hash of the page number.
• capacity_pages: page-sized regions of device memory; 0 means nsets * set_size.
• storage_pool_id: null means DefaultStoragePool(), which reads CLIO_GV_STORAGE_POOL as major.minor and falls back to the CTE core pool.

Host methods

template <typename FillFn> void PreloadPages(u64 pg_lo, u64 pg_hi, FillFn fill);
void Preload(const T *src, u64 n);
template <typename SinkFn> u64 DownloadPages(u64 pg_lo, u64 pg_hi, SinkFn sink);
u64 Download(T *dst, u64 n);
void Prefetch(u64 pg_lo, u64 pg_hi, int gpu_id = 0, u32 tables = 0);   // tables is ignored
void ClearCache(int gpu_id = 0);
void FlushResidentToCte();                 // on every GPU this vector has views for

std::unique_ptr<Vector<T>> Copy(const std::string &new_name, bool sync = false);
u64 MaterializeAll(u32 inflight = 32);
void ScoreAllPages(float score);

void RegisterPrefetcher(std::shared_ptr<Prefetcher> p);
void SetPrefetchPolicy(const PrefetchPolicy &p);  const PrefetchPolicy &GetPrefetchPolicy() const;
bool HasPrefetchers() const;  std::string PrefetcherNames() const;
clio::run::gpu::YieldObserverFn YieldObserver();   // drv.SetYieldObserver(vec.YieldObserver())
void PrefetchRunBegin();
void PrefetchNow(const std::vector<PrefetchHint> &hints);
void DrainPrefetch();  void ResetPrefetchHistory();
PrefetchStats ReadPrefetchStats() const;  void ResetPrefetchStats();

void EnableStats();  void ResetStats();  Stats ReadStats(int gpu_id) const;
std::vector<Page> ReadTable(int gpu_id) const;  u32 Occupied(int gpu_id) const;  u32 MaxPageCopies(int gpu_id) const;
clio::cte::core::TagId TagId() const;  u64 PageBytes() const;  u64 ElemsPerPage() const;  u64 NumPages() const;

• A copy made with Copy builds no device views. A kernel that must read it opens its own Vector on the copy's tag.
• A Prefetcher (in prefetch.h) implements OnYield(const YieldEvent &, PrefetchSink &), called once per block per yield round, plus OnRunBegin(), OnRoundEnd(round, sink) and Name(). Several may be registered; for a repeated page the last writer wins. Hints are deduplicated against the score last sent, batched and submitted without blocking.
• The free function FatalReport() returns a string describing why a kernel trapped, or an empty string if it did not.

gpu_vector: DeviceVector<T>

Defined in device_vector.h. Every function is block-collective: all threads of the block call it together, and no suspension may sit inside a thread-predicated branch. The awaited verbs exist when the toolchain compiles device coroutines (CLIO_HAS_YCORO).

void Init(u32 block);                                             // once, with yv.Block()

__device__ YCoroTask Fetch(u64 generation, Args... off_count_pairs);       // BeginFetch + AwaitFetch
__device__ YCoroTask BeginFetch(u64 generation, Args... off_count_pairs);  // submit, then compute
__device__ YCoroTask AwaitFetch();                                         // then wait
__device__ YCoroTaskT<Held<T>> HoldPage(u64 off, u64 count, bool write = false);
__device__ YCoroTaskT<Held<T>> UpdateRange(u64 off, u64 count, u32 from, bool write = false, u64 generation = 0);
__device__ YCoroTask Flush(u64 generation, u64 off, u64 count, Rest... more_pairs);       // BeginFlush + EndFlush
__device__ YCoroTask BeginFlush(u64 generation, u64 off, u64 count, Rest... more_pairs);
__device__ YCoroTask EndFlush();
void UnpinRange(u64 off, u64 count);
void UnpinRange2(u64 off0, u64 count0, u64 off1, u64 count1);

• YCoroTask and YCoroTaskT are clio::run::gpu::YCoroTask and clio::run::gpu::YCoroTaskT; use them with co_await inside a YCoroMain coroutine run by CLIO_YCORO_RUN.
• BeginFetch returns once the fetch is submitted; pair it with AwaitFetch before holding. Only the named elements become valid, and a hold over elements the fetch did not name traps.
• HoldPage does not fault. A page that is still absent is a caller error and the kernel stops.
• UpdateRange's from argument is accepted and ignored; the verb is Fetch plus HoldPage.
• BeginFlush writes back only the named ranges, so several blocks can write different parts of one page. EndFlush waits for it. Prefer the split form in a loop so one page's put overlaps the next page's compute.
• The clio-coroc spellings CoFetch, CoBeginFetch, CoAwaitFetch, CoHoldPage (returning PageRef<T>), CoFlush, CoBeginFlush and CoEndFlush are used with CO_AWAIT.

Device geometry: size(), PageBytes(), ElemsPerPage(), NumPages(), NumSets(), SetSize(), Regions(), PageOf, PageLo, PageSpan. The device Held<T> has the same ptr(), run(), absolute operator[] and Rescore(rank) as the host one.

Configuration

Environment variables

• CLIO_GV_STORAGE_POOL: the pool a gpu_vector uses when its constructor is given none, as major.minor. For example 563.0 sends pages through the CTE cache module for node-local caching of remote-owned pages. cpu_vector does not read it.
• CLIO_SERVER_CONF: the runtime configuration file, which must compose a clio_cte_core pool with at least one storage tier.
• CLIO_WITH_RUNTIME: 0 attaches to a running runtime; 1 starts one in the process.

YAML keys

• compose[].storage on the clio_cte_core entry: the tiers pages can reach. Pages are written at score 0.5. Under the max_bw placement engine a tier whose score is above 0.5 is not chosen for them on its own; a PrefetchNow hint, a registered Prefetcher or ScoreAllPages moves pages into it. See placement and tiering and the CTE YAML reference.
• gpu.queue_depth: tasks per GPU queue, default 16. Set it above the number of blocks that will have fetches or flushes in flight at once. The same section also takes threads_per_block.
• Compression is not a YAML key for the vectors: pass compress_lib and compress_preset to the constructor, and every page is put and got with that context.

Where to go next

• Usage: the verbs in action on the host and in a kernel.
• Install: build options and targets.
• Context Transfer Engine API: the client the vectors are built on.