Using the paged vectors.

Start the runtime and the client

A vector needs a CTE core pool with at least one storage tier. Point CLIO_SERVER_CONF at a configuration like the one on the install page. CLIO_INIT(kClient, true) starts a runtime inside the process when none is running, which is how the adapter tests run; CLIO_INIT(kClient) attaches to one started with clio_run start. Then attach the CTE client.

#include <clio_runtime/clio_runtime.h>
#include <clio_cte/core/core_client.h>

if (!clio::run::CLIO_INIT(clio::run::RuntimeMode::kClient, true)) return 1;
if (!clio::cte::core::CLIO_CTE_CLIENT_INIT()) return 1;
// ... use vectors ...
clio::run::CLIO_RUNTIME_FINALIZE();

cpu_vector from host threads

The constructor names the tag, the page size in bytes, the logical length in elements and the cache geometry. The verbs block instead of suspending, and every verb is thread-safe, so any number of host threads can share one vector. A caller error throws std::runtime_error. This example follows the adapter's own test.

#include <clio_cte/core/core_client.h>
#include <clio_cte/cpu_vector/cpu_vector.h>
#include <clio_runtime/clio_runtime.h>
#include <vector>

namespace cv = clio::cte::cpu_vector;
using clio::run::u64;

int main() {
  if (!clio::run::CLIO_INIT(clio::run::RuntimeMode::kClient, true)) return 1;
  if (!clio::cte::core::CLIO_CTE_CLIENT_INIT()) return 1;

  const u64 page_bytes = 4096;
  const u64 epp = page_bytes / sizeof(u64);      // elements per page
  const u64 n = 6 * epp;
  cv::Vector<u64> v("my_vector", page_bytes, n, /*nsets=*/2, /*set_size=*/4);

  // Write: fetch (a page with no blob yet reads as zeros), hold for write, flush, unpin.
  v.Fetch(0, 0, n);
  for (u64 off = 0; off < n; off += epp) {
    auto h = v.HoldPage(off, epp, /*write=*/true);
    for (u64 i = off; i < off + h.run(); ++i) h[i] = i * 3;
  }
  v.Flush(0, 0, n);
  v.UnpinRange(0, n);

  // Read back through the cache.
  v.Fetch(0, 0, n);
  for (u64 off = 0; off < n; off += epp) {
    auto h = v.HoldPage(off, epp);
    for (u64 i = off; i < off + h.run(); ++i) {
      if (h[i] != i * 3) return 2;
    }
  }
  v.UnpinRange(0, n);

  // Or straight from the CTE, bypassing the cache.
  std::vector<u64> out(n);
  v.Download(out.data(), n);

  clio::run::CLIO_RUNTIME_FINALIZE();
  return 0;
}

Points to keep in mind:

• A hold indexes by absolute element offset. h[off] is valid for off in [h.begin_off(), h.begin_off() + h.run()). h.run() is how many elements of the requested range lie on this page.
• A read hold needs a fetch. Holding elements that were not fetched is an error. A write hold marks the elements valid instead.
• Flush what you wrote. There is no dirty tracking. A frame evicted before it was flushed loses its writes.
• Unpin. Six pages were fetched into a cache of eight frames above. Fetching more than the cache holds while everything is still pinned waits up to 30 seconds (kAllocWaitMs) for a frame, then throws with a message saying the cache is too small for what is held. Fetch, work, flush and unpin in page-sized steps when the data exceeds the cache.
• UpdateRange is Fetch of the whole pages under a range followed by HoldPage, for code that touches one range at a time.

Bulk movement and cache control

• Preload(src, n) writes n elements straight to the CTE, page by page, bypassing the cache. PreloadPages(pg_lo, pg_hi, fill) does the same with a callable that fills each page buffer.
• Download(dst, n) and DownloadPages(pg_lo, pg_hi, sink) read straight from the CTE. Flush first, or resident writes are not seen.
• Prefetch(pg_lo, pg_hi) makes a page range resident, best effort, and leaves it unpinned.
• FlushResidentToCte() writes every resident frame's valid elements back. ClearCache() drops every resident page and discards unflushed writes.

Lazy copies

Copy("my_vector_ckpt") flushes resident state, creates the tag my_vector_ckpt with the checkpoint module as its fault handler and my_vector as the source, and returns a new Vector on that tag with the same geometry. No bytes are copied: a page of the copy is materialized from the source the first time it is read. Because the source may keep changing, a snapshot meant as a checkpoint should be materialized while the source still holds the state you want: pass sync = true, or call MaterializeAll() on the copy, which returns the number of pages materialized. ScoreAllPages(score) rescores every page of a vector, for draining a checkpoint out of the fast tier.

