Adapters (torch → CaliperTensor)

caliper/adapters/torch.hpp is a header-only bridge from a torch::Tensor to the frozen CaliperTensor (PLATFORM.md §7.2). It compiles against the applet's libtorch and is never included by the host: torch stays out of the ABI (D11). The host links no torch — it consumes only the CaliperTensor this adapter produces, then hands it to caliper.tensor_bridge.v1.

Three entry points, one ladder:

Entry point	Sync behavior	Use when
to_tensor(t)	none	CPU tensors, or you own the ordering yourself
synced_to_tensor(t)	full device drain before handoff	device tensors on a v1 host (always correct, costs a barrier)
stream_to_tensor(t, caps)	drain or stream-ordered handoff, negotiated	device tensors; pass bridge.caps() and get the best rung the host honors

Reject, never copy

The adapter's central contract: it never silently copies. Each entry point returns std::optional<CaliperTensor> that aliases the tensor's memory (zero-copy) on success, and std::nullopt — copying nothing — when the tensor cannot be represented as a v1 CaliperTensor. Rejection reasons:

• an unsupported dtype (only f32, f16, bf16, i64, i32, u8 map; f64, bool, complex, quantized, fp8 are rejected);
• ndim > 8 (the frozen shape/strides arrays are [8]);
• non-contiguous (the bridge assumes row-major with no gaps);
• an MPS view with a nonzero storage offset (see below);
• a device other than CPU, MPS, or CUDA.

The caller repairs the tensor in applet code — an explicit .contiguous() before calling again — so the cost of any copy is visible where it is paid, not hidden inside the bridge. This is the teaching point the exemplar exists to make.

The MPS offset-0 / contiguous rule

An MPS tensor must additionally have storage_offset() == 0. The reason is structural, not conservative: the frozen CaliperTensor has no storage-offset channel, and the bridge casts storage().mutable_data() straight to an id<MTLBuffer>. A view with a nonzero offset would leave the buffer starting at the wrong address and silently address the wrong texels. So the adapter rejects it, and the caller clones the view (.clone(), or .contiguous() of a materialized tensor) so the buffer starts at offset 0.

MLScope hits exactly this: a per-kernel select() slice of the (8,1,3,3) weight carries a nonzero offset, so the worker takes an owned clone of each (3,3) filter — which also decouples the snapshot from the still-mutating live weight.

CUDA tensors (Windows)

A CUDA tensor maps to CALIPER_DEV_CUDA with data = the device pointer; the Vulkan backend imports it via external-memory interop (see ZEROCOPY.md). The CUDA branch is compiled under

#if !defined(__APPLE__) && __has_include(<c10/cuda/CUDAStream.h>)

— __has_include alone is not enough because mac libtorch ships the c10/cuda headers without the CUDA toolkit headers they include. If this guard is ever wrong in the compiled-out direction, everything stays green while the drain silently returns; the torch test suite carries a tripwire case that fails loudly instead (stream must round-trip a non-default pool stream).

synced_to_tensor — cost, honestly

synced_to_tensor is to_tensor preceded by a device drain, so the texture the host uploads this frame reflects every kernel the applet has enqueued (sync-then-update — the v1 correctness story for device textures). The honest cost: torch::mps::synchronize() / torch::cuda::synchronize() is a full device barrier — it blocks the calling CPU thread until the device drains. Pay it once at the handoff, not per frame. CPU tensors need no sync, so the call is skipped for them.

stream_to_tensor — the negotiated handoff (bridge v1.1)

const uint32_t caps = bridge.caps();          // tensor_bridge.v1_1, additive
auto ct = caliper::adapters::stream_to_tensor(t, caps);

When the host grants CALIPER_BRIDGE_CAP_STREAM_ORDERED (caps bit 0), the adapter skips the drain and instead publishes the producer's queue in CaliperTensor.stream — an MTLCommandQueue* on MPS, a cudaStream_t on CUDA. The renderer GPU-orders its copy after the producer's queued work (per-texture MTLSharedEvent on Metal; a shared timeline semaphore riding the producer stream on Vulkan+CUDA), so no CPU thread waits. Without the caps bit, stream_to_tensor(t, 0) is synced_to_tensor — byte-identical v1 behavior, which is also the negotiation pin the tests hold.

