Your Agentic Stack Waste Too Many GPU Kernels

Once an agentic product moves beyond a hosted API demo, you hit a bottleneck that dominates your agentic loop: the local neural components around the agent are often doing too much GPU work for too little useful computation.

A production agent may call an embedding model, reranker, guardrail classifier, vision encoder, policy model, reward model, or local language model during one user request.

Each component can contain long chains of simple tensor operations.

In eager PyTorch, every operation can become a separate GPU kernel launch.

That means code as ordinary as this:

tmp = x * w
tmp = tmp + b
tmp = tmp.sigmoid()

may execute as three kernels, materialize two temporary tensors, and repeatedly move data through GPU global memory.

The arithmetic is cheap whereas orchestration is expensive.

PyTorch’s Inductor compiler attacks this problem with kernel fusion.

It captures compatible operations, groups them, and generates a single optimized kernel, commonly in Triton for NVIDIA GPUs. The intermediate values can remain close to the GPU’s execution units instead of being written to and read from global memory.

PyTorch has reported workloads running up to 10x faster with compilation, although that is an upper-end result rather than a number you should assume for every model.

The more useful question is:

Can the compiler remove enough launches and memory traffic from my actual agent workload to improve steady-state latency or throughput?

This article shows how to answer that question with code.

The performance problem hiding inside agent loops

An agent request is a latency chain.

The system may parse an input, create embeddings, search a vector store, rerank documents, run safety checks, call a reasoning model, execute tools, and score the result.

Even if the main language model is remote, local PyTorch models frequently sit on the critical path.

These models do not need to be huge to be expensive.

Small models can be especially sensitive to launch overhead because each individual kernel finishes quickly.

When the useful GPU work is short, the fixed cost of repeatedly launching kernels becomes a larger share of the total runtime.

The same issue appears with memory-bound operations.

Elementwise math such as addition, multiplication, activation functions, masking, normalization, and logits transformations performs relatively little computation per byte moved.

Reading and writing full tensors can cost more than the math itself.

Consider an agentic retrieval service with a local cross-encoder reranker.

A request might score 50 candidate passages.

Inside the model, it runs normalization, linear projections, bias additions, activations, residual connections, masking, and reductions.

If every intermediate result is materialized in global memory, the GPU becomes a very expensive tensor courier.

Fusion changes that execution shape.

What torch.compile changes

The normal PyTorch experience is eager execution: Python calls a PyTorch operation, PyTorch dispatches it, and the GPU runs the corresponding work.

torch.compile introduces a compiler pipeline in front of that execution.

At a high level:

1. TorchDynamo captures regions of Python and PyTorch operations into graphs.
2. AOTAutograd can capture forward and backward computation for training workloads.
3. TorchInductor lowers and schedules the graph.
4. Inductor generates optimized code, including Triton kernels on supported GPU paths.

Unsupported or difficult-to-trace Python can create graph breaks, which reduce the compiler’s optimization surface rather than silently changing the meaning of the program.

The minimal API is deliberately boring:

compiled_fn = torch.compile(fn)

or:

@torch.compile
def fn(x):
  return x.sin() + x.cos()

The interesting part is the generated execution plan.

Vertical fusion: keep dependent operations together

Vertical fusion combines operations that depend on one another.

Picture a computation graph from top to bottom:

Each operation consumes the previous operation’s output.

In eager execution, the multiply result may be written to global memory, loaded by the add kernel, written again, and loaded by the sigmoid kernel.

With vertical fusion, the compiler can generate one kernel that:

1. Loads x, w, and b.
2. Computes multiply.
3. Computes add.
4. Computes sigmoid.
5. Stores only the final output.

The two intermediate values do not need standalone global-memory buffers.

This pattern matters because neural networks are full of dependent chains: normalization into projections, matrix multiplication into bias and activation, residual arithmetic, gating, masking, and reductions.

Editor’s note: If you want to master the agentic stack and GPU programming, you can join our Agent Foundry program for hands-on, in-depth trainings.

Pointwise fusion in ordinary PyTorch

The companion code for PyTorch’s kernel-fusion is available in the kernel-fusion GitHub gist.

Its core pointwise example is simple enough to appear in almost any neural component:

import torch

def pointwise_example(x, w, b):
  tmp = x * w
  tmp = tmp + b
  tmp = tmp.sigmoid()
  return tmp

There is nothing obviously inefficient about the Python code.

