DEV Community: Myoungho Shin

GPU Flight — System Architecture

Myoungho Shin — Mon, 16 Mar 2026 05:37:39 +0000

The previous post was about thread divergence at the SASS level. Before I move on to other optimization strategies, I think it makes sense to first review GPU Flight’s architecture. Understanding the overall structure will make the upcoming topics easier to follow, and it should also be helpful for anyone trying to run the full system locally or in the cloud.

In this post, I want to step back and show the bigger picture: what each component does, why I separated them this way, and what kind of deployment setup makes the most sense in production.

The Big Picture

This is how all the components fit together:

In short:

gpufl-client — runs on the GPU machine, hooks into CUDA activity, and writes structured logs
gpufl-agent — watches those logs and sends them to the backend through HTTP or Kafka
gpufl-backend — receives the events, stores them, and provides APIs for querying
gpufl-front — React UI for browsing sessions and checking GPU behavior

gpufl-client — Capturing GPU Events

The client is a C++ library, or you can use the gpufl Python package. You link it into your application, call gpufl::init() when the program starts, and from there GPU Flight begins collecting GPU-related events automatically.

It can capture things like:

Kernel launches — kernel name, grid/block dimensions, register count, shared memory, occupancy, timing
Memory copies — copy direction, size, duration, throughput
SASS metrics — warp-level per-instruction counters when the SassMetrics engine is enabled
PC samples — warp stall reason distribution when the PcSampling engine is enabled
System metrics — periodic NVML snapshots such as GPU utilization, VRAM usage, temperature, power draw, and fan speed

These are written into three JSONL log files:

<prefix>.device.log   ← kernel launches, memory copies, timing
<prefix>.scope.log    ← GFL_SCOPE blocks with optional SASS/PC data
<prefix>.system.log   ← periodic NVML snapshots

Log rotation is handled automatically. The agent tracks file offsets so events are never re-sent across rotations.

gpufl-agent — The Flexible Shipper

This part is important because it reflects one of the main design choices in GPU Flight.

The client just writes files. It does not know anything about HTTP, Kafka, or where the data is going. That responsibility belongs entirely to the agent.

The agent is a separate process that tails the log files and forwards the events to the backend. Because it is separated from the client, it can run anywhere as long as it can access the log folder.

┌──────────────────────────────┐
│         GPU Machine          │
│                              │
│  gpufl-client → .log files   │
│                    │         │
│             gpufl-agent      │   ← simplest: agent on the same machine
└────────────────────┼─────────┘
                     │  (or NFS / shared volume)
                     ▼
              gpufl-agent           ← or: agent on a separate machine

I wanted to keep profiling logic as isolated as possible. Once the client is mixed with transport logic, the instrumentation library becomes harder to maintain and harder to use in different environments. By splitting them, I can change how data is shipped without changing the application binary.

What the agent does:

Tails all three log streams concurrently
Tracks file positions via a cursor file — survives restarts without re-sending events
Follows log rotation, picking up renamed files automatically
Publishes events via HTTP or Kafka
Optionally archives the compressed raw log files to S3-compatible object storage

That last part is also useful. The raw .log files are the full original event stream. Keeping them in object storage can help later if you want long-term retention, auditability, or reprocess old runs with updated analysis logic.

Configuration examples:

# HTTP — simplest, good for dev or small deployments
gpufl-agent --folder=/var/log/gpuflight \
            --type=http \
            --url=http://backend:8080/api/v1/events/

# Kafka — recommended for production
gpufl-agent --folder=/var/log/gpuflight \
            --type=kafka \
            --brokers=kafka:9092

# Kafka + S3 archival
gpufl-agent --folder=/var/log/gpuflight \
            --type=kafka --brokers=kafka:9092 \
            --archiver-endpoint=https://s3.amazonaws.com \
            --archiver-bucket=gpu-traces \
            --archiver-access-key=KEY \
            --archiver-secret-key=SECRET

Why a message queue in production?

When multiple GPU machines are sending data simultaneously, publishing directly to the backend via HTTP couples the producer rate to the backend's processing speed. If the backend is slow, restarting, or temporarily down, events back up or are lost.

A message queue decouples the two sides:

  GPU node 1 ──┐
  GPU node 2 ──┼── gpufl-agent ──► Kafka ──► gpufl-backend
  GPU node 3 ──┘

