Skip to main content

GPU Client Kernels

This guide describes how CUDA, ROCm, and SYCL kernels submit tasks to the Clio CPU runtime. The GPU subsystem is producer-only: kernels never execute Module code themselves. They populate pre-allocated task structures in registered device-memory backends, push them onto a per-device gpu2cpu_queue, and wait for the CPU runtime to write back results.

If you are migrating from an older version of this guide that documented a GPU work orchestrator, persistent GPU kernels, or CPU-to-GPU dispatch, see the Removed APIs section at the end.

Overview​

              GPU device                            CPU runtime
   ┌────────────────────────────┐         ┌──────────────────────────────┐
   │ kernel (producer)          │         │ Worker::ProcessNewTaskGpu     │
   │   1. mutate POD task        │         │   1. pop gpu::Future<Task>    │
   │   2. CLIO_IPC->Send(task_fp)│ ──push──▶  2. resolve Task via SendIn   │
   │   3. future.Wait()          │         │   3. RouteTask -> chimod      │
   │                            │   ◀──   │   4. SendOut writes back POD  │
   └────────────────────────────┘  signal │      output + sets            │
                                          │      fut_.is_complete_        │
                                          └──────────────────────────────┘
                                  per-device gpu2cpu_queue (MPSC ring)

The only GPU-related routing mode is clio::run::RoutingMode::ToLocalCpu, constructed with clio::run::PoolQuery::ToLocalCpu(parallelism). Every other routing decision (Local, DirectId, DirectHash, Broadcast, Physical, Dynamic, Range) routes through CPU workers exactly as it does for non-GPU clients.

Architecture​

Server-side state​

clio::run::CLIO_INIT(clio::run::RuntimeMode::kServer) (or kRuntime, the alias) brings up the CPU runtime, which enumerates GPU devices and calls clio::run::gpu::IpcManager::ServerInitGpuQueues(queue_depth). For each physical GPU the runtime allocates:

  • A pinned-host backend large enough for the per-device GpuTaskQueue (declared as using GpuTaskQueue = ctp::ipc::multi_mpsc_ring_buffer<gpu::Future<Task>, CLIO_QUEUE_ALLOC_T>).
  • A PerGpuDeviceState struct holding the queue plus a client_backends map (AllocatorId -> ClientBackend).

A dedicated CPU worker polls each lane via Worker::ProcessNewTaskGpu. When it pops a gpu::Future<Task>, IpcGpu2Cpu::RecvIn:

  1. Resolves the task's ShmPtr into a host-readable address (direct dereference for pinned host / managed UVM, D2H ctp::GpuApi::Memcpy for pure device memory).
  2. Copies the POD task bytes into a per-thread scratch slot if it had to D2H-copy them, then calls Container::FixupAfterCopy so SSO/SVO pointers inside the task POD are valid host-side.
  3. Calls IpcManager::RouteTask so the task runs on the normal CPU worker pipeline.
  4. Once the chimod handler returns, IpcGpu2Cpu::SendOut writes the mutated POD bytes back to the original device address (H2D ctp::GpuApi::Memcpy when needed) and sets fut_.is_complete_ on the device-side task. The kernel's poll loop sees the flag and future.Wait() returns.

Client-process limitations​

Pure client processes (clio::run::CLIO_INIT(clio::run::RuntimeMode::kClient)) do not attach to GPU queues. There is no ClientInitGpuQueues on clio::run::gpu::IpcManager; the only entry point is ServerInitGpuQueues. To submit tasks from a GPU kernel, the producing process must be the runtime process (the same process that called CLIO_INIT(kServer)).

Initialization​

Host side​

#include <clio_runtime/clio_runtime.h>
#include <clio_runtime/ipc_manager.h>
#include <cstdio>

using namespace clio::run;

int example_main() {
  if (!CLIO_INIT(RuntimeMode::kServer)) {
    return 1;
  }

  auto *ipc = CLIO_CPU_IPC;
  if (ipc->GetGpuQueueCount() == 0) {
    std::fprintf(stderr, "no GPU queues initialized\n");
    return 1;
  }

  // ... allocate backends, launch kernels, run tasks ...

  CLIO_RUNTIME_FINALIZE();
  return 0;
}

