Kernel Execution Explanation Toolkit (KEET)

• Kernel Execution Explanation Toolkit (KEET) is an LLM-based framework that interprets Nsight Compute GPU profiles by linking source code with detailed performance metrics.
• It employs a multi-agent architecture to decompose the analysis into source inspection, profile evaluation, and integrated optimization suggestions including optional DrGPU input.
• Empirical evaluations on NVIDIA GPUs demonstrate enhanced diagnostic accuracy and improved optimization speedups, highlighting KEET’s practical impact in performance tuning.

Kernel Execution Explanation Toolkit (KEET) is an LLM-based agentic framework for interpreting Nsight Compute profiles to generate useful and data-grounded natural language explanations of performance issues in GPU kernels, and suggestions for optimizations. It is motivated by the fact that performance profiles of GPU kernels generated by tools such as Nsight Compute are rich in detail but are often challenging to interpret, requiring kernel developers to spend significant time analyzing and comparing profiles in the tool's graphical interface to identify and understand kernel performance bottlenecks. KEET ingests kernel source files, metadata, and one or more .ncu-rep profiles, then orchestrates multiple specialized agents to produce bottleneck diagnoses, optimization suggestions, and aggregated reports (Davis et al., 6 May 2026).

1. Problem setting and scope

KEET is positioned around a specific difficulty in GPU performance engineering: the gap between detailed profiler output and actionable explanation. The framework targets CUDA kernels profiled with Nsight Compute and evaluated on NVIDIA H100 GPUs, with V100 and A100 used for multi-profile ablations. Its stated objective is not merely to summarize counters, but to interpret them in relation to source structure and to generate suggestions for optimization.

The input space is explicitly defined. KEET ingests kernel source files in .cu or .cpp form together with descriptive metadata such as launch parameters and input sizes. It also loads one or more Nsight Compute .ncu-rep files via the Python Report Interface (PRI) to extract per-kernel, per-metric data, including occupancy, stall reasons, and bandwidth counters. Line-level metrics for DrGPU are imported from CSV exports. This workflow ties source-level context to profiler-level evidence rather than treating the profile as an isolated artifact.

The framework is evaluated on several CUDA kernels of varying complexity, and the paper reports that the generated explanations, when provided as context, improve the quality of LLM code optimization and multiple-choice question answering in downstream tasks. A plausible implication is that KEET is intended both as an interpretive interface for human analysis and as a structured context generator for subsequent LLM-based optimization steps (Davis et al., 6 May 2026).

2. Multi-agent architecture and reasoning decomposition

KEET is organized as a multi-agent framework centered around two iterative stages, followed by two final roles that aggregate and review the results.

Component	Agents	Role
Source Code Inspection Stage	Code File Selector; Code Summarizer; Performance Hypothesizer	Build an algorithm summary and propose performance hypotheses
Profile Inspection Stage	Profile Selector; Metric Selector; Profile Analyzer (LLM); DrGPU Evaluator (optional)	Select informative profiles and metrics, generate bottleneck analysis, optionally merge DrGPU suggestions
Final roles	Analysis Aggregator; Explanation Reviewer (optional)	Stitch analyses into a coherent report and mark hypotheses as confirmed / refuted / inconclusive

The data flow is explicit. First, source files and metadata are ingested. Second, Nsight Compute reports are loaded through PRI, and line-level metrics are imported from CSV exports for DrGPU. Third, the Source Code Inspection agents build an “algorithm summary” and propose a small set of performance hypotheses, such as “this kernel may be memory-latency bound.” Fourth, the Profile Inspection agents select profiles and metric subsets that are most informative relative to the current hypotheses, then prompt an LLM to produce a data-grounded bottleneck analysis; DrGPU suggestions can optionally be merged at this stage. Finally, the Analysis Aggregator stitches together per-profile analyses into a coherent report, while the Explanation Reviewer contrasts the final findings against the initial hypotheses and marks each as “confirmed / refuted / inconclusive.”

The framework decomposes the overall reasoning task into dozens of specialized agents, each driven by a small template prompt. Agents are given only the context they need, and they emit structured JSON or Markdown snippets that subsequent agents ingest rather than free-form text. Inline citations such as [Metric: dram__throughput = 1.77%] are enforced by prompting for them explicitly to curb hallucination. This suggests that KEET treats performance explanation as a controlled composition of intermediate artifacts rather than as a single long-form generation problem.

