Forem: Ingero Team

Agent + MCP + eBPF: 10,869 CUDA Kernel Events, Now Queryable

Ingero Team — Tue, 21 Apr 2026 13:30:00 +0000

A vLLM inference server handles hundreds of requests per second. Then one request with n_completions=8 and logprobs=20 arrives, and every other request blocks for 9-11 seconds. GPU utilization monitors stay green. Kubernetes reports healthy pods. Latency dashboards show a spike but no why. An eBPF trace of every CUDA call is the only view that catches this.

Originally published at ingero.io.

This is a real issue (vLLM #37343). We reproduced it on an RTX 4090 running vLLM 0.18.0 with Qwen3.5 with 27B params. Ingero's eBPF trace captured everything: 10,869 events, 550 cudaLaunchKernel calls, 7,757 context switches, 6 causal chains. All stored in a 1.2 MB SQLite database.

Then we pointed Claude at it via MCP.

The session

The entire investigation below is real. These are actual MCP tool calls, actual data, actual conclusions. Nothing was edited except formatting for readability.

Claude's first move - it called get_trace_stats:

10,869 events traced:
  7,757 sched_switch (context switches)
    550 cudaLaunchKernel   p50: 20us  p99: 5,079us
  1,672 cuLaunchKernel     p50: 15us  p99: 892us
    734 cudaMemcpyAsync
     10 cudaDeviceSync     p50: 19us  p99: 4,720us (239x median)

Two numbers jump out: cudaLaunchKernel p99 is 254x the p50, and there are 7,757 context switches. That ratio is not normal.

Claude's second move - it called get_causal_chains:

The most severe chain:

cudaLaunchKernel p99 = 13.1ms (784x the p50 of 16.7us), correlated with 428 context switches totaling 8.9 seconds off-CPU.

This is the signal. The kernel launch itself takes microseconds. But the CPU thread that feeds the GPU is getting preempted 428 times, accumulating 8.9 seconds of idle time. The GPU sits waiting.

Claude's third move - without being asked, it called get_stacks to look at the actual code paths:

Kernel	Calls	Avg Latency	Max Latency
fused_add_rms_norm	118	54 us	3.0 ms
reshape_and_cache_flash	75	42 us	161 us
silu_and_mul	69	44 us	156 us
rotary_embedding	62	38 us	222 us

All running under VLLM::EngineCore, the single-threaded engine coroutine.

Claude's diagnosis (47 seconds from first tool call):

"The vLLM engine coroutine is running on a CPU shared with other processes. When those processes get scheduled, the coroutine is preempted, the GPU gets no new work, and throughput collapses. The CPU thread driving the GPU was context-switched out 5,347 times, causing 18.4 seconds of CPU unavailability. The root cause is CPU contention, not GPU contention."

Claude's fix recommendation:

Pin the engine thread to dedicated CPU cores (taskset)
Reduce priority of background processes (nice -n 19)
Do NOT use CUDA_LAUNCH_BLOCKING (it will make things worse)
Do NOT trust nvidia-smi (it shows utilization, not causation)

What happened here

No SSH. No log files. No dashboard hopping. No "let me check nvidia-smi on each node."

An AI agent made 4 MCP tool calls against a 1.2 MB SQLite database containing kernel-level eBPF traces. It identified the root cause (CPU scheduling contention), the specific code path (EngineCore coroutine), and the fix (CPU pinning) - all in under a minute.

The key insight: nvidia-smi would have shown 100% GPU utilization during this entire incident. The GPU was "utilized" - it was executing the work it was given. The problem was that it wasn't being given work fast enough because the CPU thread feeding it was being preempted. That distinction - between "GPU is busy" and "GPU is being fed work efficiently" - is invisible to every standard GPU monitoring tool.

What made this possible

This is not a wrapper around nvidia-smi. The eBPF trace attaches uprobes directly to libcudart.so (CUDA Runtime) and libcuda.so (CUDA Driver), plus tracepoints on the Linux kernel scheduler (sched_switch, sched_wakeup), memory allocator (mm_page_alloc), and I/O subsystem. Every CUDA API call is captured with nanosecond precision. Every context switch that preempted a GPU-feeding thread is recorded. The causal chain engine connects them automatically.

The MCP server exposes this data through 10 tools. The AI agent decides what to query. There is no pre-aggregation layer, no dashboard, no human selecting which metrics to look at. The agent gets the raw events and builds the diagnosis.

Try the eBPF trace yourself

The trace database is in the Ingero repo. The investigation works with any MCP-compatible AI:

# 1. Clone and build
git clone https://github.com/ingero-io/ingero.git
cd ingero && make build

# 2. With Claude Code
claude --mcp-config <(echo '{"mcpServers":{"ingero":{"command":"./bin/ingero","args":["mcp","--db","investigations/vllm-37343-logprobs-amplification.db"]}}}')

# 3. With Ollama (any open model)
pip install mcp-client-for-ollama
ollmcp -m qwen3.5:27b -j /tmp/ingero-mcp.json

Type /investigate to start the guided workflow. The AI will walk through the same investigation you just read.

The pattern repeats

This is not a one-off. We have traced dozens of GPU performance issues. The pattern is consistent:

124x slower PyTorch DataLoader - kernel tracing revealed 191,000 context switches and 299,000 page allocations in 40 seconds. The GPU was starved because DataLoader workers were fighting for CPU cores.
13x PyTorch slowdown from hidden NumPy sync - a tensor.cpu().numpy() call in a masking function triggered B x 2 implicit cudaStreamSynchronize calls per forward pass. On faster GPUs, the bottleneck got worse, not better.
GPU 97% utilized but training 3x slower - nvidia-smi reported healthy utilization while Prometheus node exporter and Fluent Bit were consuming 51.7% of available CPU time through 14,504 context switches.

Every one of these follows the same pattern: the GPU is fast, the host is the bottleneck, and standard GPU metrics cannot see it. The causal chain from host event to CUDA API call is the missing link.

What this means for GPU debugging

The traditional approach: alert fires, SSH into the machine, check nvidia-smi, check dmesg, check logs, open profiler, wait for reproduction, analyze flame graphs, correlate across tools. Hours.

The MCP-native approach: point an AI agent at the kernel traces, let it query what it needs, read the diagnosis. Minutes.

We are not saying the AI is smarter than a senior SRE. We are saying it has access to data the SRE cannot see (kernel scheduling decisions, per-CUDA-call latency distributions, automated causal chains) and it can query that data faster than a human can navigate dashboards.

The investigation databases are open source. The agent is open source. Try it locally.

Ingero - open-source eBPF agent for GPU debugging. One binary, zero deps, <2% overhead. Apache 2.0 + GPL-2.0. *GitHub** star - Open an issue if you are seeing vLLM or CUDA runtime issues. Investigation DB: investigations/vllm-cuda-kernel-events.db - Original issue: vllm-project/vllm#37343.*

11-Second Time to First Token on a Healthy vLLM Server

Ingero Team — Tue, 21 Apr 2026 13:30:00 +0000

TL;DR

A vLLM health endpoint says "ok." nvidia-smi says 95% utilization. But a user just waited 11 seconds for their first token. We reproduced a real vLLM issue on an RTX 4090 and traced every CUDA API call and Linux kernel event to find the root causes: head-of-line blocking during prefix caching. This is invisible to standard monitoring. The trace databases are available in the Ingero repo for independent investigation. We traced a production case of vLLM latency spikes down to kernel-level scheduling contention.

The Problem Nobody Can See

vLLM's continuous batching is one of the best things to happen to LLM serving. It lets the engine process multiple requests simultaneously, filling GPU capacity that would otherwise sit idle between sequential requests.

But continuous batching has a dark side: when requests compete for GPU resources inside the same batch, one expensive request can silently starve all others. No error. No health check failure. No metric spike. Just users waiting 10x-250x longer than expected for their first token.

We investigated a real vLLM issue reported in the last week (#37308) to understand what happens at the kernel level during these silent latency spikes.

Setup

The investigation used the same server configuration:

python -m vllm.entrypoints.openai.api_server \
 --model Qwen/Qwen2.5-0.5B-Instruct \
 --port 8000 \
 --gpu-memory-utilization 0.95 \
 --max-model-len 32768 \
 --enable-prefix-caching

Hardware: RTX 4090 (24GB), 4 vCPUs, Ubuntu 22.04, vLLM 0.17.1.

We ran Ingero alongside each test to trace CUDA Runtime/Driver API calls and host kernel events (scheduler context switches, memory allocations) simultaneously.

Prefix Caching Head-of-Line Blocking

Issue: vllm-project/vllm#37308

What happens

6 concurrent requests arrive within 40ms. 4 are heavy (2048-token prompts, 128-512 output tokens) and 2 are light (128-token prompts, 32-64 output tokens). All share a 32-token prefix so the prefix cache groups them together.

The light requests should complete in under 100ms. Instead:

Run	r08 (128 tok)	r05 (128 tok)	r07 (2048 tok)	r02 (2048 tok)
1	1,131ms	1,406ms	1,654ms	1,851ms
2	54ms	129ms	258ms	234ms
3	66ms	177ms	175ms	156ms

Run 1 is catastrophic: the light requests are 14x over threshold. Subsequent runs settle to 2-4x because the prefix cache warms up. But that first cold-cache batch is brutal.

What the tracer shows

3 causal chains detected. The most revealing one:

[MEDIUM] cudaLaunchKernel p99=444us (6.4x p50) - 371 sched_switch events
 Timeline:
 [HOST ] 371 context switches (5.9s off-CPU)
 [CUDA ] p99=444us (6.4x p50=70us)

The per-process breakdown tells the full story:

VLLM::EngineCore (the GPU scheduling loop):

24,347 context switches, max stall 2.5 seconds
40,632 cuLaunchKernel calls, avg 29us but max 34ms
34,087 cudaLaunchKernel calls, avg 96us but max 356ms

The engine core process – the single-threaded loop that decides which requests get GPU time – was descheduled for 2.5 seconds in the worst case. During that stall, the GPU kernel queue drained and the light requests had nothing submitted on their behalf.

The 356ms cudaLaunchKernel spike (3,700x the average) is the smoking gun. That's not the GPU being slow. That's the CPU failing to submit work to the GPU because the scheduling loop was preempted.

Why nvidia-smi misses this

nvidia-smi shows high utilization because the GPU IS working – on the heavy requests' prefills. The light requests are starving, but from the GPU's perspective there's always a kernel to run. The starvation is in the CPU-side scheduling loop, not on the GPU.

What Standard Tools Show vs What Kernel Tracing Shows

Signal	nvidia-smi	vLLM /health	vLLM metrics	Kernel tracing
GPU utilization	95%+	–	–	95%+ (but wrong work)
Server health	–	"ok"	requests_running=5	–
TTFT regression	–	–	Visible in histograms	Visible + root cause
Engine stall (2.5s)	Not visible	Not visible	Not visible	24,347 sched_switch events
Kernel launch drop (80%)	Not visible	Not visible	Not visible	1,051 -> 208 ops/s
Memory pressure	Not visible	Not visible	Not visible	43,606 mm_page_alloc
Which process is blocked	Not visible	Not visible	Not visible	VLLM::EngineCore PID 2438

The key insight: GPU utilization was high because the GPU was doing work. It was just doing the wrong work – processing heavy prefills or computation while light requests starved. No GPU-side metric can distinguish "GPU is busy computing my request" from "GPU is busy computing someone else's request while mine waits."

Implications for Production vLLM

If you're running vLLM in production with mixed workloads (different prompt sizes, some requests with or ), you're likely experiencing these silent regressions:

Monitor TTFT per-request, not just aggregate throughput. Aggregate metrics hide the tail – your p99 might be 100x worse than p50 during batch contention.
Be careful with . A single request with n=8 and =20 can block your entire server for 11+ seconds on a cold cache. Consider routing these to dedicated instances.
First-request-after-idle is the worst case. This issue showed the most extreme regression on Run 1 (cold prefix cache). If your traffic is bursty, the first batch after a quiet period will hit hardest.
GPU utilization is not a proxy for request health. Your dashboards might show 95% utilization while individual users experience 256x TTFT regression.

Investigate It Yourself

The trace database from this investigations are in the Ingero repo:

git clone https://github.com/ingero-io/ingero.git
cd ingero && make build

# View the causal chains
./bin/ingero explain --db investigations/vllm---amplification.db --since 5m

# Per-process breakdown
./bin/ingero explain --db investigations/vllm---amplification.db --per-process --since 5m

# Connect your AI assistant for interactive investigation
./bin/ingero mcp --db investigations/vllm---amplification.db

Investigate with AI (recommended)

You can point any MCP-compatible AI client at the trace database and ask questions directly. No code required.

First, create the MCP config file at /tmp/ingero-mcp-vllm.json:

{
  "mcpServers": {
    "ingero": {
      "command": "./bin/ingero",
      "args": ["mcp", "--db", "investigations/vllm-37308-hol-blocking.db"]
    }
  }
}

With Ollama (local & free: no data sent outside):

# Install ollmcp (MCP client for Ollama)
pip install ollmcp

# Investigate with a local model (no data leaves your machine)
ollmcp -m qwen3.5:27b -j /tmp/ingero-mcp-vllm.json

With Claude Code (with data sent to remote models / Anthropic):

claude --mcp-config /tmp/ingero-mcp-vllm.json

Then type /investigate and let the model explore. Follow up with questions like "what was the root cause?" or "which kernel calls had the highest latency spikes?"

Ask your AI assistant: "What caused the 80% throughput drop?" or "Which process had the most context switches?" The trace data has the full story.

The investigation database from this post is available for download. Investigations performed on TensorDock RTX 4090 (24GB), Ubuntu 22.04, vLLM 0.17.1, Qwen/Qwen2.5-0.5B-Instruct with prefix caching enabled.

GitHub (give us a star!): github.com/ingero-io/ingero. No NVIDIA SDK, no code changes, production-safe by design.

If you are seeing vLLM issues in your own workloads, we'd love to take a look. Drop an issue on GitHub and we will gladly dive into it together.

Ingero is free & open source software licensed under Apache 2.0 (user-space) + GPL-2.0/BSD-3 (eBPF kernel-space). One binary, zero dependencies, <2% overhead.

What Happens When an AI Agent Gets Kernel-Level GPU Traces

Ingero Team — Thu, 16 Apr 2026 16:48:26 +0000

TL;DR

A GPU trace of a PyTorch DataLoader bottleneck (114x slower than direct indexing) was loaded into an MCP server and handed to Claude for investigation. The AI identified the root cause in under 30 seconds: 3,676 CPU context switches starving the GPU of data. Below is the full investigation session with the trace database available for independent reproduction. We walked through a real case of Claude MCP GPU debugging, from raw eBPF traces to root cause identification.

The Idea

GPU performance debugging usually goes like this: training is slow, nvidia-smi shows nothing useful, print statements get added, hours pass. What happens when raw trace data gets handed to an AI assistant with the question “what went wrong?”

That’s what the MCP server enables. The tracer traces CUDA API calls and Linux kernel events, stores them in a SQLite database, then exposes them to AI assistants via the Model Context Protocol (MCP). The AI can query the data, read causal chains, inspect per-process breakdowns, and run custom SQL through natural conversation.

We tested this on a real investigation: a PyTorch DataLoader bottleneck where DataLoader was 114x slower than direct tensor indexing on an RTX 4090. The trace database is in the Ingero repo for independent reproduction.

Setup (2 minutes)

1. Get the trace database

git clone https://github.com/ingero-io/ingero.git
cd ingero && make build

The investigation database is at investigations/pytorch-dataloader-starvation.db (764KB).

2. Start the MCP server

./bin/ingero mcp --db investigations/pytorch-dataloader-starvation.db

This starts the MCP server in stdio mode. To connect Claude Desktop, add this to your MCP config (~/.config/claude/claude_desktop_config.json):

{
  "mcpServers": {
    "ingero": {
      "command": "/path/to/ingero",
      "args": ["mcp", "--db", "/path/to/investigations/pytorch-dataloader-starvation.db"]
    }
  }
}

For Claude Code (CLI), the MCP server is automatically available when running from the repo directory.

The Investigation Session

Here’s the actual conversation. We started with a simple question and let the AI dig deeper.

“What happened in this trace?”

The AI calls get_trace_stats and gets back:

3,807 events over 77 seconds
CUDA Runtime: cudaMemcpyAsync (14 calls, avg 31.5ms, max 65ms)
              cudaLaunchKernel (16 calls, avg 1.7ms, max 25.8ms)
              cudaStreamSync (21 calls, avg 31us)
              cudaDeviceSync (3 calls, avg 17.2ms)
Host: sched_switch (3,676 calls, avg 401us, max 23.9ms)
      process_fork (11 calls)
      process_exit (5 calls)
Driver: cuLaunchKernel (41 calls, avg 40us)
        cuMemAlloc (10 calls, avg 307us)

Right away, two things stand out:

cudaMemcpyAsync averages 31.5ms per call. That’s host-to-device data transfer. It should be microseconds for the batch sizes involved.

3,676 sched_switch events in 77 seconds. That’s the Linux scheduler preempting processes, something is fighting for CPU.

“What are the causal chains?”

The AI calls get_causal_chains:

40 causal chains detected (all MEDIUM severity)

Recurring pattern:
  cudaLaunchKernel p99=25.8ms (352x p50=73us) - N sched_switch events
  Context switches escalate from 34 -> 106 -> 268 -> 308 -> 319 over time

The chains show a clear escalation: context switches grow from 34 to 319 over the trace window. The DataLoader workers are creating more and more scheduling pressure as the benchmark runs.

“Which processes are involved?”

The AI calls run_sql:

SELECT pn.name, COUNT(*) as events
FROM events e
JOIN process_names pn ON e.pid = pn.pid
GROUP BY pn.name
ORDER BY events DESC

python3          3,799 events
pt_data_worker     (visible in fork/exit events)
pt_data_pin        (visible in fork/exit events)
systemd-journal       17 events

The main python3 process owns almost all events. But the DataLoader workers (pt_data_worker, pt_data_pin) are visible through fork and scheduling events, they’re being created, fighting for CPU, and getting descheduled.

“Why is cudaMemcpyAsync so slow?”

The AI connects the dots:

The 14 cudaMemcpyAsync calls average 31.5ms (max 65ms). These are host-to-device data transfers that should complete in microseconds for batch sizes under 100MB.

The 3,676 sched_switch events show the main process being preempted 47 times per second. Each preemption delays the data pipeline. The DataLoader workers (pt_data_worker) are competing with the main process for CPU, on a 4-vCPU machine, 4 workers + the main process + system processes means constant context switching.

The GPU is receiving cudaMemcpyAsync calls, but the source data isn’t ready in host memory because the DataLoader worker that was supposed to prepare it got descheduled. The GPU waits.

Root cause: CPU scheduling contention between DataLoader workers and the main training process. The GPU is fast, it’s the host that can’t keep up.

“What would you recommend?”

Reduce num_workers to 2 (on a 4-core machine, 4 workers causes over-subscription)

For in-memory datasets, skip DataLoader entirely, direct tensor indexing is 114x faster because it avoids the multiprocessing overhead

If DataLoader is required, use persistent_workers=True to avoid fork overhead on each epoch

Total investigation time: under 30 seconds from first question to root cause + fix.

What the AI Has Access To

The MCP server exposes 7 tools:

Tool	What It Does
`get_check`	System diagnostics (kernel, GPU, CUDA, driver)
`get_trace_stats`	CUDA + host statistics (p50/p95/p99 per operation)
`get_causal_chains`	Automated root cause chains with severity ranking
`get_stacks`	Resolved call stacks (symbols + Python source lines)
`run_demo`	Run synthetic demo scenarios
`get_test_report`	GPU integration test results
`run_sql`	Read-only SQL against the trace database

The run_sql tool is the most powerful: the AI can write arbitrary queries against the event table, joining with process names, ops, and sources.

Try It Yourself

The trace database from this investigation is in the repo:

git clone https://github.com/ingero-io/ingero.git
cd ingero && make build

# Quick analysis (no MCP needed)
./bin/ingero explain --db investigations/pytorch-dataloader-starvation.db --since 5m

# Interactive AI investigation via MCP
./bin/ingero mcp --db investigations/pytorch-dataloader-starvation.db

With Claude Desktop

Add to ~/.config/claude/claude_desktop_config.json:

{
  "mcpServers": {
    "ingero": {
      "command": "./bin/ingero",
      "args": ["mcp", "--db", "investigations/pytorch-dataloader-starvation.db"]
    }
  }
}

Then ask Claude: “What caused the GPU performance problem in this trace?”

With Any MCP Client

The MCP server works with any MCP-compatible client: Cursor, Windsurf, or custom implementations. The stdio transport is universal.

Investigate with AI (recommended)

# With Ollama (local, free)
pip install ollmcp
ollmcp -m qwen3.5:27b -j /tmp/ingero-mcp-dataloader.json

# With Claude Code
claude --mcp-config /tmp/ingero-mcp-dataloader.json

Type /investigate and let the model explore.

Why This Matters

Traditional GPU debugging is manual: run nvidia-smi, add print statements, read logs, guess. The AI-assisted approach is different:

The tracer captures everything at the kernel level: CUDA API calls, host scheduling, memory events, with zero code changes
The trace database is self-contained: no need to reproduce the issue, no need for the original hardware
The AI asks the right follow-up questions: it sees the context switches, connects them to CUDA latency, and identifies the root cause pattern

This turns GPU debugging from “spend hours staring at logs” into “ask a question, get an answer.”

Investigation DB: investigations/pytorch-dataloader-starvation.db

Original issue: pytorch/pytorch#154318

GitHub: github.com/ingero-io/ingero. No NVIDIA SDK, no code changes, production-safe by design.

Ingero is free & open source software licensed under Apache 2.0 (user-space) + GPL-2.0/BSD-3 (eBPF kernel-space). One binary, zero dependencies, <2% overhead.

MCP as Observability Interface: Connecting AI Agents to Kernel Tracepoints

Ingero Team — Thu, 16 Apr 2026 07:35:33 +0000

TL;DR

MCP is becoming the interface between AI agents and infrastructure
data. Datadog shipped an MCP
Server connecting dashboards to AI agents.
Qualys flagged MCP servers as the new shadow IT risk.
We think both are right, and we think the architecture should
go further: the MCP server should not wrap an existing observability
platform. It should BE the observability layer. This post explores how
MCP can serve as a direct observability interface to kernel
tracepoints, bypassing traditional metric pipelines entirely.

Three signals in one week

Three things happened in the same week of March 2026 that signal where
observability is headed.

Datadog shipped an MCP Server
Their implementation connects real-time observability data to AI agents for automated detection and remediation. An AI agent can now query Datadog dashboards, pull metrics, and trigger responses through the Model Context Protocol. This is a big company validating a small protocol.

Qualys published a security analysis of MCP
servers.
Their TotalAI team called MCP servers “the new shadow IT for AI” and
found that over 53% of servers rely on static secrets for
authentication. They recommended adding observability to MCP servers:
logging capability discovery events, monitoring invocation patterns,
alerting on anomalies.

Cloud Native Now covered eBPF for Kubernetes network observability.
Microsoft Retina deploys as a DaemonSet, captures network telemetry via
eBPF without application changes, and provides kernel-level drop reasons. The article draws a clear line between “monitoring” (predefined questions) and “observability” (asking questions nobody planned for).

The thread connecting all three: AI agents need direct access to
infrastructure telemetry, and MCP is becoming the way they get it.

Two approaches to MCP observability

There are two ways to connect observability data to AI agents via MCP.

Approach 1: Wrap existing platforms. Datadog’s strategy. Take
existing metrics, logs, and traces, already collected and aggregated,
and expose them through MCP tools. The AI agent queries the dashboard
API, gets pre-processed data, and acts on it. This makes sense for teams
with a mature observability stack that want to add AI-powered automation
on top.

Approach 2: Build MCP-native observability. This is what we did with
the tracer. Instead of wrapping an existing platform, we built an eBPF
agent that traces CUDA Runtime and Driver APIs via uprobes, stores the
results in SQLite, and exposes everything through 7 MCP tools. The MCP
interface is not an adapter layer; it is the primary interface.

Neither approach is wrong. They solve different problems.

The wrapper approach works well for aggregate analysis: “What was the
p99 latency for service X over the last hour?” The data is already
summarized, indexed, and queryable.

The native approach works better for root-cause investigation: “Why did
this specific GPU request take 14.5x longer than expected?” That
requires raw kernel events, CUDA call stacks, and causal chains – not
summaries. The AI agent needs to drill down, not roll up.

What MCP-native observability looks like in practice

Here is a concrete example. We traced a vLLM TTFT regression where the
first token took 14.5x longer than baseline. The trace database captured
every CUDA API call, every kernel context switch, every memory
allocation.

When Claude connects to the MCP server and loads this database, it can:

get_trace_stats – See the full trace summary: 12,847 CUDA events, 4 causal chains, total GPU time
get_causal_chains – Read the causal chains that explain why latency spiked, in plain English
run_sql – Run custom queries against the raw event data (“show me all cudaMemcpyAsync calls over 100ms”)
get_stacks – Inspect call stacks for any flagged event

Claude identified the root cause in under 30 seconds: logprobs
computation was blocking the decode loop, creating a 256x slowdown on
the critical path. That root cause was not visible in any aggregate
metric. It only appeared in the raw causal chain between specific CUDA
API calls.

A dashboard MCP adapter could not have found this. The data granularity
does not survive aggregation.

The security angle matters too

Qualys raised valid concerns about MCP server security. Their finding
that 53% of servers rely on static secrets is alarming. Their
recommendation to log discovery and invocation events is exactly right.

For MCP servers that touch GPU infrastructure, the attack surface is
different. An MCP server with access to CUDA traces can expose timing
information, memory layouts, and model architecture details. The
security model needs to account for this.

In Ingero, the MCP server runs inside the same process as the eBPF tracing pipeline. There is no separate data layer between the AI agent and the kernel-level telemetry - the MCP tools query the same event store that the eBPF probes write to. This is why Ingero can answer causal questions in real time: the AI agent has direct access to raw kernel and CUDA events, not a pre-aggregated summary.

Try It Yourself

The project is open source. The investigation database from this post is available for download. Claude (or any MCP client) can connect to it and run an investigation:

git clone https://github.com/ingero-io/ingero.git
cd ingero && make build
./bin/ingero mcp --db investigations/pytorch-dataloader-starvation.db

Investigate with AI (recommended)

You can point any MCP-compatible AI client at the trace database and ask questions directly. No code required.

First, create the MCP config file at /tmp/ingero-mcp-dataloader.json:

{
  "mcpServers": {
    "ingero": {
      "command": "./bin/ingero",
      "args": ["mcp", "--db", "investigations/pytorch-dataloader-starvation.db"]
    }
  }
}

With Ollama (local, free):

# Install ollmcp (MCP client for Ollama)
pip install ollmcp

# Investigate with a local model (no data leaves your machine)
ollmcp -m qwen3.5:27b -j /tmp/ingero-mcp-dataloader.json

With Claude Code:

claude --mcp-config /tmp/ingero-mcp-dataloader.json

Then type /investigate and let the model explore. Follow up with questions like “what was the root cause?” or “which processes were competing for CPU time?”

The MCP server exposes 7 tools. Claude will figure out the rest.

Ingero is free & open source software licensed under Apache 2.0 (user-space) + GPL-2.0/BSD-3 (eBPF kernel-space). One binary, zero dependencies, <2% overhead. Give us a star at GitHub!

One Query, Four GPUs: Tracing a Distributed Training Stall Across Nodes

Ingero Team — Mon, 13 Apr 2026 17:18:24 +0000

TL;DR

A single straggling node held up a 4-node distributed training job. We found it by fanning out one SQL query to all four nodes and getting the answer in under a second. This is distributed GPU training debugging with eBPF – no central service, no Prometheus, no time-series database, just the same single-binary agent already running on each machine.

The problem we kept hitting

We’ve been building Ingero – an eBPF agent that traces CUDA API calls and host kernel events to explain GPU latency. Until v0.9, it was single-node only. Trace one machine, explain what happened on that machine. For single-GPU inference or training, that worked well.

But distributed training spreads the debugging surface across machines. When a 4-node DDP job slows down, the question is always: which node? And then: why? nvidia-smi on each machine reports healthy utilization. dstat shows nothing obvious. The typical workflow is SSH-ing into each box, eyeballing logs, diffing timestamps across terminals, and hoping the issue is still happening.

We wanted cross-node investigation without adding infrastructure. The question was: what’s the simplest architecture that works?

What we shipped in v0.9.1

Three features, all built on top of the existing per-node agent. No new services, no new daemons, no new ports.

1. Node identity

Every event now carries a node tag. The agent stamps each event with a name from a --node flag, an ingero.yaml config value, or the hostname as fallback:

sudo ingero trace --node gpu-node-01

Event IDs become node-namespaced (gpu-node-01:4821) so databases from different nodes can merge without collisions. For torchrun workloads, rank and world size are auto-detected from environment variables (RANK, LOCAL_RANK, WORLD_SIZE) – no extra configuration needed.

2. Fleet fan-out queries

Each Ingero agent already exposes a dashboard API over HTTPS (TLS 1.3, auto-generated ECDSA P-256 cert if no custom cert is provided). The new fleet client sends the same query to every node in parallel, collects the results, and concatenates them with a node column prepended. For production clusters, the client supports mTLS – --ca-cert, --client-cert, --client-key – so both sides authenticate. Plain HTTP is available via --no-tls but requires an explicit opt-in, and even then it’s intended for trusted VPC networks only.

The --nodes flag works for ad-hoc queries, but for anything beyond a handful of nodes, the node list goes into ingero.yaml once and every command picks it up automatically:

fleet:
  nodes:
    - gpu-node-01:8080
    - gpu-node-02:8080
    - gpu-node-03:8080
    - gpu-node-04:8080

A full example config is in configs/ingero.yaml.

Here’s what it looked like when we ran it against a 4-node cluster where one node was misbehaving:

$ ingero query --nodes gpu-node-01:8080,gpu-node-02:8080,gpu-node-03:8080,gpu-node-04:8080 \
    "SELECT node, source, count(*) as cnt, avg(duration)/1000 as avg_us
     FROM events GROUP BY node, source"

node              source  cnt    avg_us
----------------  ------  -----  ------
gpu-node-01       4       11009  5.2
gpu-node-01       3       847    18400  # ← 9x higher than peers
gpu-node-02       4       10892  5.1
gpu-node-02       3       412    2100
gpu-node-03       4       10847  5.3
gpu-node-03       3       398    1900
gpu-node-04       4       10901  5.0
gpu-node-04       3       421    2200

  8 rows from 4 node(s)

Node 1 jumps out immediately: 847 host events at 18.4ms average, while the other three sit around 2ms. One more command to see the causal chains:

$ ingero explain --nodes gpu-node-01:8080,gpu-node-02:8080,gpu-node-03:8080,gpu-node-04:8080

FLEET CAUSAL CHAINS - 2 chain(s) from 4 node(s)

[HIGH] [gpu-node-01] cuLaunchKernel p99=843us (63.9x p50) - 847 sched_switch events + heavy block I/O
  Root cause: 847 sched_switch events + heavy block I/O
  Fix: Pin training process to dedicated cores with taskset; Add nice -n 19 to background jobs

[MEDIUM] [gpu-node-01] cuMemAlloc p99=932us (5.0x p50) - 855 sched_switch events + heavy block I/O
  Root cause: 855 sched_switch events + heavy block I/O
  Fix: Pin training process to dedicated cores with taskset

Both chains are on gpu-node-01. The other three nodes have zero issues. The root cause: CPU contention from block I/O – checkpoint writes preempting the training process.

Two commands to go from “distributed training is slow” to “pin the training process on node 1 and investigate the I/O source.”

3. Offline merge and Perfetto export

Not every environment allows live HTTP queries between nodes. Air-gapped clusters, locked-down VPCs, compliance constraints – there are real reasons the network path isn’t always available.

For those cases, ingero merge combines SQLite databases from each node into a single queryable file:

# 1. Collect traces from each node
scp gpu-node-01:~/.ingero/ingero.db node-01.db
scp gpu-node-02:~/.ingero/ingero.db node-02.db

# 2. Merge and analyze
ingero merge node-01.db node-02.db -o cluster.db
ingero explain -d cluster.db

Stack traces are deduplicated by hash. Events keep their node-namespaced IDs. Old databases that predate the node column work with --force-node.

For visual timeline analysis, ingero export --format perfetto produces a Chrome Trace Event Format JSON that opens in ui.perfetto.dev. Each node gets its own process track. Causal chains show up as severity-colored markers. The straggler is visible at a glance in the timeline.

Why we built it this way

The obvious approach to multi-node observability is a central collector: ship events to a time-series database, build dashboards, set up alerts. Prometheus, Datadog, Honeycomb – the well-trodden path.

We deliberately avoided that.

No new infrastructure. Ingero is a zero-config, single-binary agent with no dependencies. Adding a central collector contradicts that. The fleet client is 400 lines of Go in the existing binary. It reuses the HTTPS API the agent already exposes. Nothing new to deploy, nothing new to secure – the same TLS 1.3 + mTLS configuration that protects a single node’s dashboard protects the entire fleet.

Client-side fan-out is simple and sufficient. The CLI sends concurrent HTTP requests, collects results, and merges them locally. A sync.WaitGroup, some JSON decoding, column concatenation. No distributed query planning, no consensus protocol, no coordinator election. For 4-50 nodes, this is the right level of complexity.

Partial failure is first-class. If one node is unreachable, results from the others still come back, plus a warning. No all-or-nothing semantics. In practice, the unreachable node is often the one in trouble – and knowing which nodes failed is diagnostic information in itself.

Clock skew is measured, not ignored. eBPF timestamps come from bpf_ktime_get_ns() (CLOCK_MONOTONIC), which is per-machine. When correlating events across nodes, clock differences matter. The fleet client runs NTP-style offset estimation in parallel with the actual query – 3 samples per node, median filter. On a typical LAN with sub-millisecond RTT, precision should be well under 10ms. If skew exceeds a threshold, it warns. This adds zero latency since it runs concurrently with the data query.

Offline merge covers air-gapped environments. Some production GPU clusters have no internal HTTP connectivity between nodes. SCP the databases, merge locally, investigate. The merge path also serves as a permanent record of the cluster state at investigation time.

MCP: AI-driven fleet investigation

The fleet is also accessible through Ingero’s MCP server via the query_fleet tool. Here’s what the raw tool output looks like for a chains query across the same 4-node cluster:

query_fleet(action="chains", since="5m")

Fleet Chains: 2 chain(s)
[HIGH] gpu-node-01 | cuLaunchKernel p99=843us (63.9x p50) | 847 sched_switch events + heavy block I/O
[MEDIUM] gpu-node-01 | cuMemAlloc p99=932us (5.0x p50) | 855 sched_switch events + heavy block I/O

That’s the complete response – an AI assistant gets this back from one tool call, no SSH access to each node, no manual SQL. The tool supports four actions: chains (causal analysis), sql (arbitrary queries), ops (operation breakdown per node), and overview (event counts). Clock skew warnings are prepended automatically when detected.

Where this stands

v0.9.1 is the initial step in cluster-level tracing, not the destination.

What we have now works well for the reactive investigation workflow: something went wrong, we need to find out what and where. Fan-out queries, offline merge, Perfetto export – these are diagnostic tools for after the fact.

We’re actively working on cross-node correlation and straggler detection – more updates coming soon. And since the instrumentation sits on host-level eBPF rather than vendor-specific hooks, none of this is limited to a specific GPU vendor.

The bet is that client-side fan-out scales to 50+ nodes before anything centralized is needed. When it doesn’t, the node-namespaced ID scheme and offline merge path ensure the architecture can evolve without breaking existing deployments.

We’re stress-testing the fan-out architecture against larger clusters and would welcome feedback from teams running multi-node training. Open an issue on GitHub.

The investigations/ directory has ready-to-query databases for trying this without a GPU cluster:

sample-gpu-node-01.db, sample-gpu-node-02.db, sample-gpu-node-03.db – individual node traces from a 3-node cluster
sample-cluster.db – all three merged into one (600 events, 6 chains, 9 stacks)

GitHub (give us a star!): github.com/ingero-io/ingero. No NVIDIA SDK, no code changes, production-safe by design.

If you are facing distributed training issues in your own workloads, we’d love to take a look. Drop an issue on GitHub and we will gladly dive into it together.

Ingero is free & open source software licensed under Apache 2.0 (user-space) + GPL-2.0/BSD-3 (eBPF kernel-space). One binary, zero dependencies, <2% overhead.

CUDA Graphs: The 8-Year Overnight Success and the Observability Gap

Ingero Team — Wed, 08 Apr 2026 13:40:15 +0000

CUDA Graphs: The 8-Year Overnight Success and the Observability Gap

TL;DR

CUDA graphs shipped in 2018 but only became critical infrastructure in the past two years, driven by LLM inference demands and framework automation. They also create an observability blind spot: hundreds of kernel launches collapse into one opaque cudaGraphLaunch. A March 2025 study found 25% of CUDA graphs in PyTorch workloads actually degrade performance. We traced graph lifecycle events (capture, instantiate, launch) via eBPF uprobes and correlated them with CPU scheduling and I/O pressure to detect graph pool exhaustion, re-capture storms, and CPU contention during dispatch. The investigation database and reproduction steps are included below.

CUDA graphs shipped in 2018. For five years, almost nobody used them. Today, they power every token generated by vLLM, SGLang, and TensorRT-LLM. The technology didn’t change. The world did.

We’ve been investigating what changed, why it matters, and what the CUDA graphs observability gap looks like in practice: the industry’s most widely adopted GPU optimization hides hundreds of kernel launches behind a single API call, and existing tools can’t see into it.

CUDA Graphs: Quick Context

Every CUDA kernel launch costs 20-200 microseconds of CPU-side work: Python interpreter overhead, framework dispatch, driver processing, hardware submission. CUDA graphs record a sequence of GPU operations into a DAG, instantiate it once, and replay it with a single API call. The CPU tax is paid once for the entire graph instead of per kernel.

The mechanism has been available since CUDA 10. The interesting question is why adoption stayed flat for five years and then went vertical.

Five Forces: Why Now?

1. GPUs Got Too Fast for CPUs to Keep Up

GPU FP16 throughput grew 47x from Pascal (GP100, 21.2 TFLOPS) to Hopper (H100, ~1,000 TFLOPS with Tensor Cores). Kernel execution times collapsed from milliseconds to single-digit microseconds. CPU-side launch overhead stayed at 20-140µs per operation (higher in Python frameworks).

The PyGraph team measured a segment of DALL-E 2 inference that launches 740+ kernels with a combined GPU time of 3.4ms. End-to-end latency: 14ms. 75% of the wall-clock time, the GPU sits idle, waiting for the CPU to submit the next operation.

CUDA graphs collapse all of that per-kernel overhead into a single graph launch, as low as ~2.5µs since CUDA 12.6.

2. The Workload That Needs Graphs Most Didn’t Exist Yet

LLM autoregressive decode (generating one token at a time with fixed compute shapes and small batch sizes) is the perfect CUDA graph workload. Static shapes, repeated execution, CPU overhead dominance.

This workload category barely existed before ChatGPT launched in November 2022. By 2026, inference is projected to account for roughly two-thirds of all AI compute spending (Deloitte TMT Predictions). The economic pressure to optimize it became enormous.

3. Frameworks Made It a Flag Flip

Before PyTorch 2.x, using CUDA graphs meant writing manual capture code in C++/CUDA: stream capture semantics, fixed memory addresses, graph instantiation. Expert territory.

torch.compile(mode="reduce-overhead") lowered that barrier to one line of Python. vLLM and SGLang built graph capture directly into their serving pipelines. Adoption shifted from “CUDA expert” to “set a flag.”

4. NVIDIA Made the API Actually Usable

The original CUDA graph API was rigid: no conditional logic, no dynamic control flow. Real workloads with variable batch sizes or branching paths couldn’t use graphs without ugly workarounds.

NVIDIA shipped a steady stream of fixes:

CUDA 12.4 (2024): Conditional nodes (IF, WHILE)
CUDA 12.6: Constant-time graph launch (~2.5µs + ~1ns/node)
CUDA 12.8: IF/ELSE, SWITCH nodes; Blackwell support
Nsight Compute 2025.3+: CUDA Graph Viewer and graph-aware profiling

By 12.8, the API covered the majority of real-world control flow patterns.

5. The Economics Made It Non-Optional

At the scale of billions of inference requests per day, a 2.3x throughput improvement from CUDA graphs (measured on LLaMA-2 7B by Fireworks AI) translates to cutting the GPU fleet, or the cloud bill, nearly in half. That’s millions of dollars.

Within two years, CUDA graphs moved from optional optimization to baseline infrastructure for inference serving.

The Observability Gap

The performance story is well-documented. The observability story is not.

CUDA graphs create an observability black hole.

When a graph launches, hundreds of individual kernel launches, memory copies, and synchronization points collapse into a single cudaGraphLaunch call. From any external observer (profiler, monitoring, eBPF probes) there is one event where there used to be hundreds.

In our investigation work, we’ve found several ways this gap manifests in production:

Graphs can silently hurt performance. A March 2025 paper from the PyGraph project found that 25% of CUDA graphs (29 of 116 analyzed) in PyTorch workloads actually degraded performance, with individual graphs reaching up to ~5x slowdown (397% degradation). The culprits: parameter copy overhead eating up to 24% of execution time, memory garbage collection after replay, RNG state resets. Without graph-aware tracing, these costs are invisible.

Graph re-capture is expensive and hard to detect. (NVIDIA best practices, PyGraph 2025) When a new batch size arrives that doesn’t match any pre-captured graph, the framework re-captures. That’s a costly operation that blocks inference. In vLLM, this can cause latency spikes that cascade to all co-scheduled requests. From standard monitoring, it looks like a random latency blip.

CPU contention during graph dispatch is invisible. (NVIDIA CUDA Graphs blog, PyGraph 2025) A CUDA graph launch is fast (~2.5µs on CUDA 12.6+), but only if the CPU thread gets to run uninterrupted. If logrotate, a DataLoader worker, or a noisy neighbor preempts the thread during dispatch, the graph launch stalls. nvidia-smi sees nothing. The GPU utilization dashboard stays green. But your p99 latency just spiked.

Graph pool exhaustion has no standard alert. (NVIDIA memory troubleshooting, CUDA pool API) When the pool of instantiated graphs fills up, launch rates drop. In one of our traces, graph launch rate dropped from 163 to 2 launches/second, a 99% collapse, with no warning from any standard monitoring tool. The root cause: a batch size change triggered re-capture, and the pool couldn’t keep up.

The irony is precise: the optimization that eliminates per-kernel CPU overhead also eliminates per-kernel visibility.

Tracing Graph Lifecycle Events with eBPF

To close this observability gap, we added CUDA graph lifecycle tracing to Ingero (open-source, eBPF-based GPU observability) in v0.9.0. The approach uses eBPF uprobes on the CUDA runtime, which means no CUPTI dependency, no Nsight session, and no application code changes.

Live Graph Tracing

ingero trace attaches uprobes to cudaStreamBeginCapture, cudaStreamEndCapture, cudaGraphInstantiate, and cudaGraphLaunch, alongside the standard CUDA API, driver API, and host kernel events.

Above: a torch.compile inference workload traced live. The CUDA Runtime table shows graphLaunch (985 launches, p50=9.9µs), graphBeginCapture (2 captures), graphEndCapture (2 completions), right alongside cudaLaunchKernel, cudaStreamSync, and the rest of the standard CUDA operations. Host context shows 50,000+ sched_switch events with CPU at 100%.

The bottom of the display fires real-time anomaly correlation: “cudaStreamSync p99=1.4ms (171.7x p50), correlated with 46,507 sched_switch events.” And the graph-specific finding: “[MEDIUM] CPU contention delaying graph dispatch (985 launches, 184,619 sched_switch).”

This kind of cross-layer correlation (CUDA API latency tied to host scheduler pressure) is what we found missing from existing GPU profiling workflows.

Causal Chain Diagnosis

After tracing, ingero explain reads the recorded events from SQLite and assembles causal chains: cross-layer correlations that explain why something went wrong, not just what.

The incident report finds 8 causal chains (6 HIGH, 2 MEDIUM). The graph-specific findings:

[MEDIUM] CPU contention delaying graph dispatch (12 launches, 2,251 sched_switch)

  Fix: Pin the inference process to dedicated CPU cores;
       reduce background CPU load during inference.

[MEDIUM] Graph launch rate dropped 99% (exec 0x0, PID 1789)
  Rate dropped from 163 to 2 launches/sec

  Root cause: graph pool exhaustion, likely re-capture triggered by new batch size.

  Fix: Pre-warm all expected batch sizes during model startup;
       set max_num_batched_tokens to limit batch size variability.

The causal chain links four layers: system context (CPU 100%) to host events (124,241 context switches, 7.9s off-CPU) to CUDA API latency spikes (cudaDeviceSync p99=20.8ms, 436x normal) to graph dispatch stalls. One chain, one root cause.

AI-Assisted Investigation

Ingero also exposes an MCP (Model Context Protocol) server. Any MCP-compatible AI assistant (Claude Code, Cursor, local models via Ollama) can query a trace database directly.

Above: Claude Code runs broad diagnostics first (get_check, get_trace_stats, get_causal_chains in parallel), then drills into graph_lifecycle, graph_frequency, and get_stacks. Then the agent correlates the results across all layers and produces a root cause analysis with fix recommendations, without manual SQL or log parsing.

Try It Yourself

Two ways to reproduce or extend this investigation.

Path 1: Instant Investigation (No GPU Needed)

Download the pre-captured investigation database and query it immediately:

# Download the CUDA graph investigation database
curl -fsSL -o cuda-graph.db \
  https://raw.githubusercontent.com/ingero-io/ingero/main/investigations/cuda-graph-cpu-contention.db

# View causal chains
ingero explain --db cuda-graph.db

This database contains a real torch.compile inference workload under CPU contention: graph captures, instantiations, 985 launches, pool exhaustion, and the full causal chain explaining the 99% launch rate drop.

Investigate with AI

# With Claude Code:
claude mcp add -s local ingero -- ingero mcp --db cuda-graph.db

claude
# Ask: "Use ingero tools to investigate this CUDA graph trace"

# Install ollmcp (MCP client for Ollama):
pip install ollmcp

# Or with Ollama:
cat > /tmp/ingero-mcp.json << 'EOF'
{
  "mcpServers": {
    "ingero": {
      "command": "ingero",
      "args": ["mcp", "--db", "cuda-graph.db"]
    }
  }
}
EOF

# Local model (no data leaves your machine)
ollmcp -m qwen3.5:27b -j /tmp/ingero-mcp.json

# Or use a cloud-hosted model via Ollama (faster, data sent to provider)
ollmcp -m minimax-m2.7:cloud -j /tmp/ingero-mcp.json

Path 2: Full End-to-End (Any NVIDIA GPU + Linux)

Reproduce the entire investigation from scratch:

# Install Ingero
VERSION=0.9.1
curl -fsSL "https://github.com/ingero-io/ingero/releases/download/v${VERSION}/ingero_${VERSION}_linux_amd64.tar.gz" | tar xz
sudo mv ingero /usr/local/bin/

# Run the CUDA graph demo workload (requires PyTorch 2.x)
python tests/workloads/cuda_graph_demo.py &

# Add CPU contention to trigger the interesting behavior
stress-ng --cpu 2 --timeout 30s &

# Trace
sudo ingero trace --pid $(pgrep -f cuda_graph_demo) --db demo.db --duration 30s

# Investigate
ingero explain --db demo.db

# Or let AI investigate
claude mcp add -s local ingero -- sudo ingero mcp --db demo.db

This reproduces the same graph lifecycle events, causal chains, and root cause analysis on your own hardware and workload.

What’s Next

CUDA graphs aren’t going away. Every new NVIDIA toolkit release makes them more capable: conditional nodes, device-side launch, tighter framework integration. The workloads that depend on them (LLM inference, diffusion models, real-time serving) are only growing.

Observability for these workloads needs to keep pace. If you’re running torch.compile or serving models with vLLM, your GPU workload is already using CUDA graphs. The community needs better tooling to see what they’re doing under the hood. Furthermore, the eBPF architecture we are building at the Linux kernel level sets the foundation for tracking these same host-side bottlenecks across heterogeneous hardware in the future.

The investigation database from this post is available for download.
Investigation performed on EC2 g4dn.xlarge (Tesla T4), Ubuntu 24.04, kernel 6.17, NVIDIA 580.126.09, PyTorch 2.10+CUDA 12.0. Also validated on RTX 4090, A100, H100, and GH200.

GitHub (give us a star!): github.com/ingero-io/ingero. No NVIDIA SDK, no code changes, production-safe by design.

Interested in CUDA Graphs tracing? Drop us a message at info(@)ingero.io or create an issue on GitHub and we will gladly dive into it together.

Ingero is free & open source software licensed under Apache 2.0 (user-space) + GPL-2.0/BSD-3 (eBPF kernel-space). One binary, zero dependencies, <2% overhead.

124x Slower: What PyTorch DataLoader Actually Does at the Kernel Level

Ingero Team — Wed, 01 Apr 2026 16:40:55 +0000

TL;DR: PyTorch's DataLoader can be 50-124x slower than direct tensor indexing for in-memory GPU workloads. We reproduced a real PyTorch issue on an RTX 4090 and traced every CUDA API call and Linux kernel event to find the root cause. The GPU wasn't slow - it was starving. DataLoader workers generated 200,000 CPU context switches and 300,000 page allocations in 40 seconds, leaving the GPU waiting an average of 301ms per data transfer that should take microseconds.

The Problem

A PyTorch user reported that DataLoader was 7-22x slower than direct tensor indexing for a simple MLP inference workload. Even with num_workers=12, pin_memory=True, and prefetch_factor=12, the gap remained massive. GPU utilization sat at 10-20%.

We reproduced it. The gap was even worse on our hardware:

Method	Time	vs Direct
Direct tensor indexing	0.39s	1x
DataLoader (shuffle=True)	48.49s	124x slower
DataLoader (optimized, 4 workers, pin_memory)	43.29s	111x slower

The workload is trivial: 7M samples, 100 features, 2-layer MLP, batch size 1M. The model processes a batch in milliseconds. So where does the time go?

What nvidia-smi Shows

Nothing useful. GPU utilization flickers between 0% and 30%. Memory usage is stable. Temperature is fine. The GPU is clearly underutilized, but nvidia-smi can't explain why.

What torch.profiler Shows

The reporter tried PyTorch's built-in profiler and "obtained no meaningful trace data." This is a common frustration - application-level profilers can show what CUDA kernels are running, but they cannot see the host-side scheduling, memory, and process lifecycle events that determine whether data arrives at the GPU on time.

What Kernel-Level Tracing Shows

We ran the benchmark while tracing both CUDA API calls (via eBPF uprobes on libcudart.so) and Linux kernel events (scheduler context switches, memory page allocations, process forks) simultaneously. The results tell the complete story.

4 HIGH-severity causal chains

The causal chain engine detected 4 high-severity patterns, all with the same root cause:

[HIGH] cudaStreamSync p99=42ms (1,638x p50=25us) - CPU 100% + 1,880 sched_switch events
Timeline:
  [SYSTEM] CPU 100%
  [HOST  ] 1,880 context switches (21s off-CPU)
  [CUDA  ] p99=42ms (1,638x p50=25us)
Root cause: DataLoader workers fighting for CPU, massive page allocation pressure

[HIGH] cudaLaunchKernel p99=24.67ms (349x p50=70us) - CPU 100%
Root: 34 sched_switch events

[HIGH] cuMemAlloc p99=627us (4.0x p50) - CPU 100%

[HIGH] cuLaunchKernel p99=106us (4.0x p50) - CPU 100%

The cudaStreamSync p99 is 1,638 times the p50. That's not GPU slowness - that's the GPU waiting for data that never arrives on time.

The Per-Process Breakdown

This is where it gets clear. The main process and its 4 DataLoader workers are visible as separate entities:

Main process:

cudaMemcpyAsync (host-to-device transfer): avg 301ms, max 2.9 seconds
cudaStreamSync: p99 = 42ms (normally 25us)
1,567 context switches, avg 16ms off-CPU, worst stall 5 seconds
799,018 page allocations

DataLoader worker 1: 52,863 context switches, 89,338 page allocations, worst stall 5s
DataLoader worker 2: 50,638 context switches, 83,509 page allocations, worst stall 5s
DataLoader worker 3: 49,361 context switches, 70,035 page allocations, worst stall 5s
DataLoader worker 4: 38,862 context switches, 56,354 page allocations, worst stall 5s

Total across workers: ~191,000 context switches and ~299,000 page allocations in 40 seconds.

What This Means

The DataLoader workers are doing three expensive things that direct indexing avoids entirely:

Shuffling and indexing. DataLoader with shuffle=True generates a random permutation of indices, then each worker selects its chunk. This requires random memory access across the full 7M-sample tensor - terrible for cache locality and triggers page faults.
Collation and copying. Each worker gathers scattered samples into a contiguous batch tensor. This means allocating new memory (page allocations), copying data from random locations (cache misses), and serializing the result back to the main process via shared memory or a queue.
Competing for CPU. Four workers + the main process on a 4-vCPU machine means constant preemption. Each worker gets descheduled 50,000 times. The worst-case stall is 5 seconds - during which the GPU has nothing to process.

With direct indexing: X[i:i+batch_size] is a zero-copy view of a contiguous tensor already in memory. .to(device) triggers one DMA transfer from a single contiguous region. No workers, no shuffling, no collation, no cross-process copies, no context switches. The GPU gets data in microseconds, not hundreds of milliseconds.

The Fix

For in-memory GPU workloads where the entire dataset fits in RAM:

1. Don't use DataLoader. Direct indexing with a pre-shuffled index array is simpler and 100x faster:

indices = torch.randperm(num_samples)
for i in range(0, num_samples, batch_size):
    batch = X[indices[i:i+batch_size]].to(device)
    output = model(batch)

2. When DataLoader is necessary, match num_workers to the actual CPU core count minus 1. On a 4-core machine, num_workers=2 reduces contention. Add persistent_workers=True to avoid fork overhead.

3. For larger-than-memory datasets where DataLoader is necessary, the real bottleneck shifts to disk I/O. Use prefetch_factor=2 (not higher - more prefetching means more memory pressure) and ensure the storage subsystem can keep up.

The Bigger Picture

This investigation illustrates a pattern we see constantly in GPU workloads: the GPU is fast, the host is the bottleneck, and GPU metrics can't see it. nvidia-smi reported low utilization but couldn't explain why. torch.profiler captured CUDA kernels but missed the 200,000 context switches happening in userspace.

The only way to see the full picture was to trace both sides simultaneously - CUDA API calls at the library level and Linux kernel scheduling events - and correlate them by time and process ID.

The causal chain CPU 100% -> 1,880 sched_switch -> cudaMemcpyAsync 301ms -> cudaStreamSync 42ms tells the complete story in one line. Without cross-stack tracing, this would have remained a mystery - as it was for the original reporter who spent weeks debugging it.

Try It Yourself

Reproduce the benchmark:

import torch, time
from torch.utils.data import DataLoader

X = torch.randn(7_000_000, 100)
model = torch.nn.Sequential(
    torch.nn.Linear(100, 512),
    torch.nn.ReLU(),
    torch.nn.Linear(512, 512),
    torch.nn.ReLU(),
    torch.nn.Linear(512, 10)
).cuda()

# Fast path
start = time.time()
with torch.no_grad():
    for i in range(0, len(X), 1_048_576):
        model(X[i:i+1_048_576].cuda())
    torch.cuda.synchronize()
print(f"Direct: {time.time()-start:.3f}s")

# Slow path
loader = DataLoader(X, batch_size=1_048_576, shuffle=True)
start = time.time()
with torch.no_grad():
    for batch in loader:
        model(batch.cuda())
    torch.cuda.synchronize()
print(f"DataLoader: {time.time()-start:.3f}s")

Trace with Ingero to see what's happening under the hood:

git clone https://github.com/ingero-io/ingero.git
cd ingero && make build
sudo ./bin/ingero trace --duration 60s   # in one terminal
python3 benchmark.py                      # in another terminal
./bin/ingero explain --since 60s          # after benchmark completes

GitHub: github.com/ingero-io/ingero

Original issue: pytorch/pytorch#154318

Investigation performed on TensorDock RTX 4090 (24GB), Ubuntu 22.04, PyTorch 2.10.0+cu128.

Tracing a 13x PyTorch Slowdown to a Hidden NumPy Synchronization

Ingero Team — Tue, 31 Mar 2026 14:53:32 +0000

TL;DR: A .cpu().numpy() call buried inside a forward pass was forcing a full CPU-GPU synchronization on every batch, every loop iteration. The GPU would finish its work in milliseconds, then sit idle for ~2 seconds waiting for Python and NumPy to catch up. Replacing the NumPy logic with pure PyTorch ops gave a 6.4x speedup on a T4 and 13x on an RTX 5080. The fix is two lines of code.

The bug

Swin-MAE is a masked autoencoder built on Swin Transformers. A user training on 5.2 million images with an RTX 5090 noticed that GPU utilization kept dropping to ~30% during the forward pass. The model would spike, stall, spike, stall.

The problem was in window_masking, the function that decides which image patches to mask during training. Here is what the code looked like:

# The hot loop: runs once per batch, every forward pass
for i in range(B):
    index_mask[i] = np.setdiff1d(index_all, index_keep.cpu().numpy()[i])

for i in range(B):
    x_masked[i, index_mask.cpu().numpy()[i, :], :] = mask_token

Two Python for loops. Two .cpu().numpy() calls per iteration. On a batch size of B, that is B x 2 implicit cudaStreamSynchronize calls per forward pass.

Why this kills performance

Every time you call .cpu().numpy() on a CUDA tensor, PyTorch has to:

Flush the GPU pipeline. Any queued CUDA operations must finish before the data can be read back.
Transfer data over PCIe. The tensor moves from GPU VRAM to system RAM.
Block the Python thread. Nothing else happens until the transfer completes.

On older GPUs, this penalty was small enough to go unnoticed. The Swin-MAE maintainer confirmed the code was written four years ago. But on modern hardware (RTX 5080, 5090), the GPU finishes its batch computation in milliseconds. The NumPy detour takes ~2 seconds. The faster the GPU gets, the worse this anti-pattern becomes.

The np.setdiff1d call is doing set-difference math on the CPU to figure out which patches to mask. This is something PyTorch can do entirely on the GPU without ever leaving CUDA.

The fix

Drop NumPy entirely. Keep everything as PyTorch tensor operations:

# Before: B x 2 implicit cudaStreamSync per forward pass
for i in range(B):
    index_mask[i] = np.setdiff1d(index_all, index_keep.cpu().numpy()[i])

for i in range(B):
    x_masked[i, index_mask.cpu().numpy()[i, :], :] = mask_token

# After: zero CPU transfers
index_mask = ids_shuffle[:, -mask_len:]
x_masked.scatter_(1, index_mask.unsqueeze(-1).expand(-1, -1, C), mask_token)

The key insight: ids_shuffle already contains the full permutation on the GPU. The masked indices are just the tail end of that shuffle. No need to compute a set difference at all, and scatter_ handles the masked token assignment without leaving CUDA.

Results

Hardware	Before	After	Speedup
T4 (AWS EC2)	baseline	6.4x faster	6.4x
RTX 5080	baseline	13x faster	13x

The RTX 5080 benefits more because its raw compute is faster, which makes the CPU round-trip penalty proportionally larger.

This pattern is everywhere

The .cpu().numpy() anti-pattern is not unique to Swin-MAE. We see it constantly in training code, especially in:

Custom masking logic (MAE, BEiT, any masked pretraining)
Dynamic batching and padding (NLP sequence collation)
Custom augmentation pipelines that mix PyTorch and NumPy/SciPy
Metric computation mid-training (accuracy checks that pull tensors to CPU every N steps)

The tell is always the same: GPU utilization that spikes and drops in a sawtooth pattern. The GPU is not slow. It is waiting.

How to find this deterministically

The manual way to find this bug is to grep the codebase for .cpu() and .numpy() and hope the culprit is actually inside the hot loop. The slightly better way is to run a standard profiler, stare at a wall of timelines, and try to manually correlate host-side Python threads with GPU stream synchronizations.

To catch these implicit syncs deterministically without code changes, we take it to the kernel level. Ingero, an open-source eBPF tracer, attaches uprobes directly to libcudart.so and libcuda.so. Instead of polling metrics, it captures every CUDA API call with nanosecond precision and builds causal chains connecting host OS events directly to GPU stalls.

Investigate with AI

After running ingero trace, you can point any MCP-compatible AI client at the resulting trace database and ask questions directly. No code required.

Create the MCP config file at /tmp/ingero-mcp-pytorch.json:

{
  "mcpServers": {
    "ingero": {
      "command": "./bin/ingero",
      "args": ["mcp", "--db", "ingero-trace.db"]
    }
  }
}

Replace ingero-trace.db with the path to the trace database created by ingero trace.

With Ollama (local, free):

ollmcp -m qwen3.5:27b -j /tmp/ingero-mcp-pytorch.json

With Claude Code:

claude --mcp-config /tmp/ingero-mcp-pytorch.json

Then type /investigate and let the model explore.

The takeaway

If you are mixing NumPy and PyTorch inside a training loop, you are probably paying a synchronization tax on every batch. Modern GPUs are fast enough that a single .cpu().numpy() call can dominate your total training time.

Check your forward pass. Check your masking logic. Check your custom collation functions. If the GPU is waiting, the fix might be two lines.

GitHub: github.com/ingero-io/ingero
Original issue: Zian-Xu/Swin-MAE#24

The code is open source (Apache 2.0). Star it on GitHub.

Forem: Ingero Team

Agent + MCP + eBPF: 10,869 CUDA Kernel Events, Now Queryable

The session

What happened here

What made this possible

Try the eBPF trace yourself

The pattern repeats

What this means for GPU debugging

11-Second Time to First Token on a Healthy vLLM Server

TL;DR

The Problem Nobody Can See

Setup

Prefix Caching Head-of-Line Blocking

What happens

What the tracer shows

Why nvidia-smi misses this

What Standard Tools Show vs What Kernel Tracing Shows

Implications for Production vLLM

Investigate It Yourself

Investigate with AI (recommended)

Related reading

What Happens When an AI Agent Gets Kernel-Level GPU Traces

TL;DR

The Idea

Setup (2 minutes)

1. Get the trace database

2. Start the MCP server

The Investigation Session

“What happened in this trace?”

“What are the causal chains?”

“Which processes are involved?”

“Why is cudaMemcpyAsync so slow?”

“What would you recommend?”

What the AI Has Access To

Try It Yourself

With Claude Desktop

With Any MCP Client

Investigate with AI (recommended)

Why This Matters

Related reading

MCP as Observability Interface: Connecting AI Agents to Kernel Tracepoints

TL;DR

Three signals in one week

Two approaches to MCP observability

What MCP-native observability looks like in practice

The security angle matters too

Try It Yourself

Investigate with AI (recommended)

Related reading

One Query, Four GPUs: Tracing a Distributed Training Stall Across Nodes

TL;DR

The problem we kept hitting

What we shipped in v0.9.1

1. Node identity

2. Fleet fan-out queries

3. Offline merge and Perfetto export

Why we built it this way

MCP: AI-driven fleet investigation

Where this stands

Related reading

CUDA Graphs: The 8-Year Overnight Success and the Observability Gap

CUDA Graphs: The 8-Year Overnight Success and the Observability Gap

TL;DR

CUDA Graphs: Quick Context

Five Forces: Why Now?

1. GPUs Got Too Fast for CPUs to Keep Up

2. The Workload That Needs Graphs Most Didn’t Exist Yet

3. Frameworks Made It a Flag Flip

4. NVIDIA Made the API Actually Usable

5. The Economics Made It Non-Optional

The Observability Gap

Tracing Graph Lifecycle Events with eBPF

Live Graph Tracing

Causal Chain Diagnosis

AI-Assisted Investigation

Try It Yourself

Path 1: Instant Investigation (No GPU Needed)

Investigate with AI

Path 2: Full End-to-End (Any NVIDIA GPU + Linux)

What’s Next