CLIO_INIT is declared in clio_runtime/clio_runtime.h as an inline wrapper around ClioInitImpl:

#include <clio_runtime/clio_runtime.h>

using namespace clio::run;

void example() {
  // enum class RuntimeMode { kClient, kServer, kRuntime = kServer };
  // bool CLIO_INIT(RuntimeMode mode,
  //                bool default_with_runtime = false,
  //                bool is_restart = false);
  bool ok = CLIO_INIT(RuntimeMode::kServer);
  (void)ok;
}

Environment variable CLIO_WITH_RUNTIME overrides default_with_runtime. There is no co-located mode beyond kServer.

Kernel side​

Every kernel that submits tasks must expand CLIO_GPU_INIT(gpu_info, ipc_ptr) exactly once at entry. The macro has two backend-specific definitions in clio_runtime/gpu/gpu_ipc_manager.h:

  • CUDA / ROCm: places a clio::run::gpu::IpcManager instance in __shared__ memory (via GetBlockIpcManager()), calls ClientInitGpu(gpu_info) on lane 0, and synchronizes the block. ipc_ptr is ignored in this backend (pass nullptr).
  • SYCL: takes a USM-allocated clio::run::gpu::IpcManager * as ipc_ptr, calls ClientInitGpu(gpu_info) on it, and stores the pointer in a kernel-scope variable named g_ipc_manager_ptr so CLIO_IPC resolves correctly inside the kernel body.

In both expansions the macro introduces two kernel-scope names:

  • g_ipc_manager_ptr — a clio::run::gpu::IpcManager *.
  • g_ipc_manager — a reference to *g_ipc_manager_ptr.

You can use either directly, or use the CLIO_IPC macro (which resolves to one of them depending on the build). The stress test uses g_ipc_manager_ptr explicitly to dodge an NVCC two-pass name-resolution quirk; see context-runtime/test/unit/gpu/test_gpu_kernel_stress_gpu.cc for the reasoning.

Memory model​

Backend kinds​

clio::run::gpu::IpcManager::MemKind selects how the runtime allocates and resolves a registered backend:

MemKindUnderlying allocator (CUDA / ROCm / SYCL)CPU visibility
kPinnedHostcudaHostAlloc / hipHostMalloc / sycl::malloc_hostDirect (UVA)
kManagedUvmcudaMallocManaged / hipMallocManaged / sycl::malloc_sharedDirect (page-fault)
kDeviceMemcudaMalloc / hipMalloc / sycl::malloc_deviceD2H ctp::GpuApi::Memcpy of POD bytes

kPinnedHost is the lowest-latency option for small tasks. kDeviceMem lets you keep blob payloads on the GPU; the worker handles the round-trip POD copy automatically.

Allocating a backend​

#include <clio_runtime/clio_runtime.h>
#include <clio_runtime/ipc_manager.h>
#include <clio_runtime/gpu/gpu_ipc_manager.h>

using namespace clio::run;

void example() {
#if CTP_ENABLE_GPU
  auto *ipc = CLIO_CPU_IPC;
  char *base = nullptr;
  ctp::ipc::AllocatorId alloc_id = ipc->AllocateAndRegisterGpuBackend(
      /*gpu_id=*/0,
      gpu::IpcManager::MemKind::kPinnedHost,
      /*bytes=*/4096, &base);
  if (alloc_id.IsNull()) { /* allocation failed */ }
  // ... use base ...
  ipc->FreeGpuBackend(/*gpu_id=*/0, alloc_id);
#endif
}

The host-side helpers are declared in clio_runtime/ipc_manager.h (gated behind CTP_ENABLE_CUDA || CTP_ENABLE_ROCM || CTP_ENABLE_SYCL):

ctp::ipc::AllocatorId AllocateAndRegisterGpuBackend(
    u32 gpu_id, gpu::IpcManager::MemKind kind, size_t bytes,
    char **out_base);