Buffering — Kafka holds events if the backend restarts or lags behind
Burst handling — GPU workloads produce spiky event rates; Kafka absorbs the peaks
Multiple consumers — other services can read from the same topic independently: alerting, analytics, archival

For a single machine or a development setup, HTTP is simpler and perfectly adequate. The agent supports both modes with a single flag.

gpufl-backend — Ingestion and API

The backend is a Spring Boot REST service. It has two responsibilities: ingesting events and answering queries.

Ingestion happens via POST /api/v1/events/{eventType}, where eventType is device, scope, or system. In a Kafka deployment, a consumer bridge reads from the topic and calls the same endpoint internally.

The three event types:

device — kernel launches, memory copies, precise timing from CUPTI
scope — GFL_SCOPE timing blocks, optionally enriched with SASS metrics and PC samples
system — periodic NVML snapshots (GPU util%, VRAM, temperature, power draw, fan speed)

Query API — sessions list, system metrics over time, kernel events per session, profile samples per scope. The frontend is built entirely on top of this API; there's no special internal interface.

Auth — Bearer tokens for direct HTTP ingestion from the agent; API keys for programmatic access from the frontend or external tooling.

Data retention — configurable cleanup policy so disk usage stays bounded. Old sessions are pruned automatically once they exceed the configured retention window.

gpufl-front — The Dashboard

The frontend is a React + TypeScript SPA with a dark theme, connected to the backend REST API.

The main view is a three-tab dashboard scoped to a single profiling session:

┌────────────────────────────────────────────────────┐
│  Session: my_training_run  [host: gpu-node-01]     │
├──────────────┬─────────────┬───────────────────────┤
│   Kernels    │   Scopes    │       System           │
└──────────────┴─────────────┴───────────────────────┘

Kernels tab — timeline of kernel launches. Click any kernel to open the Inspector: occupancy, register count, shared memory, grid and block dimensions, and the CPU call stack that triggered the launch.
Scopes tab — hierarchical tree of GFL_SCOPE blocks with timing breakdown. For sessions using the SASS Metrics engine, each scope shows the divergence data from the previous post.
System tab — time-series charts for GPU util%, VRAM usage, temperature, power draw, and fan speed across the full session duration.

The session selector lets you browse by host and application name, so you can compare runs across machines or across time.

What's Next

Next post is a local setup guide: a Docker Compose file that brings up the full stack — client example, agent, backend, and frontend — on a single machine. If you want to try GPU Flight without configuring anything manually, that post will cover it end-to-end.

Cloud demo coming soon. I'm currently working on a live deployment. Once it's up, I'll post a link so you can browse real GPU profiling data in the UI without setting up anything yourself.

Detecting Thread Divergence with SASS Metrics and GPU Flight

Myoungho Shin — Tue, 10 Mar 2026 06:54:51 +0000

In the previous post I showed how to set up GPU Flight with Python and read kernel-level profiling data — occupancy, register counts, and resource bottlenecks. That tells you how well a kernel uses the hardware. But it doesn't tell you what's happening inside the kernel.

Today I want to look at one specific problem: thread divergence. When threads within a warp take different code paths, the GPU serializes execution — it runs one branch, then the other, while idle threads wait. If half the threads branch left and half branch right, you're running at 50% efficiency on those instructions.

GPU Flight's SASS Metrics engine gives you a direct way to measure this. It instruments the GPU at the assembly (SASS) level and reports two key counters per instruction:

smsp__sass_inst_executed — the number of warp-level instruction executions
smsp__sass_thread_inst_executed — the total number of thread-level instruction executions

The ratio thread_executed / (inst_executed × 32) tells you the average number of active threads per warp at each instruction. If it's 32.0, every thread was active. If it's 16.0, half were diverged. If it's 8.0, only a quarter was doing useful work.

The Demo: Five Divergence Patterns

I wrote a small CUDA program with five kernels, each demonstrating a different divergence pattern. The full source is in the GPU Flight repo at example/cuda/sass_divergence_demo.cu. Here's a summary:

Kernel	Pattern	Expected Active Threads
`uniformWork`	No divergence (baseline)	32
`branchByWarpLane`	`if (threadIdx.x % 2)` — even/odd split	16 in each branch
`branchByWarpQuad`	`if (threadIdx.x % 4 == 0)` — 1-in-4	8 in hot path
`earlyExit`	Data-dependent early return	Varies (~16)
`indirectBranch`	4-way switch on random data	Varies (~8)