The chaining of reasoning steps is described as: Code Summarizer, Hypothesizer, Profile Selector, Metric Selector, Profile Analyzer, DrGPU Evaluator, Aggregator, and Reviewer. The paper also notes an optional scoring idea in which one could sample multiple candidate analyses and select the highest-scoring one under the LLM’s own log-probabilities:

$\arg\max_i \log P(x_i)$

where $x_1, x_2, \ldots, x_n$ are candidate texts and $P(x)$ is the model’s joint probability of the token sequence $x$. KEET does not currently ensemble multiple outputs in this way, but the framework is described as permitting it (Davis et al., 6 May 2026).

3. Explanation generation and metric-grounded diagnosis

KEET’s explanation generation is centered on the Profile Analyzer. This agent uses a small “performance-analysis guidelines” document, approximately 200 words distilled from NVIDIA docs and experts, together with selected metric values. It maps raw counters to natural-language observations through explicit rules. Examples given in the paper include: if occupancy is below 25%, the explanation becomes “Low SM occupancy: only X out of Y warps were ever active”; if sm__warps_active < 30%, the explanation becomes “Thread-level underutilization: average active threads per warp is only Z”; and if dram__throughput < 5% but L1/L2 hit rates exceed 90%, the explanation becomes “Memory-latency bound: stalls are due to long scoreboard waits rather than raw bandwidth limits” (Davis et al., 6 May 2026).

A prompt excerpt illustrates how these rules are operationalized: “Given the following metrics: {…}, and the guideline that occupancy below 30% often indicates launch-configuration bottlenecks, produce a brief bottleneck diagnosis with inline citations to metric names and values.” The use of inline citations is a central control mechanism, intended to anchor each generated statement in named counters and specific values.

The gaussian Fan2 kernel provides a concrete example of the final “Summary of main bottlenecks.” KEET identifies four main issues: memory-latency bound, not bandwidth bound; very poor global memory coalescing; low warp and thread-level utilization; and algorithmic tail and wasted threads. The associated evidence is numerical and specific: long scoreboard stalls dominate at 21.39 per issue; the L2 hit rate is high while dram__throughput is only 1.77%; global memory coalescing is 8.34 B/sector versus an ideal 32 B/sector, with 374 514 excessive sectors; and blocks have only 16 threads, with average active threads approximately 15.15 (~50%). Fixed grid size is reported to lead to many threads doing no useful work at large pivot indices.

These examples show that KEET does not restrict itself to one category of bottleneck. It can connect stall behavior, cache behavior, memory access efficiency, launch geometry, and algorithmic structure inside a single explanation. A plausible implication is that the framework aims to bridge low-level hardware counter interpretation and higher-level algorithmic diagnosis.

4. Optimization suggestions and interaction with DrGPU

KEET’s Profile Analyzer also emits concrete “Suggested optimizations.” The examples reported are: “Increase block size to 128 threads to raise occupancy from 25% → 75%”; “Refactor the memory layout so that global loads are fully coalesced (32 B/sector)”; and “Use __ldg() to leverage the read-only data cache for large lookup tables.” These suggestions are grounded in simple decision rules rather than unconstrained free-form generation.

The decision rules are stated directly. If occupancy is below a threshold $X\%$, KEET advises adjusting blocks or grids. If there is excessive warp serialization, it advises loop unrolling or algorithmic reordering. If the shared-memory footprint is high, it advises register accumulation or reorganizing data structures. The framework therefore links classes of profiler evidence to classes of optimization actions in a rule-mediated way.

DrGPU enters as an optional auxiliary source of suggestions. When enabled, its leaf-node suggestions—for example, “remove redundant __syncthreads()” or “use warp intrinsics”—are passed through a “DrGPU Evaluator” agent. That agent filters them using a short LLM prompt: “Which of these DrGPU suggestions directly addresses a confirmed bottleneck?” This filtering step matters because KEET’s reviewer stage can mark hypotheses as confirmed, refuted, or inconclusive; the DrGPU Evaluator is therefore not simply appending suggestions, but selecting those aligned with the evidence-backed diagnosis (Davis et al., 6 May 2026).