void FreeGpuBackend(u32 gpu_id, const ctp::ipc::AllocatorId &alloc_id);

For kPinnedHost and kManagedUvm, *out_base is a CPU-readable pointer; the same address is also dereferenceable by the GPU. For kDeviceMem, *out_base is the raw device pointer; the host cannot dereference it.

Slot layout: a self-contained Task​

The Task is now self-contained. Its embedded FutureInfo (task->fut_) carries both the completion flag (is_complete_) and the POD size (task_size_), so there is no separate co-located FutureShm — each slot inside a backend is just sizeof(TaskT) bytes:

+---------------- slot i (one task) ----------------+
| TaskT POD (e.g. GpuSubmitTask, PutBlobTask)       |
|   ... IN/OUT fields ...                            |
|   FutureInfo fut_ { is_complete_, task_size_ }     |
+---------------------------------------------------+
   ^
   task_addr = base + i * sizeof(TaskT)

clio::run::FutureInfo lives inside every Task (see clio_runtime/task.h):

#include <clio_runtime/task.h>

using namespace clio::run;

void example() {
  Task *t = nullptr;
  // Every Task embeds a self-contained completion record (FutureInfo fut_):
  //   t->fut_.task_size_   — sizeof(concrete TaskT), stamped before Send
  //   t->fut_.is_complete_ — atomic flag the CPU worker sets on completion
  // Accessors: SetTaskSize / GetTaskSize / SetComplete / IsComplete.
  if (t) {
    t->SetTaskSize(0u);
    (void)t->IsComplete();
  }
}

The host pre-constructs each task in place and stamps fut_.task_size_. The kernel later flips POD input fields and submits.

#include <new>
#include <clio_runtime/types.h>
#include <clio_runtime/task.h>
#include <clio_runtime/pool_query.h>
#include <clio_runtime/MOD_NAME/MOD_NAME_tasks.h>

using namespace clio::run;

void example(char *base, PoolId pool_id, u32 i) {
  using TaskT = MOD_NAME::GpuSubmitTask;
  char *task_addr = base + i * sizeof(TaskT);

  auto *task = new (task_addr) TaskT(
      CreateTaskId(), pool_id, PoolQuery::ToLocalCpu(),
      /*gpu_id=*/0u, /*test_value=*/i);
  task->fut_.task_size_ = static_cast<u32>(sizeof(TaskT));   // critical
  (void)task;
}

Failing to set fut_.task_size_ is a common bug: the CPU worker reads it to know how many bytes to D2H-copy and how many to write back. If the kernel submits a task whose fut_.task_size_ is zero, the worker drops it. The host-side PlaceTaskSlots helper in the stress test stamps fut_.task_size_ for every slot before launch.

Handing slots to the kernel​

The kernel needs a ctp::ipc::FullPtr<TaskT> for each slot it intends to submit. Build them on the host and stage them where the kernel can read them (a pinned-host array works for CUDA/ROCm; SYCL kernels capture USM pointers directly):

#include <clio_runtime/types.h>
#include <clio_runtime/task.h>
#include <clio_runtime/gpu/future.h>

using namespace clio::run;

void example(char *task_ptr, size_t task_off) {
  using TaskT = Task;
  ctp::ipc::AllocatorId alloc_id;                 // from AllocateAndRegisterGpuBackend

  ctp::ipc::FullPtr<TaskT> fp;
  fp.shm_.alloc_id_ = alloc_id;                   // backend handle
  fp.shm_.off_ = task_off;                        // byte offset in backend
  fp.ptr_ = reinterpret_cast<TaskT *>(task_ptr);  // host-side raw pointer
  (void)fp;
}

For kDeviceMem backends the kernel still uses FullPtr with alloc_id set to null and off_ carrying the raw device address; the CPU worker detects this via ctp::IsDevicePointer and uses ctp::GpuApi::Memcpy to read the POD.

Writing a GPU-compatible task​