CUDA nuance — NULL can be an honored handoff: torch's default stream handle is literally nullptr (legacy default stream), so a t.stream == NULL from a CUDA producer still orders correctly — the renderer's NULL rung uses that same default stream. Only a producer on a non-default stream carries a non-NULL handle. Don't write assertions that assume otherwise.

Thread safety — the MPS serialization rule

None of torch's public MPS stream calls are internally serialized (proven by disassembly; the crashes were real: command-buffer corruption when the frame thread drained or committed while the training thread encoded). The adapter therefore runs the entire MPS portion of both rungs — the v1 drain (8b0a010) and the stream handoff (545a2f7) — as one block on torch's own stream dispatch queue. Two consequences for applet authors:

1. Calling the adapter from any thread is safe on the adapter's own operations — but the rule extends to you: any additional raw MPS/Metal calls you make must also ride that dispatch queue.
2. CUDA has no analogous rule (driver calls are thread-safe by contract), and this is verified empirically by a concurrency stress test rather than assumed — the MPS lesson was exactly that "should be safe" isn't evidence.

In the exemplar the training worker performs the handoff when it snapshots, so the frame thread never syncs during upload.

Full source

#pragma once
/* caliper/adapters/torch.hpp — header-only bridge from a torch::Tensor to the
 * frozen CaliperTensor (PLATFORM.md §7.2). It compiles against the APPLET's
 * libtorch and is NEVER included by the host: torch stays out of the ABI (D11).
 * The host links no torch; it consumes only the CaliperTensor this produces.
 *
 * Exemplar teaching points baked into the contract:
 *   - The adapter NEVER silently copies. A non-contiguous tensor is REJECTED
 *     (std::nullopt); the caller makes the copy visible in applet code with an
 *     explicit `.contiguous()` before calling again — the cost is in the applet,
 *     not hidden in the bridge.
 *   - MPS tensors additionally REQUIRE storage_offset() == 0. The frozen
 *     CaliperTensor has no storage-offset channel, and the bridge casts
 *     storage().mutable_data() straight to an id<MTLBuffer>; a view with a
 *     nonzero offset would silently address the wrong texels. Reject it — the
 *     caller clones the view (`.clone()` / `.contiguous()` of a materialized
 *     tensor) so the buffer starts at offset 0.
 *   - v1 device story: `stream == NULL`. Correctness on device textures comes
 *     from torch::mps::synchronize() (see synced_to_tensor), i.e. sync-then-
 *     update. That sync is a FULL device barrier — it blocks the calling CPU
 *     thread until every MPS stream drains — so pay it once at the handoff, not
 *     needlessly per frame. stream_to_tensor (M2/D24) supersedes this when the
 *     host's bridge-v1.1 caps() grants stream-ordered handoff.
 */
#include <optional>

#include <caliper/tensor.h>
#include <caliper/services/tensor_bridge_v1_1.h>
#include <torch/torch.h>
#if defined(__APPLE__)
#include <dispatch/dispatch.h>  // dispatch_sync_f: serialize on torch's MPS stream queue
#include <objc/message.h>       // producer-queue lookup without an ObjC++ TU
#endif
// !__APPLE__: mac libtorch ships the c10/cuda headers but no CUDA toolkit
// headers, so presence-of-header alone is not usability — and Apple torch
// builds never have CUDA anyway.
#if !defined(__APPLE__) && __has_include(<c10/cuda/CUDAStream.h>)
#include <c10/cuda/CUDAStream.h>
#endif