Each kernel is wrapped in a GFL_SCOPE so GPU Flight can attribute the SASS metrics to the right section.

Kernel 1: Uniform Work (Baseline)

__global__
void uniformWork(float* out, const float* in, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float val = in[idx];
        for (int i = 0; i < 512; ++i) {
            val = val * 1.01f + 0.001f;
        }
        out[idx] = val;
    }
}

Every thread does the same thing. No branches inside the loop, no divergence. This is the baseline — you should see thread_executed / inst_executed close to 32 for the loop body instructions.

Kernel 2: Even/Odd Divergence

__global__
void branchByWarpLane(float* out, const float* in, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float val = in[idx];
        if (threadIdx.x % 2 == 0) {
            for (int i = 0; i < 512; ++i)
                val = val * 1.01f + 0.001f;
        } else {
            for (int i = 0; i < 512; ++i)
                val = val + 0.001f * (float)i;
        }
        out[idx] = val;
    }
}

This is the classic divergence example. Within every warp, 16 threads go left, 16 go right. The GPU executes both paths sequentially with half the threads masked off each time. The SASS metrics will show ~16 active threads for instructions inside each branch.

Kernel 3: Quad Divergence

__global__
void branchByWarpQuad(float* out, const float* in, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float val = in[idx];
        if (threadIdx.x % 4 == 0) {
            for (int i = 0; i < 2048; ++i)
                val = val * 1.001f + 0.0001f;
        }
        out[idx] = val;
    }
}

Only every 4th thread enters the loop. That's 8 out of 32 threads doing the heavy work while 24 sit idle. Worse than 50/50 — 75% of the warp is wasted during the loop body.

Kernel 4: Early Exit

__global__
void earlyExit(float* out, const float* in, float threshold, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float val = in[idx];
        if (val < threshold) {
            out[idx] = val;
            return;
        }
        for (int i = 0; i < 1024; ++i)
            val = val * 1.01f - 0.005f;
        out[idx] = val;
    }
}

This is data-dependent. Threads whose input is below the threshold return early, while the rest do the expensive computation. With random inputs in [0, 1) and a threshold of 0.5, roughly half the threads will exit early. But unlike Kernel 2, the split isn't uniform across warps — some warps might have 20 threads exit, others might have 10.

Kernel 5: Data-Dependent Switch

__global__
void indirectBranch(float* out, const float* in, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float val = in[idx];
        int category = (int)(val * 4.0f) % 4;
        switch (category) {
            case 0: for (int i = 0; i < 256; ++i) val = val * 1.01f; break;
            case 1: for (int i = 0; i < 256; ++i) val = val + 0.01f; break;
            case 2: for (int i = 0; i < 256; ++i) val = val - 0.005f; break;
            case 3: for (int i = 0; i < 256; ++i) val = val * 0.99f; break;
        }
        out[idx] = val;
    }
}

A 4-way branch driven by random data. On average, each case gets ~8 threads per warp, but the GPU must execute all 4 paths sequentially. This is the worst case — 4x the instruction count for the branch body.

Running the Demo

Build and run from the GPU Flight repo:

git clone https://github.com/gpu-flight/gpufl-client.git
cd gpufl-client
cmake -B build -DCMAKE_BUILD_TYPE=Release
cmake --build build --target sass_divergence_demo
./build/example/cuda/sass_divergence_demo

The key part is in main() — initializing GPU Flight with the SASS Metrics engine:

gpufl::InitOptions opts;
opts.app_name = "sass_divergence_demo";
opts.log_path = "sass_divergence";
opts.enable_kernel_details = true;
opts.sampling_auto_start = true;
opts.profiling_engine = gpufl::ProfilingEngine::SassMetrics;

gpufl::init(opts);

Setting profiling_engine to SassMetrics tells GPU Flight to instrument every kernel at the SASS level. Each GFL_SCOPE block then collects per-instruction counters for the kernels launched inside it.

Results: RTX 3090

Here's what I got running on an NVIDIA GeForce RTX 3090 (Ampere, SM 8.6, 82 SMs) with 1M elements:

Kernel                    Weighted Avg Active Threads    Instructions
----------------------------------------------------------------------
uniformWork                                      32.0             277
branchByWarpLane                                 16.3             796
branchByWarpQuad                                  8.2             281
earlyExit                                        16.2             280
indirectBranch                                    1.5            1062