A task struct that can be created from the host and submitted by a kernel must satisfy:

  1. POD layout — no std::string, std::vector, raw pointers into host memory, etc. SSO types (with CTP_CROSS_FUN ctors) are acceptable, but only if the chimod registers a FixupAfterCopy for that task so the worker can rebase the SSO data_ pointer after the D2H copy.
  2. All constructors annotated CTP_CROSS_FUN (__host__ __device__).
  3. SerializeIn and SerializeOut annotated CTP_CROSS_FUN.
  4. A unique method id assigned to method_ in the constructor.

Real example from clio::run::MOD_NAME::GpuSubmitTask (see context-runtime/modules/MOD_NAME/include/clio_runtime/MOD_NAME/MOD_NAME_tasks.h):

#include <clio_runtime/types.h>
#include <clio_runtime/task.h>
#include <clio_runtime/pool_query.h>

using namespace clio::run;

struct GpuSubmitTask : public Task {
  enum Method { kGpuSubmit = 25 };

  IN u32 gpu_id_;
  IN u32 test_value_;
  INOUT u32 result_value_;
  OUT u32 counter_value_;

  CTP_CROSS_FUN GpuSubmitTask()
      : Task(), gpu_id_(0), test_value_(0), result_value_(0),
        counter_value_(0) {}

  CTP_CROSS_FUN explicit GpuSubmitTask(
      const TaskId &task_node,
      const PoolId &pool_id,
      const PoolQuery &pool_query,
      u32 gpu_id,
      u32 test_value)
      : Task(task_node, pool_id, pool_query, 25),
        gpu_id_(gpu_id), test_value_(test_value),
        result_value_(0), counter_value_(0) {
    task_id_ = task_node;
    pool_id_ = pool_id;
    method_ = Method::kGpuSubmit;
    task_flags_.Clear();
    pool_query_ = pool_query;
  }

  template <typename Archive>
  CTP_CROSS_FUN void SerializeIn(Archive &ar) {
    Task::SerializeIn(ar);
    ar(gpu_id_, test_value_, result_value_);
  }

  template <typename Archive>
  CTP_CROSS_FUN void SerializeOut(Archive &ar) {
    Task::SerializeOut(ar);
    ar(result_value_, counter_value_);
  }

  // ... Copy / AggregateOut ...
};

Writing the CPU-side handler​

The runtime side is plain CPU code: implement Runtime::GpuSubmit (or whatever name your method takes) just like any other CPU handler. The Module dispatcher does not distinguish GPU-produced tasks from CPU-produced ones. The handler takes a clio::run::shared_ptr<TaskT> and returns a clio::run::TaskResume:

#include <clio_runtime/clio_runtime.h>
#include <clio_runtime/task.h>
#include <clio_runtime/MOD_NAME/MOD_NAME_tasks.h>

using namespace clio::run;

// Same shape the chimod generator emits for MOD_NAME::Runtime::GpuSubmit.
TaskResume ExampleGpuSubmit(shared_ptr<MOD_NAME::GpuSubmitTask> &task) {
  CLIO_TASK_BODY_BEGIN
  // Formula: result = test_value * 2 + gpu_id
  task->result_value_ = (task->test_value_ * 2) + task->gpu_id_;
  CLIO_CO_RETURN;
  CLIO_TASK_BODY_END
}

After this returns, IpcGpu2Cpu::SendOut writes the mutated task POD back to the device-side slot and sets fut_.is_complete_.

Routing​

Inside a kernel, the only meaningful routing mode is clio::run::PoolQuery::ToLocalCpu(parallelism). From clio_runtime/pool_query.h:

#include <clio_runtime/pool_query.h>

using namespace clio::run;

void example() {
  // static CTP_CROSS_FUN PoolQuery ToLocalCpu(u32 parallelism = 32);
  PoolQuery query = PoolQuery::ToLocalCpu(/*parallelism=*/32);
  (void)query;
}

parallelism is informational — it lets the scheduler hint at how many lanes the producer expects to keep busy.

When the task arrives on the CPU side, IpcManager::RouteTask resolves it like any other task: it consults the pool's container set, picks a container, and either runs the task inline on the current worker or enqueues it on the destination worker's lane.

