---
title: 'Tessera: Kernel-Granularity GPU Disaggregation'
url: https://www.emergentmind.com/topics/tessera-475d4a79-b49c-4e6f-a6bf-e357eeef9777
type: topic
---

# Tessera: Kernel-Granularity GPU Disaggregation

Searching arXiv for the specified paper and closely related heterogeneous GPU disaggregation work.
TESSERA is a kernel-disaggregation system for large-model inference on heterogeneous GPU clusters. It addresses a specific limitation of prior disaggregation schemes: existing solutions operate at coarse granularity and are tightly coupled to specific model architectures, which leaves substantial performance headroom unused. The central claim of Tessera is that kernels within a single application exhibit diverse resource demands and therefore form the right granularity for aligning computation with heterogeneous hardware capabilities. On that basis, the system combines PTX-level offline analysis, pipelined multi-request execution, mixed-integer programming (MILP) scheduling, and lightweight runtime adaptation to improve serving throughput and cost efficiency on heterogeneous GPUs [2604.10180].

## 1. Kernel-granularity disaggregation as a systems abstraction

Disaggregation, in this setting, means mapping different parts of an AI inference workload to different GPU types. Tessera departs from phase-level and block-level partitioning by selecting the GPU assignment at the level of individual GPU kernels rather than at the level of entire model phases such as prefill and decode, or architectural blocks such as attention and FFN. The paper characterizes this as the first kernel disaggregation system for large-model inference on heterogeneous GPUs [2604.10180].

The rationale is empirical rather than merely architectural. Kernel-level heterogeneity is reported to be substantial: for the A100+L40s pair, 67% of kernels run faster on L40s vs. A100, and these kernels account for up to 53% of end-to-end time. This directly challenges the common assumption that a nominally weaker or cheaper accelerator is uniformly inferior for inference serving. Tessera’s design therefore treats the inference graph not as a monolithic program or as a small number of coarse phases, but as a sequence of kernels whose compute, bandwidth, and communication properties can differ enough to justify heterogeneous placement.

A plausible implication is that kernel granularity is not only a performance optimization but also a portability strategy. Because the assignment is driven by measured kernel behavior and inter-kernel dependencies rather than by hand-coded assumptions about attention, FFN, or modality-specific structure, the method can generalize to model architectures where prior disaggregation approaches do not apply [2604.10180].

## 2. System organization and execution model

Tessera comprises four components: an Offline Kernel Analyzer, a Policy Planner, GPU Workers with Pipelined Execution, and an Online Monitor for Adaptation [2604.10180]. This decomposition separates correctness-critical static analysis from throughput-critical runtime execution.

The runtime model is explicitly pipeline-oriented. Batched requests are dispatched across heterogeneous GPUs, with inter-GPU transfers inserted between kernels when assignments cross device boundaries. Compute and communication run on separate CUDA streams with CUDA-event synchronization. Each GPU worker maintains one compute stream per in-flight request and one dedicated communication stream. When a kernel assigned to one GPU depends on outputs produced by another GPU, the worker issues all `ncclRecv` operations on the communication stream, records a receive event, launches the kernel on the compute stream while waiting on that event, and then issues the corresponding `ncclSend` operations after kernel completion [2604.10180].

This organization supports multi-request pipelining. If one stream stalls on a receive event, the hardware scheduler can dispatch work from other streams whose dependencies are already satisfied. Tessera also introduces priority-aware staggering: later arrivals are assigned lower CUDA stream priority so that their communication phases do not align, avoiding the “all-streams-stall” pattern. The system-level goal is therefore not only to minimize raw communication volume, but to convert unavoidable communication into latency that can be hidden behind useful work.

## 3. PTX-level analysis and dependency extraction

Correct kernel-level disaggregation requires exact knowledge of inter-kernel data dependencies. Tessera obtains this information by PTX-level instrumentation. For each JIT or user-defined PTX kernel, it injects two extra global variables, `min_addr` and `max_addr`, and instruments every global load and store. The paper gives the representative pattern