It is readable, vectorized, and uses standard PyTorch operations.

Without fusion, the conceptual Triton kernels look like this.

Kernel 1: multiply

@triton.jit
def mul_kernel(in_ptr0, in_ptr1, out_ptr0, xnumel, XBLOCK: tl.constexpr):
  xoffset = tl.program_id(0) * XBLOCK
  xindex = xoffset + tl.arange(0, XBLOCK)
  xmask = xindex < xnumel
  x = tl.load(in_ptr0 + xindex, mask=xmask)
  w = tl.load(in_ptr1 + xindex, mask=xmask)
  out = x * w
  tl.store(out_ptr0 + xindex, out, mask=xmask)

Kernel 2: add

@triton.jit
def add_kernel(in_ptr0, in_ptr1, out_ptr0, xnumel, XBLOCK: tl.constexpr):
  xoffset = tl.program_id(0) * XBLOCK
  xindex = xoffset + tl.arange(0, XBLOCK)
  xmask = xindex < xnumel
  x = tl.load(in_ptr0 + xindex, mask=xmask)
  b = tl.load(in_ptr1 + xindex, mask=xmask)
  out = x + b
  tl.store(out_ptr0 + xindex, out, mask=xmask)

Kernel 3: sigmoid

@triton.jit
def sigmoid_kernel(in_ptr0, out_ptr0, xnumel, XBLOCK: tl.constexpr):
  xoffset = tl.program_id(0) * XBLOCK
  xindex = xoffset + tl.arange(0, XBLOCK)
  xmask = xindex < xnumel
  x = tl.load(in_ptr0 + xindex, mask=xmask)
  out = tl.sigmoid(x)
  tl.store(out_ptr0 + xindex, out, mask=xmask)

The syntax follows Triton’s block-oriented programming model: identify a program instance, create offsets, mask out-of-bounds elements, load values, compute, and store.

Triton’s vector-add tutorial is a useful primer, while the language reference documents primitives such as program_id, arange, load, store, and sigmoid.

The important detail is the memory behavior.

• Multiply reads x and w, then writes the multiply result.
• Add reads the multiply result and b, then writes the add result.
• Sigmoid reads the add result, then writes the final output.
• Total: five full-tensor reads and three full-tensor writes.

That is eight full-tensor memory operations across three launches.

What the fused kernel looks like

Inductor can combine the three pointwise operations into one kernel similar to this:

@triton.jit
def triton_poi_fused_add_mul_sigmoid_0(
  in_ptr0,
  in_ptr1,
  in_ptr2,
  out_ptr0,
  xnumel,
  XBLOCK: tl.constexpr,
):
  xoffset = tl.program_id(0) * XBLOCK
  xindex = xoffset + tl.arange(0, XBLOCK)
  xmask = xindex < xnumel
  x = tl.load(in_ptr0 + xindex, mask=xmask)
  w = tl.load(in_ptr1 + xindex, mask=xmask)
  b = tl.load(in_ptr2 + xindex, mask=xmask)
  multiplied = x * w
  biased = multiplied + b
  activated = tl.sigmoid(biased)
  tl.store(out_ptr0 + xindex, activated, mask=xmask)

Now the kernel performs three reads and one write.

• Unfused: 3 kernel launches, 2 intermediate output buffers, 8 full-tensor memory operations.
• Fused: 1 kernel launch, 0 intermediate output buffers, 4 full-tensor memory operations.

That is a 50% reduction in the simple memory-operation count, plus two avoided kernel launches.

The variables multiplied and biased still exist logically, but they no longer need to become globally visible tensors.

The compiler can keep them in registers or otherwise within the generated kernel’s local execution.

This is the key mental model:

Fusion is not mainly about doing fewer mathematical operations. It is about making fewer trips through expensive memory and paying fewer launch costs.

Inspect fusion on your own GPU

The following workflow is deliberately small. Do it before trying to compile an entire agent stack.

1. Create an environment

Use a Linux environment with a supported NVIDIA GPU if you want to reproduce Triton CUDA code similar to the examples. Install the PyTorch build that matches your driver and CUDA setup using the official PyTorch local installation selector.

python -m venv .venv
source .venv/bin/activate
python -m pip install --upgrade pip
# Use the command produced by the official PyTorch install selector.
# A generic installation may look like this:
pip install torch

Verify the environment:

python - <<'PY'
import torch
print("PyTorch:", torch.__version__)
print("CUDA available:", torch.cuda.is_available())
if torch.cuda.is_available():
  print("GPU:", torch.cuda.get_device_name(0))
  print("CUDA runtime:", torch.version.cuda)
PY

torch.compile exists in PyTorch 2.x.

2. Create fusion_example.py

This is a corrected, self-contained version of the reduction workflow shown in the companion gist:

import torch

def reduction_example(x: torch.Tensor) -> torch.Tensor:
  # Pointwise prologue
  tmp = x * 2.0
  # Reduction
  result = tmp.sum(dim=-1)
  # Pointwise epilogue
  result = result + 1.0
  return result

def main() -> None:
  if not torch.cuda.is_available():
      raise SystemExit(
          "CUDA is required for this Triton GPU demonstration."
      )
  torch.manual_seed(0)
  x = torch.randn(1024, 1024, device="cuda")
  eager_result = reduction_example(x)
  compiled_fn = torch.compile(reduction_example)
  compiled_result = compiled_fn(x)
  torch.testing.assert_close(
      compiled_result, eager_result, rtol=1e-4, atol=1e-5
  )
  print("Correctness check passed.")
  print("Output shape:", tuple(compiled_result.shape))

if __name__ == "__main__":
  main()

Run it normally first:

python fusion_example.py

Compiler work should always begin from a known-good eager implementation, and the compiled output should be compared against it.

3. Print Inductor’s generated code

PyTorch’s logging system includes an output_code channel that prints code generated by Inductor. Run:

TORCH_LOGS="output_code" python fusion_example.py