Removed APIs​

The following names appear in older code and documentation but have been removed. They will not compile against the current headers; if you have code that uses them, the migration is to recast the call as a producer-only submission via ToLocalCpu and have the chimod do the work on the CPU side.

Removed nameNotes
RoutingMode::ToLocalGpu / PoolQuery::ToLocalGpu(gpu_id)CPU→GPU dispatch removed
RoutingMode::LocalGpuBcast / PoolQuery::LocalGpuBcast()CPU→GPU broadcast removed
IpcManager::RouteToGpuSee deletion note in context-runtime/src/ipc_manager.cc
LaunchGpuOrchestrator, persistent GPU kernelGPU orchestrator removed
gpu::FutureShmFolded into Task::fut_ (clio::run::FutureInfo)
ClientInitGpuQueues on clio::run::gpu::IpcManagerPure-client processes have no GPU surface
CHIMAERA_INIT / ChimaeraModeRenamed to CLIO_INIT / RuntimeMode
CHI_CLIENT_GPU_INIT / CHIMAERA_GPU_INIT macrosReplaced by CLIO_GPU_INIT(gpu_info, ipc_ptr)
CHIMAERA_GPU_ORCHESTRATOR_INIT macroOrchestrator no longer exists
ChimaeraMode::kColocatedUse RuntimeMode::kServer (the runtime process is the producer)
IpcGpu2Cpu::ClientSend / RuntimeSendRenamed to SendIn / SendOut

clio::run::WorkOrchestrator still exists as a CPU thread-pool singleton — that is unrelated to the removed GPU orchestrator. Do not confuse the two.

CMake integration​

The runtime build has three independent GPU options, declared in the top-level CMakeLists.txt:

option(CLIO_CORE_ENABLE_CUDA "Enable CUDA support" OFF)
option(CLIO_CORE_ENABLE_ROCM "Enable ROCm support" OFF)
option(CLIO_CORE_ENABLE_SYCL "Enable Intel GPU support via SYCL/oneAPI (icpx -fsycl)" OFF)

When any of these is ON, the corresponding CLIO_CTP_ENABLE_* is forced on as well, which defines the CTP_ENABLE_CUDA, CTP_ENABLE_ROCM, or CTP_ENABLE_SYCL preprocessor macros consumed by the GPU headers. Any one of them also implies CTP_ENABLE_GPU.

Configure the project:

# CUDA
cmake -S . -B build -DCLIO_CORE_ENABLE_CUDA=ON

# ROCm
cmake -S . -B build -DCLIO_CORE_ENABLE_ROCM=ON

# SYCL (Intel oneAPI)
cmake -S . -B build -DCLIO_CORE_ENABLE_SYCL=ON \
    -DCMAKE_CXX_COMPILER=icpx

The test tree uses three helper functions (add_cuda_executable, add_rocm_gpu_executable, add_sycl_executable) to build the same .cc files with the right compiler. See context-runtime/test/unit/CMakeLists.txt for usage examples.

For your own kernels you can either:

  • Put the kernel in a *_gpu.cc source file and compile it with NVCC / HIPCC, or
  • Use a SYCL .cc file that includes <sycl/sycl.hpp> and is built with icpx -fsycl.

Either way, link against the runtime library (libclio_runtime) and any chimod client libraries you need. The headers you include from a kernel file are the same ones host code uses; the difference is purely the compiler driver.

Worked example​

This walks through the canonical pattern used by test_gpu_kernel_stress_*. It launches N blocks (CUDA / ROCm) or N single_task submissions (SYCL), each of which mutates one task POD and submits it.

1. Bring up the runtime and pool​

#include <string>
#include <clio_runtime/clio_runtime.h>
#include <clio_runtime/ipc_manager.h>
#include <clio_runtime/pool_query.h>
#include <clio_runtime/MOD_NAME/MOD_NAME_tasks.h>
#include <clio_runtime/MOD_NAME/MOD_NAME_client.h>

using namespace clio::run;