auto ckpt = v.Copy("my_vector_ckpt", /*sync=*/true);   // std::unique_ptr<cv::Vector<u64>>
ckpt->ScoreAllPages(0.0f);                             // push the snapshot to the slowest tier

gpu_vector from a CUDA kernel

The host class Vector<T> owns the allocations and hands out a DeviceVector<T> that the kernel takes by value. The device side does the paging: a block fetches, holds, flushes and unpins from inside the kernel, and the runtime's CPU workers complete the transfers. The kernel is a device coroutine run by the yield driver, so the fetching verbs are awaited. Every DeviceVector call is block-collective: all threads of the block call it together, and no suspension may sit inside a thread-predicated branch.

This is the adapter's smoke test, cut down. One block fills a vector page by page, flushing as it goes, then a second kernel reads it back. gv::ResumeWhenComplete holds a parked block back until its transfer has landed.

#include <clio_runtime/clio_runtime.h>
#include <clio_runtime/gpu/gpu_ipc_manager.h>
#include <clio_runtime/gpu/yield_stack.h>
#include <clio_runtime/gpu/yieldable.h>
#include <clio_cte/core/core_client.h>
#include <clio_cte/gpu_vector/gpu_vector.h>

namespace gv = clio::cte::gpu_vector;
namespace gy = clio::run::gpu;
using clio::run::u32;
using clio::run::u64;

__device__ gy::YCoroMain FillCoro(gv::DeviceVector<u32> v, u64 n) {
  for (u64 i = 0; i < n;) {
    u64 run = 0;
    {
      co_await v.Fetch(0, v.PageLo(i), v.PageSpan(i, 1));
      auto h = co_await v.HoldPage(i, n - i, /*write=*/true);
      run = h.run();
      for (u64 k = threadIdx.x; k < run; k += blockDim.x) {
        h[i + k] = static_cast<u32>((i + k) * 7 + 1);
      }
    }
    co_await v.Flush(0, i, run);                  // the vector never writes back on its own
    v.UnpinRange(v.PageLo(i), v.PageSpan(i, 1));  // Fetch pinned; give the pin back
    i += run;
  }
}

__global__ void FillKernel(clio::run::IpcManagerGpuInfo info,
                           gv::DeviceVector<u32> v, u64 n,
                           gy::YieldableView<> yv, gy::YieldStackView ys) {
  CLIO_GPU_INIT(info, nullptr);
  v.Init(yv.Block());                 // the logical block id, not blockIdx.x
  gy::YieldTlsPublish(ys, yv.Y(), yv.Block());
  __syncthreads();
  CLIO_YCORO_RUN(FillCoro(v, n));
}

__device__ gy::YCoroMain CheckCoro(gv::DeviceVector<u32> v, u64 n,
                                   unsigned long long *bad) {
  for (u64 i = 0; i < n;) {
    co_await v.Fetch(0, v.PageLo(i), v.PageSpan(i, 1));
    auto h = co_await v.HoldPage(i, n - i);
    for (u64 k = threadIdx.x; k < h.run(); k += blockDim.x) {
      if (h[i + k] != static_cast<u32>((i + k) * 7 + 1)) atomicAdd(bad, 1ull);
    }
    const u64 run = h.run();
    v.UnpinRange(v.PageLo(i), v.PageSpan(i, 1));
    i += run;
  }
}

__global__ void CheckKernel(clio::run::IpcManagerGpuInfo info,
                            gv::DeviceVector<u32> v, u64 n,
                            unsigned long long *bad,
                            gy::YieldableView<> yv, gy::YieldStackView ys) {
  CLIO_GPU_INIT(info, nullptr);
  v.Init(yv.Block());
  gy::YieldTlsPublish(ys, yv.Y(), yv.Block());
  __syncthreads();
  CLIO_YCORO_RUN(CheckCoro(v, n, bad));
}

#if !CTP_IS_DEVICE_PASS
template <typename LaunchT>
static void RunYieldable(unsigned nblocks, LaunchT &&launch) {
  gy::Yieldable<> drv(nblocks, 32);
  gy::YieldStack stack(nblocks, 32, 8192);
  drv.RunToCompletion(
      [&](dim3 g, dim3 b, gy::YieldableView<> view) { launch(g, b, view, stack.View()); },
      [] {}, /*max_rounds=*/200000, gv::ResumeWhenComplete);
}