namespace caliper::adapters {

namespace detail {

// v1 dtype map. Anything not named here (f64, bool, complex, quantized, fp8…)
// is rejected — the bridge only understands these six.
inline std::optional<CaliperDType> map_dtype(at::ScalarType st) {
    switch (st) {
        case at::kFloat:    return CALIPER_DT_F32;
        case at::kHalf:     return CALIPER_DT_F16;
        case at::kBFloat16: return CALIPER_DT_BF16;
        case at::kLong:     return CALIPER_DT_I64;
        case at::kInt:      return CALIPER_DT_I32;
        case at::kByte:     return CALIPER_DT_U8;
        default:            return std::nullopt;
    }
}

#if defined(__APPLE__)
// [cb commandQueue] via the C ObjC runtime, so this header stays compilable
// as plain C++ (applet TUs are .cpp, not .mm). The queue is torch's global
// MPS command queue — process-lifetime, safe to hand across the ABI.
inline void* mtl_command_queue_of(void* command_buffer) {
    using Send = void* (*)(void*, SEL);
    return command_buffer
        ? ((Send)objc_msgSend)(command_buffer, sel_registerName("commandQueue"))
        : nullptr;
}

// The whole MPS handoff — command-buffer peek AND commit — as ONE block on
// torch's MPS stream dispatch queue. NONE of the torch::mps stream calls are
// internally serialized (verified by disassembly of this libtorch:
// commitStream/deviceSynchronize tail-call MPSStream::synchronize, which is
// straight-line objc_msgSends — no dispatch_sync anywhere), while every
// torch-internal kernel encode runs as a block on get_dispatch_queue() (its
// documented purpose is exactly this synchronization). Touching the stream
// from a frame thread while a training thread encodes therefore corrupts the
// MPSCommandBuffer/encoder state — MPS aborts with 'command buffer already
// committed' / AGX encoder-coalescing crashes (the EmbedScope SIGABRT). The
// same disassembly is what makes nesting commit() inside our block safe: it
// cannot deadlock, because it never dispatches. Plain C dispatch API — no
// ObjC blocks — keeps the header .cpp-compilable.
inline void mps_handoff_probe(void* ctx) {
    void** queue = static_cast<void**>(ctx);
    *queue = mtl_command_queue_of(torch::mps::get_command_buffer());
    if (*queue != nullptr)
        torch::mps::commit();   // enqueue-not-drain, atomic with the peek
}
inline void* mps_commit_and_get_queue() {
    void* queue = nullptr;
    if (void* dq = torch::mps::get_dispatch_queue())
        dispatch_sync_f(static_cast<dispatch_queue_t>(dq), &queue,
                        &mps_handoff_probe);
    return queue;
}

// The drain-path twin of the above, for synced_to_tensor and the null-queue
// fallback: torch::mps::synchronize() is just as unserialized as commit()
// (same disassembly), so a frame-thread drain races worker encodes the same
// way — the third face of that race is an AGX encoder-coalescing SIGSEGV.
// Same recipe: run the drain as ONE block on torch's stream dispatch queue,
// atomic with worker encodes; nesting synchronize() inside is deadlock-free
// because it never dispatches.
inline void mps_sync_block(void*) { torch::mps::synchronize(); }
inline void mps_synchronize_serialized() {
    if (void* dq = torch::mps::get_dispatch_queue())
        dispatch_sync_f(static_cast<dispatch_queue_t>(dq), nullptr,
                        &mps_sync_block);
    else
        torch::mps::synchronize();
}
#endif

}  // namespace detail

// Build a CaliperTensor that aliases `t`'s memory (zero-copy). Returns nullopt
// — and copies nothing — when `t` cannot be represented as a v1 CaliperTensor.
// Rejection reasons: unsupported dtype, ndim > 8, non-contiguous, an MPS view
// with a nonzero storage offset, or a device other than CPU/MPS. The caller
// logs / repairs (e.g. `t.contiguous()`); the adapter never hides a copy.
inline std::optional<CaliperTensor> to_tensor(const at::Tensor& t) {
    const auto dt = detail::map_dtype(t.scalar_type());
    if (!dt) return std::nullopt;

    const int64_t nd = t.dim();
    if (nd < 0 || nd > 8) return std::nullopt;   // frozen shape/strides are [8]

    // Contiguity is mandatory on every device: the bridge assumes row-major with
    // no gaps. Reject rather than copy so the cost is visible in applet code.
    if (!t.is_contiguous()) return std::nullopt;

    CaliperTensor out{};
    out.struct_size = sizeof(CaliperTensor);
    out.dtype = *dt;
    out.ndim = static_cast<int32_t>(nd);
    for (int64_t i = 0; i < nd; ++i) {
        out.shape[i]   = t.size(i);      // elements
        out.strides[i] = t.stride(i);    // elements
    }
    out.stream = nullptr;                // v1: no stream channel (sync explicitly)

    if (t.is_cpu()) {
        out.data = t.data_ptr();
        out.device = CALIPER_DEV_CPU;
        out.device_index = 0;
        return out;
    }
    if (t.is_mps()) {
        // No offset channel in CaliperTensor and the bridge casts this pointer
        // straight to id<MTLBuffer>; a nonzero offset would mis-address texels.
        if (t.storage_offset() != 0) return std::nullopt;
        out.data = t.storage().mutable_data();   // the MTLBuffer bridge pointer
        out.device = CALIPER_DEV_METAL;
        out.device_index = static_cast<int32_t>(t.device().index());
        return out;
    }
    if (t.is_cuda()) {
        // Unlike MPS there is no buffer-object cast: data_ptr() IS the device
        // address (storage offset already applied), and the Vulkan backend
        // copies from it in-VRAM (ZEROCOPY.md). Contiguity was enforced above;
        // views with a nonzero storage offset are therefore fine here.
        out.data = t.data_ptr();
        out.device = CALIPER_DEV_CUDA;
        out.device_index = static_cast<int32_t>(t.device().index());
        return out;
    }
    // Any other device is not a v1 target for this adapter.
    return std::nullopt;
}

// Same as to_tensor, but first drains the MPS device so the texture the host
// uploads THIS frame reflects every kernel the applet has enqueued. This is the
// v1 correctness story for device textures (sync-then-update). Cost, honestly:
// torch::mps::synchronize() is a full device barrier that blocks the CPU until
// all MPS streams complete — the price of a stream-free v1 ABI. CPU tensors
// need no sync, so the call is skipped for them.
inline std::optional<CaliperTensor> synced_to_tensor(const at::Tensor& t) {
#if defined(__APPLE__)
    // Serialized on torch's MPS stream dispatch queue — a bare synchronize()
    // races concurrent training-thread encodes (see the detail helper).
    if (t.is_mps()) detail::mps_synchronize_serialized();
#else
    if (t.is_mps()) torch::mps::synchronize();
#endif
    if (t.is_cuda()) torch::cuda::synchronize();   // same contract, CUDA form
    return to_tensor(t);
}

// M2 (docs/metal-pipelining.md §4, D24): hand over ORDER instead of a drained
// device. When the host's bridge-v1.1 caps carry
// CALIPER_BRIDGE_CAP_STREAM_ORDERED, populate t.stream with the producer's
// stream/queue and SKIP the full-device drain; the renderer GPU-orders its
// update after the producer's already-enqueued work. Without the bit (v1
// host, GL fallback, headless) this is exactly synced_to_tensor — the adapter
// never skips a drain the host didn't promise to replace. Thread-safety story
// unchanged from v1: the caller still hands over a tensor it owns at a
// quiescent point in its own logic (spec §4, last paragraph).
//
// FRAME-THREAD WARNING: without CALIPER_BRIDGE_CAP_STREAM_ORDERED this degrades
// to synced_to_tensor (a full device barrier). Gate on the cap before calling
// from draw_ui.
inline std::optional<CaliperTensor> stream_to_tensor(const at::Tensor& t,
                                                     uint32_t bridge_caps) {
    if (!(bridge_caps & CALIPER_BRIDGE_CAP_STREAM_ORDERED))
        return synced_to_tensor(t);
#if defined(__APPLE__)
    if (t.is_mps()) {
        auto out = to_tensor(t);
        if (!out) return out;
        // One dispatch-serialized block does the queue peek AND the commit —
        // atomic against concurrent training-thread encodes (see the detail
        // helper's comment for why nothing here may touch the stream outside
        // that block). Enqueue, not drain: pending torch kernels become
        // committed GPU work the renderer's producer-queue signal is ordered
        // after (M2b). The queue outlives the handoff (process-lifetime).
        void* queue = detail::mps_commit_and_get_queue();
        if (queue == nullptr) { detail::mps_synchronize_serialized(); return out; }
        out->stream = queue;
        return out;
    }
#endif
#if !defined(__APPLE__) && __has_include(<c10/cuda/CUDAStream.h>)
    if (t.is_cuda()) {
        auto out = to_tensor(t);
        if (!out) return out;
        // Stream order puts the renderer's DtoD copy after the producer's
        // kernels — torch::cuda::synchronize() elided entirely (M2a).
        out->stream = (void*)at::cuda::getCurrentCUDAStream(
                          t.device().index()).stream();
        return out;
    }
#endif
    return synced_to_tensor(t);   // CPU: no sync needed; unknown devices drain
}

}  // namespace caliper::adapters