void example() {
  constexpr u32 kNumTasks = 64;
  // One slot per task; the Task carries its own completion record (fut_),
  // so there is no separate FutureShm.
  constexpr size_t kSlotBytes = sizeof(MOD_NAME::GpuSubmitTask);
  constexpr size_t kBackendBytes = kNumTasks * kSlotBytes + 256;

  PoolId pool_id(20002, 1);

  if (!CLIO_INIT(RuntimeMode::kServer)) return;

  auto *ipc = CLIO_CPU_IPC;
  if (ipc->GetGpuQueueCount() < 1u) return;

  MOD_NAME::Client client(pool_id);
  using CreateTask = MOD_NAME::CreateTask;
  using CreateParams = MOD_NAME::CreateParams;
  auto task = ipc->NewTask<CreateTask>(
      CreateTaskId(), kAdminPoolId,
      PoolQuery::Dynamic(),
      CreateParams::chimod_lib_name,
      std::string("gpu_demo_pool"),
      pool_id, &client);
  ipc->Send(task).Wait();
  (void)kBackendBytes;
}

2. Allocate a backend and place tasks​

#include <new>
#include <vector>
#include <clio_runtime/clio_runtime.h>
#include <clio_runtime/ipc_manager.h>
#include <clio_runtime/pool_query.h>
#include <clio_runtime/gpu/gpu_ipc_manager.h>
#include <clio_runtime/MOD_NAME/MOD_NAME_tasks.h>

using namespace clio::run;

void example() {
#if CTP_ENABLE_GPU
  auto *ipc = CLIO_CPU_IPC;
  const u32 gpu_id = 0;
  using TaskT = MOD_NAME::GpuSubmitTask;
  constexpr u32 kNumTasks = 64;
  constexpr size_t kSlotBytes = sizeof(TaskT);

  char *base = nullptr;
  auto alloc_id = ipc->AllocateAndRegisterGpuBackend(
      gpu_id, gpu::IpcManager::MemKind::kPinnedHost,
      kNumTasks * kSlotBytes + 256, &base);
  if (alloc_id.IsNull()) return;

  std::vector<ctp::ipc::FullPtr<TaskT>> handles;
  handles.reserve(kNumTasks);
  for (u32 i = 0; i < kNumTasks; ++i) {
    size_t off = static_cast<size_t>(i) * kSlotBytes;
    char *task_addr = base + off;
    auto *t = new (task_addr) TaskT(
        CreateTaskId(), PoolId(20002, 1), PoolQuery::ToLocalCpu(),
        gpu_id, /*test_value=*/i);
    t->fut_.task_size_ = static_cast<u32>(sizeof(TaskT));

    ctp::ipc::FullPtr<TaskT> fp;
    fp.shm_.alloc_id_ = alloc_id;
    fp.shm_.off_ = off;
    fp.ptr_ = t;
    handles.push_back(fp);
  }
#endif
}

3. Launch (CUDA / ROCm)​

#include <vector>
#include <clio_runtime/clio_runtime.h>
#include <clio_runtime/ipc_manager.h>
#include <clio_runtime/gpu/gpu_ipc_manager.h>
#include <clio_runtime/MOD_NAME/MOD_NAME_tasks.h>
#include <clio_ctp/util/gpu_api.h>

using namespace clio::run;

#if CTP_ENABLE_CUDA || CTP_ENABLE_ROCM
using TaskT = MOD_NAME::GpuSubmitTask;

__global__ void DemoKernel(IpcManagerGpuInfo info,
                           ctp::ipc::FullPtr<TaskT> *handles,
                           u32 num) {
  CLIO_GPU_INIT(info, /*ipc_ptr=*/nullptr);
  if (threadIdx.x != 0) return;
  u32 slot = blockIdx.x;
  if (slot >= num) return;
  auto fp = handles[slot];
  auto fut = g_ipc_manager_ptr->Send(fp);
  fut.Wait();
  (void)g_ipc_manager;
}