int main() {
  // CLIO_SERVER_CONF names a config whose compose has a clio_cte_core pool
  // with at least one storage tier, and a gpu: queue_depth large enough for
  // the blocks you launch (the default is 16).
  if (!clio::run::CLIO_INIT(clio::run::RuntimeMode::kClient, true)) return 1;
  if (!clio::cte::core::CLIO_CTE_CLIENT_INIT()) return 1;
  auto gpu_info = CLIO_CPU_IPC->GetGpuIpcManager()->GetGpuInfo(0);

  const u64 page_bytes = 4096, n = 4096;   // 4 pages of u32
  gv::Vector<u32> vec("gv_example", {0}, page_bytes,
                      /*nblocks=*/1, /*set_size=*/4, n);

  RunYieldable(1, [&](dim3 g, dim3 b, gy::YieldableView<> vw, gy::YieldStackView sv) {
    FillKernel<<<g, b, CLIO_YIELD_SMEM_BYTES>>>(gpu_info, vec.GetDevice(0), n, vw, sv);
  });
  ctp::GpuApi::Synchronize();

  auto *d_bad = ctp::GpuApi::Malloc<unsigned long long>(sizeof(unsigned long long));
  ctp::GpuApi::Memset(d_bad, 0, sizeof(unsigned long long));
  RunYieldable(1, [&](dim3 g, dim3 b, gy::YieldableView<> vw, gy::YieldStackView sv) {
    CheckKernel<<<g, b, CLIO_YIELD_SMEM_BYTES>>>(gpu_info, vec.GetDevice(0), n, d_bad, vw, sv);
  });
  ctp::GpuApi::Synchronize();

  unsigned long long bad = 0;
  ctp::GpuApi::Memcpy(&bad, d_bad, sizeof(bad));
  ctp::GpuApi::Free(d_bad);
  clio::run::CLIO_RUNTIME_FINALIZE();
  return bad == 0 ? 0 : 2;
}
#endif

Rules the test comments state and that apply to any kernel:

• Slice by yv.Block(), not blockIdx.x. The driver relaunches a compacted grid, so a block resumes on a different CUDA block than it started on. Init(block) with a block number beyond the nblocks the vector was built for traps.
• Every lane runs the loop. A hold synchronizes the block internally; returning early from some threads deadlocks the barrier.
• Flush as you go. Writing more pages than the cache holds without flushing fills the cache with pages nobody can evict.
• Several blocks writing one page flush only the element ranges they wrote, so a full-page flush does not clobber a neighbor's part.
• Overlap transfers with compute by splitting the verbs: BeginFetch then AwaitFetch, and BeginFlush then EndFlush, with work between them.

When a kernel traps, the host-side free function gv::FatalReport() returns the reason the device latched: page not resident, range never fetched, set full, fetch failed, vector used before Init, and so on.

Host-side control of a gpu_vector

The host never pages. It moves data in and out with Preload, PreloadPages, Download and DownloadPages, warms the cache with Prefetch(pg_lo, pg_hi, gpu_id), drops it with ClearCache(gpu_id), and writes resident frames back with FlushResidentToCte(). Copy, MaterializeAll and ScoreAllPages work as on the CPU. A copy made with Copy builds no device views; a kernel that must read it opens its own Vector on the copy's tag.

Tier placement hints

Pages are written at score 0.5, so with tiers scored HBM 1.0, host RAM 0.2 and storage 0.0 a page never lands in HBM on its own. A hint at 1.0 is what moves a page into such a tier. PrefetchNow(hints) takes a vector of PrefetchHint{page_, score_} and submits them as batched rescores; DrainPrefetch() waits for them. On gpu_vector, RegisterPrefetcher(p) installs a Prefetcher whose OnYield runs once per block per yield round and emits hints from what the kernel just touched; YieldObserver() gives the driver the callback that feeds it.

SYCL

A SYCL kernel body is a lambda rather than a __global__ function, so the launch code differs, but the DeviceVector verbs are the same. The SYCL tests in context-transfer-engine/adapter/gpu_vector/test/sycl are the reference for the launch sequence.

clio-coroc

Under the clio-coroc transpiler (-DCLIO_COROC) the same verbs are spelled CoFetch, CoBeginFetch, CoAwaitFetch, CoHoldPage (returning a PageRef<T>), CoFlush, CoBeginFlush and CoEndFlush, used with CO_AWAIT instead of co_await.

Where to go next

• API reference: every constructor argument, method and configuration key.
• Context Transfer Engine usage: reading pages as blobs from any client.
• Placement and tiering: how scores decide where a page lives.