The "Weighted Avg Active Threads" is thread_inst_executed / inst_executed across all SASS instructions in each kernel, weighted by execution count. "Instructions" is the number of unique PC offsets (SASS instructions) instrumented.

Let's walk through what this tells us:

uniformWork — 32.0 active threads. Perfect. Every warp runs at full width. This is the expected baseline for a kernel with no divergence.

branchByWarpLane — 16.3 active threads. Very close to the theoretical 16. The slight overshoot comes from instructions outside the branch (the if (idx < n) guard, loop control, and the final store) where all 32 threads are active. The 796 unique instructions — nearly 3x the baseline — show the cost: the compiler generates separate code for each branch, and both paths must be executed.

branchByWarpQuad — 8.2 active threads. Again close to the theoretical 8 (only 1 in 4 threads enters the loop). Similar instruction count to the baseline since there's only one branch path — but every instruction in the hot loop runs with 75% of threads idle.

earlyExit — 16.2 active threads. Matches the expectation for a 50% threshold with random data. Threads that exit early become inactive for the remaining instructions.

indirectBranch — 1.5 active threads, 1062 instructions. This is the most striking result. A 4-way switch on random data drops the weighted average to just 1.5 active threads per warp — far worse than the other kernels. It also generates the highest instruction count at 1062, nearly 4x the baseline. This is a crucial insight: divergence doesn't just halve your throughput — multi-way branching on random data can drop you below 5% when measured at the instruction level.

What This Means in Practice

Thread divergence is easy to create and hard to notice. Your kernel still produces correct results. But you might be leaving 50-95% of your GPU's compute on the table.

Here are the common patterns to watch for:

Lane-based branching — if (threadIdx.x % N). This is almost always unintentional. Consider rearranging your data so that threads within a warp take the same path.

Data-dependent branches — like the earlyExit kernel. If your input distribution is skewed, some warps diverge heavily while others don't. The average might look okay, but the worst warps are bottlenecks.

Switch statements on computed values — like indirectBranch. This was the worst offender in our test — each additional case multiplies the predicated instruction overhead.

The fix depends on the situation:

Sort or bin your data so threads in the same warp hit the same branch
Replace branches with predicated arithmetic — branchless code runs all threads at full width
Restructure your algorithm so the branch happens at the warp or block level, not the thread level

Profiling GPU (CUDA) — Getting Started with GPU Flight's Python Package

Myoungho Shin — Mon, 09 Mar 2026 03:59:53 +0000

In the previous posts I've been showing how to investigate GPU occupancy utilization and optimize kernels that aren't using the hardware fully. That was just one case — I'll cover more occupancy scenarios in future posts.

Today, I want to go through how to use GPU Flight in Python, especially with PyTorch. Since GPU Flight is still in active development, the current version is v0.1.0.dev7. You can install it with:

pip install gpufl==0.1.0.dev7

However, I highly recommend building from source inside a CUDA container. There are two reasons:

Prerequisite libraries — GPU Flight's backend needs CUPTI, the CUDA runtime, and NVML headers at compile time. Getting these right on a bare system is fiddly.
NVML support — the pre-built PyPI wheel is compiled in a minimal CI environment that doesn't include NVML stubs. This means the wheel works for kernel profiling, but can't collect runtime GPU utilization or VRAM usage. Building from source inside the nvidia/cuda:*-devel image picks up NVML automatically.

In this post, I'll show how to use Docker to set up an environment that's ready to go — with GPU Flight built from source, PyTorch, and Jupyter Lab all pre-installed.

The Dockerfile

Here's the full Dockerfile. It's straightforward — CUDA 13.1 base, PyTorch, GPU Flight, and Jupyter Lab:

FROM nvidia/cuda:13.1.0-devel-ubuntu24.04

ENV DEBIAN_FRONTEND=noninteractive