void launch(std::vector<ctp::ipc::FullPtr<TaskT>> &handles, u32 gpu_id) {
  auto *ipc = CLIO_CPU_IPC;
  const u32 kNumTasks = static_cast<u32>(handles.size());

  ctp::ipc::FullPtr<TaskT> *handle_dev =
      ctp::GpuApi::MallocHost<ctp::ipc::FullPtr<TaskT>>(kNumTasks);
  for (u32 i = 0; i < kNumTasks; ++i) handle_dev[i] = handles[i];

  IpcManagerGpuInfo info = ipc->GetGpuIpcManager()->GetGpuInfo(gpu_id);
  DemoKernel<<<kNumTasks, 32>>>(info, handle_dev, kNumTasks);
  ctp::GpuApi::Synchronize();

  ctp::GpuApi::FreeHost(handle_dev);
}
#endif

4. Launch (SYCL)​

The SYCL source additionally does #include <sycl/sycl.hpp> and is compiled with icpx -fsycl:

#include <new>
#include <vector>
#include <clio_runtime/clio_runtime.h>
#include <clio_runtime/ipc_manager.h>
#include <clio_runtime/gpu/gpu_ipc_manager.h>
#include <clio_runtime/MOD_NAME/MOD_NAME_tasks.h>

using namespace clio::run;

#if CTP_ENABLE_SYCL
void launch_sycl(
    std::vector<ctp::ipc::FullPtr<MOD_NAME::GpuSubmitTask>> &handles,
    u32 gpu_id) {
  using TaskT = MOD_NAME::GpuSubmitTask;
  auto *ipc = CLIO_CPU_IPC;
  const u32 kNumTasks = static_cast<u32>(handles.size());

  sycl::queue q{sycl::gpu_selector_v};
  auto *info_storage = sycl::malloc_shared<IpcManagerGpuInfo>(1, q);
  *info_storage = ipc->GetGpuIpcManager()->GetGpuInfo(gpu_id);

  auto *ipc_storage = sycl::malloc_shared<gpu::IpcManager>(1, q);
  new (ipc_storage) gpu::IpcManager();

  auto *handle_storage =
      sycl::malloc_shared<ctp::ipc::FullPtr<TaskT>>(kNumTasks, q);
  for (u32 i = 0; i < kNumTasks; ++i) handle_storage[i] = handles[i];

  auto *slot_storage = sycl::malloc_shared<u32>(1, q);

  for (u32 i = 0; i < kNumTasks; ++i) {
    *slot_storage = i;
    q.submit([&](sycl::handler &cgh) {
      cgh.single_task<class clio_demo_kernel>([=]() {
        CLIO_GPU_INIT(*info_storage, ipc_storage);
        auto fp = handle_storage[*slot_storage];
        auto fut = CLIO_IPC->Send(fp);
        fut.Wait();
        (void)g_ipc_manager;
      });
    }).wait_and_throw();
  }
}
#endif

Unlike CUDA / ROCm, the SYCL path expects each task submission to be its own single_task. The CPU GPU worker pops them concurrently from the multi-MPSC ring buffer regardless of producer concurrency.

5. Verify and clean up​

#include <vector>
#include <clio_runtime/clio_runtime.h>
#include <clio_runtime/ipc_manager.h>
#include <clio_runtime/gpu/gpu_ipc_manager.h>
#include <clio_runtime/MOD_NAME/MOD_NAME_tasks.h>

using namespace clio::run;

void example(
    std::vector<ctp::ipc::FullPtr<MOD_NAME::GpuSubmitTask>> &handles,
    u32 gpu_id) {
  for (u32 i = 0; i < handles.size(); ++i) {
    // chimod formula: result = test_value * 2 + gpu_id
    u32 expected = (i * 2u) + gpu_id;
    if (handles[i]->result_value_ != expected) { /* mismatch */ }
  }

#if CTP_ENABLE_GPU
  auto *ipc = CLIO_CPU_IPC;
  ctp::ipc::AllocatorId alloc_id;  // returned by AllocateAndRegisterGpuBackend
  ipc->FreeGpuBackend(gpu_id, alloc_id);
#endif
  CLIO_RUNTIME_FINALIZE();
}