```text
ld.global.f32   %f1, [%rd1]                // original
atom.global.min.u64 [%min_buf], %rd1       // track
atom.global.max.u64 [%max_buf], %rd1
```

After one run, the `(min_addr, max_addr)` pair per instruction yields the exact `[base, base+size]` of each buffer, while instruction type distinguishes read and write sets [2604.10180].

From these traces Tessera constructs a data-dependency graph (DDG). In simplified form, the procedure iterates over kernels in launch order, checks each accessed buffer, adds an edge from the last writer to the current reader when a read-after-write dependency exists, and updates the last writer on writes. The resulting graph is a DAG $G=(K,E)$ where $K$ contains kernels plus async-memcpy nodes and $E$ contains RAW edges $(i \to j)$ annotated with transfer size $d_{ij}$ [2604.10180]. In the same profiling pass, Tessera measures $t_{k,1}, t_{k,2}, \ldots, t_{k,|G|}$, the execution time of kernel $k$ on each GPU.

This PTX-derived DDG is the correctness backbone of the system. Kernel placement is only meaningful if the inserted communication preserves the true dependency structure; otherwise, any throughput gain would be invalidated by semantic errors. The paper therefore positions dependency extraction not as a heuristic graph recovery mechanism but as a precise prerequisite for safe kernel migration across devices.

## 4. Scheduling, MILP objectives, and online policy switching

Tessera’s scheduler is workload-aware and uses MILP formulations tailored to two operating regimes. Let $G$ be the set of GPUs, $K$ the set of kernels, and $E$ the set of dependency edges. The measured kernel time on GPU $g$ is $t_{k,g}$; the data size on edge $(i \to j)$ is $d_{ij}$; and inter-GPU bandwidth and base latency from GPU $u$ to GPU $g$ are $bw_{u,g}$ and $\ell_{u,g}$, respectively. The binary placement variable is $x_{k,g} \in \{0,1\}$ with $\sum_g x_{k,g}=1$, and $y_{ij}^{u,g}$ linearizes the product $x_{i,u}\times x_{j,g}$ [2604.10180].

For throughput-oriented scheduling, Tessera defines the computation load on GPU $g$ as
$$
T_g = \sum_k t_{k,g}\cdot x_{k,g},
$$
the communication load as
$$
M_g = \sum_{(i\to j)\in E}\sum_{u\neq g} \left(\ell_{u,g} + d_{ij}/bw_{u,g}\right)\cdot y_{ij}^{u,g},
$$
and the stage time as
$$
W_g = \max(T_g, M_g).
$$
The objective is
$$
\min \max_{g\in G} W_g.
$$
This objective is explicitly aligned with a pipelined serving regime in which communication can be overlapped and the effective bottleneck is the slowest stage across devices [2604.10180].

For latency-oriented scheduling, used under low concurrency where communication cannot be hidden, Tessera minimizes the total computation-plus-communication time:
$$
\sum_k \sum_g t_{k,g} x_{k,g}
+
\sum_{(i\to j)} \sum_{u\neq g}
\left(\ell_{u,g} + d_{ij}/bw_{u,g}\right) y_{ij}^{u,g}.
$$

The online monitor selects between these two policies every $W$ ms by comparing end-to-end queuing-sensitive latency to pure execution-plus-communication latency. Specifically, it measures $L_{\text{req}}$ and $L_{\text{exec}}$ and switches to the throughput-oriented policy when
$$
L_{\text{req}} / L_{\text{exec}} > \beta,
$$
otherwise using the latency-oriented policy. The defaults are $W=300$ ms and $\beta=1.5$ [2604.10180]. This adaptation mechanism is intentionally lightweight: it does not re-profile kernels online, but it does react to concurrency changes that alter the value of overlap.

## 5. Experimental scope and quantitative results