# System dependencies (Ubuntu 24.04 ships Python 3.12)
# NOTE: cmake/ninja come from pip (build-system.requires needs >=3.31, apt has 3.28)
RUN apt-get update && apt-get install -y \
    python3 \
    python3-venv \
    python3-dev \
    python3-pip \
    git \
    curl \
    && rm -rf /var/lib/apt/lists/*

# Create venv to avoid PEP 668 issues
RUN python3 -m venv /opt/venv
ENV PATH="/opt/venv/bin:$PATH"

# Upgrade pip
RUN pip install --upgrade pip

# Install PyTorch with CUDA 13.1 support
RUN pip install torch --index-url https://download.pytorch.org/whl/cu130

# Build gpufl from source so it picks up NVML from the CUDA devel image
ARG GPUFL_VERSION=main
RUN git clone --depth 1 --branch ${GPUFL_VERSION} \
        https://github.com/gpu-flight/gpufl-client.git /tmp/gpufl-client \
    && CMAKE_ARGS="-DBUILD_TESTING=OFF" \
       pip install -v "/tmp/gpufl-client[analyzer,viz]" \
    && rm -rf /tmp/gpufl-client

# Install Jupyter
RUN pip install jupyterlab

# Working directory for notebooks
WORKDIR /workspace

# Expose Jupyter port
EXPOSE 8888

# Start Jupyter Lab
CMD ["jupyter", "lab", "--ip=0.0.0.0", "--port=8888", "--no-browser", \
     "--allow-root", "--NotebookApp.token=''", "--NotebookApp.password=''"]

A few things to note:

Ubuntu 24.04 — ships Python 3.12 natively, which is what GPU Flight requires. No PPA hacks needed.
devel image — we use nvidia/cuda:13.1.0-devel-ubuntu24.04 because the devel variant includes CUPTI, CUDA headers, and NVML stubs that GPU Flight's backend needs at compile time.
Building from source — we clone the repo and build with pip install rather than using the pre-built PyPI wheel. This is important: the devel image has NVML stubs at /usr/local/cuda/lib64/stubs/libnvidia-ml.so, so CMake detects them and compiles in the NVML collector. The pre-built wheel doesn't have this, which means no GPU utilization or VRAM monitoring.
PyTorch cu130 — at the time of writing, PyTorch doesn't publish a cu131 wheel yet. The cu130 build is forward-compatible with the CUDA 13.1 runtime in the container, so this works fine.
No token — Jupyter starts without authentication. This is fine for local development; don't expose this to the internet.

Building and Running

Prerequisites

You need two things on your host machine:

Docker — any recent version
NVIDIA Container Toolkit — this lets Docker containers access your GPU

Important: Having an NVIDIA driver installed on your host is not enough. Docker doesn't know how to talk to your GPU on its own — you need the NVIDIA Container Toolkit to bridge that gap. Without it, --gpus all will fail with:

docker: Error response from daemon: could not select device driver "" with capabilities: [[gpu]]

You can check if it's already installed by running nvidia-ctk --version. If not, here's how to set it up:

# Add the NVIDIA container toolkit repo
curl -fsSL https://nvidia.github.io/libnvidia-container/gpgkey \
  | sudo gpg --dearmor -o /usr/share/keyrings/nvidia-container-toolkit-keyring.gpg

curl -s -L https://nvidia.github.io/libnvidia-container/stable/deb/nvidia-container-toolkit.list \
  | sed 's#deb https://#deb [signed-by=/usr/share/keyrings/nvidia-container-toolkit-keyring.gpg] https://#g' \
  | sudo tee /etc/apt/sources.list.d/nvidia-container-toolkit.list

# Install and configure
sudo apt-get update && sudo apt-get install -y nvidia-container-toolkit
sudo nvidia-ctk runtime configure --runtime=docker
sudo systemctl restart docker

That last systemctl restart is easy to forget — Docker needs to be restarted after the runtime is configured, or it won't pick up the new GPU capability.

You can verify it worked with:

docker run --rm --gpus all nvidia/cuda:13.1.0-base-ubuntu24.04 nvidia-smi

If you see your GPU listed, you're good to go.

Build the Image

docker build -t gpufl-python .

This will take a few minutes the first time — mostly downloading PyTorch.

Run the Container

docker run --gpus all -p 8888:8888 -v $(pwd)/notebooks:/workspace gpufl-python

Breaking that down:

Flag	What it does
`--gpus all`	Passes all GPUs into the container
`-p 8888:8888`	Maps Jupyter's port to your host
`-v $(pwd)/notebooks:/workspace`	Mounts a local folder so your notebooks persist

Connect

Open your browser and go to:

http://localhost:8888

You'll land in Jupyter Lab with GPU Flight, PyTorch, and a CUDA-capable GPU ready to go.

Quick Smoke Test

Create a new notebook and run this to verify everything is working:

import torch
import gpufl
from gpufl import ProfilingEngine

print(f"PyTorch: {torch.__version__}")
print(f"CUDA available: {torch.cuda.is_available()}")
print(f"Device: {torch.cuda.get_device_name(0)}")

# Initialize GPU Flight
gpufl.init("smoke-test",
           log_path="./smoke_test",
           sampling_auto_start=True,
           enable_kernel_details=True,
           enable_stack_trace=True,
           profiling_engine = ProfilingEngine.RangeProfiler)

# Run a simple operation
with gpufl.Scope("RandomGeneration"):
    a = torch.randn(1024, 1024, device="cuda")
    b = torch.randn(1024, 1024, device="cuda")
with gpufl.Scope("a @ b"):
    c = a @ b
    torch.cuda.synchronize()

gpufl.shutdown()
print("GPU Flight logs written!")

After running this, you should see *.log files in your working directory. These are your GPU Flight recordings — every kernel launch, memory copy, and timing event that happened during that matrix multiply.

Analyzing the Results

GPU Flight's Python analyzer can load those logs directly in the notebook:

from gpufl.analyzer import GpuFlightSession

session = GpuFlightSession(".", log_prefix="smoke_test")
session.print_summary()

GpuFlightSession takes two main arguments: the directory where logs live, and the log_prefix matching your log_path from init. It automatically finds and loads smoke_test.device.log, smoke_test.scope.log, and smoke_test.system.log.

print_summary() gives you a quick dashboard — total duration, kernel count, GPU busy time, average utilization, and peak VRAM.

Now let's look at the kernel hotspots:

session.inspect_hotspots(top_n=10)

This gives you a Rich-formatted table of your hottest kernels with occupancy, register usage, shared memory, and the per-resource occupancy breakdown showing exactly what's limiting each kernel.

Here's what that actually looks like — this is real output from the matrix multiply we just ran:

Now let's look at 33.3% occupancy. That sounds bad, right? Let's break it down.

The kernel is ampere_sgemm_128x64_nn — cuBLAS's single-precision matrix multiply. It uses 122 registers per thread. That's a lot. Let's trace through what happens on an Ampere SM:

128 threads per block = 4 warps per block
122 regs/thread × 32 threads/warp = 3,904 → rounded up to the hardware allocation granularity of 256 → 4,096 regs/warp
4 warps × 4,096 = 16,384 registers per block
An Ampere SM has 65,536 registers total → 65,536 / 16,384 = 4 blocks max
4 blocks × 4 warps = 16 active warps out of 48 max = 33.3%

The breakdown confirms it: reg 33.3% is the bottleneck, while shared memory (66.7%), warps (100%), and block count (100%) all have headroom.

But Is This Actually a Problem?

Not necessarily. If the algorithm itself doesn't require all those registers, high register usage might be a problem — but it could also be by design. This is a good example of why occupancy alone doesn't tell the whole story — you need to understand what's limiting it and whether that tradeoff makes sense for the workload.

If you saw 33% occupancy with limiting_resource: shared_mem on your own custom kernel, that might be worth investigating.

What's Next

Now that you have a working environment, you can start profiling your own models. The occupancy breakdown makes it easy to spot which kernels are underutilizing the GPU and — more importantly — why. Not every low-occupancy kernel is a problem, but when one is, you'll know exactly which resource to optimize.

In the next post, I'll cover GPU Flight's profiling engines — PC sampling, SASS metrics, and the range profiler — which let you go beyond kernel metadata and collect hardware-level data about what's happening inside the GPU while your kernels run.

Profiling GPU (CUDA) — What Is Actually Limiting Your Kernel?

Myoungho Shin — Mon, 02 Mar 2026 01:19:03 +0000

In my last post I introduced GPU Flight — a lightweight CUDA observability tool that acts like a flight recorder for your GPU. We covered what it collects: system metrics, device capabilities, and per-kernel events.

Today I want to talk about one specific metric that GPU Flight captures: occupancy. It's one of the most important numbers for understanding GPU performance, and also one of the most misunderstood.

What Is Occupancy?

A GPU is organized around Streaming Multiprocessors (SMs). Each SM can run many threads simultaneously — not by context-switching like a CPU, but by actually running them in parallel. The unit of scheduling on an SM is a warp: a group of 32 threads that execute the same instruction in lockstep.

An SM has a fixed warp budget — say, 48 warps on a typical Ampere GPU. When you launch a kernel with blocks of 256 threads (8 warps each), the SM can hold up to 6 blocks concurrently to fill those 48 warp slots. If something prevents that — too many registers, too much shared memory — fewer blocks fit, and some warp slots sit idle.

Occupancy measures how well those warp slots are filled:

occupancy = active warps / maximum warps per SM

A value of 1.0 means every slot is in use. A value of 0.5 means half the SM's compute capacity is being wasted while your kernel runs.

How GPU Flight Captures It

GPU Flight records occupancy automatically for every kernel launch. No code changes needed — just initialize with enableKernelDetails: true and it shows up in the log:

{
  "type": "kernel_event",
  "name": "_Z18block_reduce_naivePKfPfi",
  "occupancy": 0.833333,
  "num_regs": 16,
  "static_shared_bytes": 16384,
  "dyn_shared_bytes": 0,
  "block": "(256,1,1)",
  "grid": "(16384,1,1)",
  "max_active_blocks": 5,
...
}

Under the hood, GPU Flight calls cudaOccupancyMaxActiveBlocksPerMultiprocessor at kernel launch time to get max_active_blocks, then divides by the SM's warp budget to compute occupancy. This happens inside the CUPTI callback — zero overhead to your kernel execution.

That 0.833333 immediately tells you something is off. This kernel only fills 5 out of 6 possible concurrent blocks on each SM. Some compute is being left on the table.

But What Is Actually Causing It?

Here's where a single number hits its limit.

Is it registers? Shared memory? The hardware block count cap? Looking at the log fields, you can make an educated guess — static_shared_bytes: 16384 is 16 KB of shared memory per block, which is pretty large. But you still have to do the math yourself against your specific GPU's properties to confirm.

That manual detective work is exactly what I wanted to eliminate. So GPU Flight now also computes a per-resource occupancy breakdown and identifies the limiting resource automatically. Let me show what this looks like with a concrete kernel.

The kernel

Here's a simple parallel block reduction — it sums an array by having all 256 threads in a block cooperate through shared memory:

__global__ void block_reduce_naive(const float* in, float* out, int n) {
    __shared__ float smem[4096]; // 16 KB — statically reserved

    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + tid;

    // Load one element per thread into shared memory
    smem[tid] = (gid < n) ? in[gid] : 0.0f;
    __syncthreads();

    // Reduce in shared memory — each step halves the active threads
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) smem[tid] += smem[tid + s];
        __syncthreads();
    }

    // Thread 0 writes the block's result
    if (tid == 0) out[blockIdx.x] = smem[0];
}

Launched with 256 threads per block across 4M elements:

const int BLOCK = 256;
const int GRID  = (N + BLOCK - 1) / BLOCK; // ~16384 blocks
block_reduce_naive<<<GRID, BLOCK>>>(d_in, d_out, N);

Nothing unusual here — this is a textbook reduction. But GPU Flight flags a problem immediately.

What GPU Flight sees

{
...,
  "occupancy":         0.833333,
  "reg_occupancy":     1.0,
  "smem_occupancy":    0.833333,
  "warp_occupancy":    1.0,
  "block_occupancy":   1.0,
  "limiting_resource": "shared_mem"
}

Each *_occupancy field answers: "if only this constraint existed, what would occupancy be?" The limiting_resource field names the one that's actually binding. Here — smem_occupancy matches occupancy and everything else is 1.0 — shared memory is definitively the culprit.

Why

The problem is __shared__ float smem[4096]. Static shared memory is sized at compile time and reserved in full for the block's entire lifetime — even if the kernel only uses part of it. With 256 threads per block, this reduction only ever touches smem[0] through smem[255], but all 4096 floats (16 KB) are locked up on the SM regardless. Every block is paying a 16 KB reservation it doesn't actually need, and that prevents the SM from scheduling as many concurrent blocks as the warp budget would otherwise allow.

The fix

Switch to dynamic shared memory, which is sized at launch time rather than compiled in:

__global__ void block_reduce_optimized(const float* in, float* out, int n) {
    extern __shared__ float smem[]; // size comes from the launch call

    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + tid;

    smem[tid] = (gid < n) ? in[gid] : 0.0f;
    __syncthreads();

    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) smem[tid] += smem[tid + s];
        __syncthreads();
    }

    if (tid == 0) out[blockIdx.x] = smem[0];
}

The kernel body is completely unchanged. The only differences are extern __shared__ instead of a fixed-size array, and passing the size as the third launch argument:

size_t smem_bytes = BLOCK * sizeof(float); // 256 × 4 = 1 KB
block_reduce_optimized<<<GRID, BLOCK, smem_bytes>>>(d_in, d_out, N);

The shared memory footprint drops from 16 KB to 1 KB per block — 16× smaller — and now the SM can fit all 6 concurrent blocks instead of 5.

GPU Flight confirms the fix worked:

{
  "occupancy":         1.0,
  "limiting_resource": "warps"
}

"warps" as the limiting resource means full occupancy — every SM warp slot is filled and shared memory is no longer in the way.

Full Sample Code: GitHub Repo

Profiling GPU (CUDA) — Introducing GPU Flight

Myoungho Shin — Tue, 24 Feb 2026 01:32:26 +0000

Last year, I took a GPU programming course at Johns Hopkins University as part of my graduate studies, where I learned CUDA programming. For my final project, I built a lightweight GPU monitoring and profiling tool focused on CUDA.

I enjoyed the process so much that I decided to continue developing it beyond the course.

In this post, I’d like to briefly introduce the project:

GPU Flight — a 100% open-source GPU observability tool

GitHub: https://github.com/gpu-flight/gpufl-client

Why I Started GPU Flight

When profiling a CUDA application, you typically:

Install profiling tools such as Nsight
Or manually integrate CUPTI into your application, which often makes the code complex and difficult to manage
Deal with additional complexity in cloud or containerized environments

This workflow can be inconvenient — especially in production systems.

I wanted something lighter.

Something that works more like a flight recorder for GPUs.

So I built GPU Flight.

Instead of requiring heavy tooling at runtime, GPU Flight writes structured profiling logs directly on the host machine. A separate component (GPUFL Agent) crawls these log files and forwards them to a backend service or other destinations.

This makes GPU observability more flexible and easier to integrate into distributed systems.

What is GPU Flight?

GPU Flight is designed to be lightweight and modular.

If you only need monitoring, the overhead is minimal.
Enabling deeper profiling provides more detailed metrics.

The goal is to expose useful GPU metrics so you can clearly understand:

How the GPU manages resources
How your program utilizes GPU resources
Where performance bottlenecks occur

Project Structure

GPU Flight currently consists of several components:

1️⃣ gpufl-client

https://github.com/gpu-flight/gpufl-client

The client library that users embed into their applications for monitoring and profiling.

2️⃣ gpufl-agent

https://github.com/gpu-flight/gpufl-agent

Despite the name, this is not an AI agent 🙂

It tracks log files and forwards profiling data to the configured destination.

3️⃣ gpufl-desktop

https://github.com/gpu-flight/gpufl-desktop

Originally, I planned to build a desktop viewer. Due to time constraints and the need for better cross-platform accessibility, I pivoted to a web-based frontend.
I am currently keeping the web frontend and backend repositories private as I develop them into a hosted cloud platform. To ensure the open-source community can still easily parse and utilize the trace logs locally, I am providing a lightweight Python viewer alongside the open-source C++ client.

What Metrics Does GPU Flight Support?

GPU Flight captures observability at multiple layers.

1️⃣ System & GPU Monitoring (NVML)

Host memory usage
GPU memory usage (used/free/total)
GPU utilization
Memory utilization
Temperature
Power consumption
Clock speeds (GFX / SM / Memory)
PCIe RX/TX bandwidth
Power and thermal throttling flags

Example JSON snippet:

{
  "type": "system_sample",
  "util_gpu": 57,
  "temp_c": 39,
  "power_mw": 54415,
  "clk_sm": 1740
}

2️⃣ CUDA Device Capabilities

Static architectural information:

Compute capability
L2 cache size
Shared memory per block
Registers per block
SM count
Warp size

3️⃣ CUDA API & Kernel Events (CUPTI)

API enter/exit timestamps
Kernel execution start/end timestamps
Grid/block dimensions
Shared memory usage
Register usage
Occupancy
Correlation IDs
Memory copy events (HtoD, DtoH)

Python Support

GPU Flight is also being extended to support Python applications that use CUDA (e.g., PyTorch).

Example:

https://github.com/gpu-flight/gpufl-client/blob/main/example/python/03_pytorch_benchmark.py

This allows profiling GPU-heavy ML workloads without deeply modifying existing code.

What’s Next?

In the next post, I’ll walk through a minimal CUDA example and show how to:

Integrate gpufl-client
Run a kernel
Inspect generated profiling logs
Interpret stall reasons and metrics

Thanks for reading — this is just the beginning