Tests​

The reference tests live under context-runtime/test/unit/gpu/:

FileBackendCTest name
test_gpu_kernel_stress_gpu.ccCUDAcr_gpu_kernel_stress_cuda
test_gpu_kernel_stress_gpu.ccROCmcr_gpu_kernel_stress_rocm
test_gpu_kernel_stress_sycl.ccSYCLcr_gpu_kernel_stress_sycl

Both .cc files share test_gpu_kernel_stress_common.h for host-side setup. Run them via:

cd build
ctest -R cr_gpu_kernel_stress -V

For the CTE-specific GPU device-memory round-trip, see context-transfer-engine/test/unit/gpu/test_cte_devmem_putget.cc (CTest name cte_devmem_putget_cuda), which exercises the same producer-only path with kDeviceMem backends and a real Module (PutBlobTask / GetBlobTask).

Troubleshooting​

Task dropped with fut_.task_size_ == 0​

Cause: the kernel pushed a slot whose fut_.task_size_ was zero. The host forgot to stamp task->fut_.task_size_ = sizeof(TaskT) when it placement-new'd the task.

Fix: always stamp task->fut_.task_size_ = sizeof(TaskT) (the PlaceTaskSlots helper does this) and submit through CLIO_IPC->Send(task_fp) (or IpcGpu2Cpu::SendIn directly).

CUDA error 700 ("an illegal memory access was encountered")​

Usually means the kernel dereferenced a FullPtr whose off_ pointed into a backend the runtime never registered. Re-check that:

  • AllocateAndRegisterGpuBackend returned a non-null AllocatorId.
  • The FullPtr you pass to Send carries either the returned alloc_id (when slots live inside the backend) or a null alloc_id with the raw device address in off_ (when you registered the backend as kDeviceMem and the kernel addresses it directly — the worker detects it via ctp::IsDevicePointer).

Kernel hangs in future.Wait()​

The worker dropped or refused to dispatch the task. Inspect the runtime log for any Worker {}: ProcessNewTaskGpu: errors — common causes are task_pod_size {} exceeds scratch capacity {} (POD too large; the worker's thread-local scratch is 4 KiB), or Container not found (pool=...) (the pool was not created before the kernel launched).

GetGpuQueueCount() == 0 after CLIO_INIT​

ServerInitGpuQueues runs as part of CLIO_INIT(RuntimeMode::kServer) only when the build includes at least one of CLIO_CORE_ENABLE_CUDA, CLIO_CORE_ENABLE_ROCM, or CLIO_CORE_ENABLE_SYCL. Verify the option is on, and that ipc->GetGpuQueueCount() > 0 before launching kernels.

SYCL Unexpected kernel lambda size​

DPC++ requires the host pass and device pass to lay out captures identically. Touch every captured pointer in both passes — e.g. (void)info_storage; (void)ipc_storage; — even if the host pass body otherwise does nothing. See the comment in test_gpu_kernel_stress_sycl.cc for context.

  • context-runtime/include/clio_runtime/clio_runtime.h — CLIO_INIT + RuntimeMode enum.
  • context-runtime/include/clio_runtime/pool_query.h — routing modes.
  • context-runtime/include/clio_runtime/gpu/gpu_ipc_manager.h — the central GPU API (clio::run::gpu::IpcManager, CLIO_GPU_INIT).
  • context-runtime/include/clio_runtime/gpu/future.h — gpu::Future / GpuTaskQueue.
  • context-runtime/include/clio_runtime/gpu/gpu_info.h — IpcManagerGpuInfo.
  • context-runtime/include/clio_runtime/task.h — Task, FutureInfo (fut_), TaskResume.
  • context-runtime/include/clio_runtime/ipc/ipc_gpu2cpu.h — SendIn / RecvIn / SendOut.
  • context-transport-primitives/include/clio_ctp/util/gpu_api.h — ctp::GpuApi (Malloc/Memcpy/Free/Synchronize).
  • context-runtime/src/worker.cc — Worker::ProcessNewTaskGpu.