The evaluation covers five heterogeneous GPUs and four model architectures, scaling up to 16 GPUs [2604.10180]. Single-node GPU-pair setups include A100+L40s, H100+RTX Pro 6000, and B200+H100, all with 200 Gbps RDMA. Cluster-scale configurations include 2 A100 + 1 L40s and 8 B200 + 8 H100, with cross-node 200/400 Gbps and within-node NVLink [2604.10180].

The workloads span four architecture classes: LLMs (Llama3-8B, GPT-oss-20B), SSM (Mamba-Codestral-7B), MLLM (Qwen2.5-VL-7B), and diffusion (Stable Diffusion 3.5). Baselines are homogeneous execution without disaggregation, phase-level Prefill–Decode disaggregation, and block-level Attention–FFN disaggregation where applicable [2604.10180].

The main reported results are as follows. Offline throughput, measured as tokens/sec or images/min, improves by +1.5× on average and up to 2.3× versus Prefill–Decode disaggregation, and by +1.35× on average and up to 1.7× versus Attention–FFN disaggregation. Cost efficiency improves by up to 1.6× compared to existing disaggregation methods [2604.10180].

The cost-efficiency table reported in the paper, normalized to the left GPU alone, further illustrates the effect. For A100 + L40s, Tessera reaches 1.21 versus 0.87 for Prefill–Decode and 1.02 for Attention–FFN. For H100 + RTX Pro 6000, Tessera reaches 1.64 versus 1.00 and 1.07. For B200 + H100, Tessera reaches 1.01 versus 0.70 and 0.74 [2604.10180]. These numbers formalize the paper’s broader argument that heterogeneous pairing can improve performance-per-dollar rather than merely salvage underutilized hardware.

Under online latency evaluation with GPT-oss-20B, at an SLO of 50 ms/token, Tessera supports 1.3× more QPS before violation. At cluster scale, for Qwen-3 235B on 8 B200 + 8 H100, Tessera composed with tensor parallelism outperforms phase- and block-level baselines by 1.4–1.5× [2604.10180].

## 6. Ablations, limitations, and broader significance

The ablation studies isolate the effect of the pipeline and the online monitor. Pipeline plus priority-aware staggering hides 96.6% of communication idle time, compared with approximately 50% without pipelining. For online adaptation, switching too frequently with $W=30$ ms increases latency by 1.3×, while switching too slowly with $W=1.5$ s increases latency by 1.55× [2604.10180]. These results indicate that policy adaptation is useful only when it operates on a timescale commensurate with workload variation.

Robustness to degraded interconnect is also evaluated. Even under 25 Gbps throttling, offline throughput drops by less than 6%, and the online policy can revert to homogeneous execution when needed [2604.10180]. This is notable because kernel-granularity disaggregation might intuitively appear highly sensitive to communication bandwidth; the reported measurements suggest that accurate dependency extraction and overlap can substantially reduce that sensitivity.

The limitations are explicit. Indirect memory accesses via globals not present in kernel arguments fall back to no disaggregation, although the paper describes such cases as very rare in modern frameworks. The current MILP does not model per-GPU memory capacity, and the full model is replicated widely; future work is proposed around pruning unused weights or integrating capacity constraints. The implementation focus is NVIDIA PTX, so extending the approach to AMD or NPUs would require unified interconnect and software APIs [2604.10180].

In broader systems terms, Tessera reframes heterogeneous inference scheduling around the microstructure of kernel behavior rather than around coarse model semantics. Its most counterintuitive result—that a heterogeneous GPU pair under Tessera can exceed the throughput of two homogeneous high-end GPUs at lower cost—does not imply that weaker devices dominate stronger ones in general. Rather, it suggests that end-to-end serving efficiency depends on the match between kernel resource demands and hardware specialization, plus the ability to extract and exploit enough parallel slack to hide communication [2604.10180]. That interpretation places Tessera at the intersection of compiler-level dependency recovery, multi-accelerator runtime design, and optimization-based scheduling for inference systems.

Source: https://www.emergentmind.com/topics/tessera-475d4a79-b49c-4e6f-a6bf-e357eeef9777