A common misconception would be that any optimization suggestion attached to a profile is equally justified. The framework’s design argues against that interpretation: suggestions are tied to specific metric patterns, and optional DrGPU guidance is further screened against confirmed bottlenecks.

5. Evaluation protocol, downstream tasks, and quantitative results

The evaluation uses NVIDIA H100 as the primary target hardware, with V100 and A100 for multi-profile ablations. The benchmark suite consists of 9 Rodinia kernels, LULESH, and XSBench, for 11 kernels total. One representative Nsight Compute profile per kernel was collected, while multi-profile ablations used up to 48 profiles. Two downstream tasks are defined. The first is multiple-choice question answering (MCQ), with 20 questions per kernel and score@1 as the metric, described as expected percent correct with a single attempt. The second is code optimization (OPT), in which an LLM generates optimized kernel code that is compiled and tested; the metrics are pass@1 and speedup@1 over the original (Davis et al., 6 May 2026).

The paper gives the following metric definitions:

$$\text{pass@k} = \frac{1}{|T|} \sum_{t \in T} \left[1 - \frac{C(N-c_t, k)}{C(N, k)}\right]$$

$$\text{speedup@k} = \frac{1}{|P|} \sum_{p \in P} \sum_{j=1}^N \left[\frac{C(j-1, k-1)}{C(N, k)}\right]\cdot \left(\frac{T^*_p}{T_{p,j}}\right)$$

The reported results are specific. In MCQ, KEET and KEET+DrGPU both achieve average score@1 ≈ 82%, compared with 79% for the best non-KEET baseline. In OPT, KEET alone yielded the highest maximum speedup in 5/11 cases, while KEET+DrGPU did so in 3/11. Example speedups are: gaussian, 14.0× for KEET versus 12.4× for the baseline; LULESH, 2.87× versus 2.81×; and b+tree, 1.64× versus 1.28×.

The ablation results are also reported in a focused way: Metric Selector and Profile Selector both raise speedup@1 by approximately 10–15%, measured as a harmonic mean across applications. This suggests that selective context construction—choosing which profiles and which metrics to expose to the LLM—is not merely a convenience for context-window management, but an empirically relevant contributor to optimization quality.

6. Multi-profile scaling, limitations, and future work

KEET includes a multi-profile mode that reuses the same agent pipeline. In this setting, the Profile Selector chooses which profile or profiles to analyze next by comparing current hypotheses against the metadata of untouched profiles and seeking maximal “hypothesis coverage.” The Metric Selector picks new metrics across the batch to avoid re-analysis of already-covered data. The Profile Analyzer receives an expanded context of multiple profiles only when necessary, such as when correlating a block-size sweep with occupancy changes (Davis et al., 6 May 2026).

The empirical sweet spot reported for multi-profile input is not monotonic with profile count. Providing approximately 4 profiles for LULESH and approximately 16 for XSBench maximized speedup@1. Beyond that, additional profiles added noise and slightly degraded optimization quality. This directly counters the simple assumption that more profiler data always improves explanation or optimization.

The limitations are stated explicitly. First, hallucination risk remains: despite inline citations, LLMs can still invent explanations for metrics not seen in the profile data, so ongoing work is needed to enforce strict metric grounding. Second, KEET currently recommends block and grid changes but does not search them automatically; future versions could integrate an autotuner loop, with Bayesian optimization given as an example, for knob selection. Third, the present implementation relies on Nsight Compute; porting to AMD ROCm or Intel GPUs would require replacing the ingestion layer while retaining the same agentic architecture. Fourth, very large kernels exceeding 1 000 SLoC may exceed model context windows, motivating more aggressive file chunking or retrieval-augmented prompting. Fifth, KEET’s quality varies with the chosen model; smaller open-source LLMs produce weaker suggestions, although integration with DrGPU can partially compensate.

Taken together, these points characterize KEET as an extensible, interpretable pipeline that turns raw Nsight Compute data and CUDA source code into human-readable performance diagnoses and targeted optimization advice, orchestrated by transparent LLM agents. The reported evaluation on H100 GPUs indicates improved understanding in MCQ tasks and larger code-level speedups in downstream optimizations compared to prior rule-based tools and simpler code+data baselines, while also making clear that grounding, autotuning, portability, context limits, and model dependence remain active constraints (Davis et al., 6 May 2026).