This is the same workflow documented in PyTorch’s kernel-fusion article and [TORCH_LOGS](https://docs.pytorch.org/tutorials/recipes/torch_logs.html) recipe. Search the output for a name similar to:

triton_per_fused_add_mul_sum_0

The exact string is not guaranteed but name containing fused, mul, sum, and add is the useful signal.

In the example, the generated reduction kernel conceptually performs:

x = tl.load(...)
scaled = x * 2.0
reduced = tl.sum(scaled, axis=1)
out = reduced + 1.0
tl.store(..., out)

The pointwise multiplication before the reduction and the addition after it are represented inside one reduction-oriented kernel.

Benchmark fusion without lying to yourself

GPU benchmarking is easy to get wrong.

CUDA work is asynchronous and Python timer can measure how quickly the CPU enqueues work instead of how long the GPU takes to finish it.

Compilation also makes the first invocation intentionally expensive because the graph must be captured, lowered, compiled, and cached.

Both synchronization and warm-up are essential concerns.

Create benchmark_fusion.py:

import statistics
import torch

def pointwise_example( x: torch.Tensor,
  w: torch.Tensor,
  b: torch.Tensor,) -> torch.Tensor:
  return torch.sigmoid(x * w + b)

@torch.inference_mode()
def benchmark_cuda( fn,
  *args: torch.Tensor,
  warmup: int = 25,
  iterations: int = 100,) -> tuple[float, float]:
  # Warm-up triggers compilation for compiled functions and allows
  # runtime libraries/caches to reach a steady state.
  for _ in range(warmup):
      fn(*args)
  torch.cuda.synchronize()
  samples_ms: list[float] = []
  for _ in range(iterations):
      start = torch.cuda.Event(enable_timing=True)
      end = torch.cuda.Event(enable_timing=True)
      start.record()
      fn(*args)
      end.record()
      end.synchronize()
      samples_ms.append(start.elapsed_time(end))
  return statistics.median(samples_ms), statistics.mean(samples_ms)

def main() -> None:
  if not torch.cuda.is_available():
      raise SystemExit("CUDA is required for this benchmark.")
  torch.manual_seed(0)
  # Large enough to expose memory traffic, but adjust for your GPU.
  shape = (4096, 4096)
  x = torch.randn(shape, device="cuda")
  w = torch.randn(shape, device="cuda")
  b = torch.randn(shape, device="cuda")
  compiled = torch.compile(pointwise_example)
  eager_out = pointwise_example(x, w, b)
  compiled_out = compiled(x, w, b)  # compilation happens here
  torch.testing.assert_close(
      compiled_out, eager_out, rtol=1e-4, atol=1e-5
  )
  eager_median, eager_mean = benchmark_cuda(pointwise_example, x, w, b)
  compiled_median, compiled_mean = benchmark_cuda(compiled, x, w, b)
  print(f"Eager median:    {eager_median:.3f} ms")
  print(f"Compiled median: {compiled_median:.3f} ms")
  print(f"Eager mean:      {eager_mean:.3f} ms")
  print(f"Compiled mean:   {compiled_mean:.3f} ms")
  print(f"Median speedup:  {eager_median / compiled_median:.2f}x")

if __name__ == "__main__":
  main()

Run it:

python benchmark_fusion.py

Do not publish the resulting speedup without recording at least:

• GPU model
• PyTorch version
• CUDA version
• tensor shape
• tensor dtype
• compile mode
• warm-up count
• measurement method

For a long-lived inference worker, steady-state speed can dominate but for a short-lived job that handles one request and exits, compilation may make total latency worse.

The fusion patterns you will see in real models

Pointwise fusion is the easiest pattern to understand, but Inductor’s useful optimization surface is broader.

• Pointwise vertical fusion combines dependent elementwise operations, as in scale -> bias -> activation.
• Reduction fusion combines pointwise work around sum, mean, or max, as in normalization, scoring, and loss fragments.
• GEMM epilogue fusion combines work that follows a matrix multiplication, as in matmul -> bias -> activation.
• Prologue fusion combines preprocessing while inputs are loaded, such as a cast, scale, or normalize step ahead of heavier compute.
• Horizontal fusion combines independent operations that share inputs or launch structure, such as parallel tensor-list updates and related transforms.

Reduction fusion

A reduction combines many values into fewer values.

Normalization, pooling, score aggregation, and loss calculations use reductions heavily.

The earlier example:

result = (x * 2.0).sum(dim=-1) + 1.0

contains a pointwise operation before sum and one after it.

Fusing these stages avoids separate buffers and launches around the reduction.

GEMM plus epilogue

Large language and vision models spend substantial time in matrix multiplication.

The matrix multiply itself is compute-heavy, but the result is often immediately followed by bias, activation, scaling, or residual arithmetic.

Conceptually:

out = x @ weight
out = out + bias
out = torch.relu(out)

An epilogue fusion attaches compatible operations to the end of the matrix computation, avoiding a round trip where the matrix result is stored and then reloaded just to apply simple math.

Prologue fusion

A prologue moves compatible preprocessing into the data-loading side of a heavier operation. That can include scaling, casting, or normalization-like transformations that do not need a standalone tensor.

Horizontal fusion

Horizontal fusion combines independent work rather than a dependency chain. PyTorch now also documents explicit horizontal fusion with [foreach_map](https://docs.pytorch.org/tutorials/recipes/foreach_map.html), although that API is described as a prototype and should be treated accordingly. The broader principle is the same: expose enough related work to the compiler that it can amortize launches and data access.

Where this helps in agentic AI systems

torch.compile does not accelerate an HTTP request to a hosted LLM. It does not make a database query faster. It does not reduce tool latency. It helps where your agent executes PyTorch graphs.

1. Embedding services

Local embedding models often run at high request volume with relatively small batches. They contain the same normalization, projection, activation, and pooling patterns as larger networks. Fusion can improve GPU utilization when launch overhead and memory traffic are significant relative to compute.

2. Rerankers

A retrieval agent may rerank dozens or hundreds of candidates. Cross-encoders can become a major portion of end-to-end latency, especially when the generation model is fast or cached.

Compile the reranker as its own measured service boundary. Keep input shapes controlled where possible, warm it before accepting traffic, and compare tail latency.

3. Guardrail and policy models

Safety classifiers, intent routers, and policy networks are often small. Small does not automatically mean efficient. These models can be launch-bound at low batch sizes. The right question is whether compilation improves the real batching strategy used by the service.

4. Local model inference

When an agent uses a local language or multimodal model, compilation can optimize compatible graph regions around attention, normalization, projections, logits processing, and other tensor work. Do not assume the whole model becomes one kernel. Real models mix generated kernels with specialized libraries and backend-specific implementations. The win comes from improving the overall execution plan.

5. Training agent components

Preference models, reward models, routers, and domain-specific encoders may be fine-tuned in-house. Compilation can cover forward and backward regions, but training introduces more shape, autograd, optimizer, and memory considerations. Begin with a representative training step, validate gradients and loss curves, and benchmark after warm-up.

Compiler-friendly engineering matters more than the decorator

Adding torch.compile is one line. Getting durable production value is an engineering workflow.

Find graph breaks early

Graph breaks divide a program into smaller captured regions. Smaller graphs mean fewer opportunities for fusion and scheduling across operation boundaries. PyTorch recommends using [fullgraph=True](https://docs.pytorch.org/docs/stable/user_guide/torch_compiler/compile/programming_model.fullgraph_true.html) to identify and eliminate graph breaks:

compiled = torch.compile(model, fullgraph=True)

This is especially useful during development because it fails when the function cannot be captured as one graph instead of quietly continuing around a break.

Do not blindly enable fullgraph=True in production and call the work finished. Use it as a diagnostic pressure test.

Some dynamic Python behavior genuinely needs restructuring or a smaller compilation boundary.

Useful logging commands include:

TORCH_LOGS="graph_breaks" python app.py
TORCH_LOGS="recompiles" python app.py
TORCH_LOGS="guards" python app.py
TORCH_LOGS="output_code" python app.py

For larger programs, PyTorch’s [TORCH_TRACE](https://docs.pytorch.org/docs/stable/user_guide/torch_compiler/compile/programming_model.observability.html) and [tlparse](https://docs.pytorch.org/docs/stable/user_guide/torch_compiler/compile/programming_model.observability.html) workflow can produce a navigable compilation report:

TORCH_TRACE="/tmp/agent-trace" python app.py
pip install tlparse
tlparse /tmp/agent-trace --latest

Control shape variability

Agent traffic is messy. Prompt lengths, candidate counts, image sizes, and batch sizes vary.

Compilers use guards to determine when a previously compiled graph remains valid. New shapes or behaviors can trigger recompilation. Recompilation is not free, and uncontrolled variability can erase steady-state gains.

At service boundaries, consider bucketing sequence lengths, limiting supported image sizes, or compiling stable repeated regions rather than an entire highly dynamic pipeline.

Separate cold-start and steady-state metrics

Track at least two measurements:

1. Time to first compiled result.
2. Steady-state latency after compilation and warm-up.

For autoscaled workers, cold start can affect user experience. For persistent workers, it may be amortized over thousands of requests.

Keep eager mode as the reference implementation

The eager function is your correctness oracle and fallback path. For every compiled region:

eager = model(inputs)
compiled = compiled_model(inputs)
torch.testing.assert_close(compiled, eager)

Use representative dtypes and shapes. For training, compare gradients and short-run optimization behavior as well as outputs.

Optimize the request path, not an isolated kernel

A fused kernel may be faster while the service remains unchanged because the actual bottleneck is tokenization, host-to-device transfer, queueing, retrieval, or a remote tool call. Profile the full request. Then zoom into PyTorch regions that are both expensive and compilable.

Choosing a compile mode

The default torch.compile mode is a good starting point. PyTorch also exposes modes such as reduce-overhead and max-autotune in the [torch.compile](https://docs.pytorch.org/docs/stable/generated/torch.compile.html) API:

compiled_default = torch.compile(model)
compiled_low_overhead = torch.compile(model, mode="reduce-overhead")
compiled_tuned = torch.compile(model, mode="max-autotune")

Use them as hypotheses, not magic settings.

• reduce-overhead can be useful for small batches where Python and launch overhead matter, and may use CUDA graphs when applicable.
• max-autotune spends more compile-time effort evaluating implementation choices to improve runtime performance.
• The default mode often gives a better cold-start/steady-state balance for general use.

Benchmark all candidate modes on the hardware, shapes, dtypes, and concurrency level that your service actually uses.

Concluding Thoughts

Agentic AI has encouraged developers to think in workflows: one model calls another model, which calls a tool, which updates memory, which triggers another decision. GPU compilers apply the same idea at a lower level.

Instead of treating every tensor operation as an isolated action, Inductor sees a graph.

It asks which actions can be scheduled together, which intermediate states do not need to become durable tensors, and which launches can disappear.

That shift from individual operations to execution graphs is why fusion is so important.

Your Python code can remain expressive:

return torch.sigmoid(x * w + b)

while the runtime becomes much less literal about executing every operation independently.

The practical takeaway is straightforward:

Do not rewrite standard PyTorch math as custom CUDA before checking what the compiler already generates.

In many cases, the compiler can remove the most obvious waste without forcing the team to own low-level GPU code.

And in an agent system where one request may touch several local models, those avoided launches and memory round trips can compound across the entire loop.