100 DAYS OF GPU Course preview
Day 41Module 5
in-technical-review

Reading an Nsight Systems timeline

An unmarked trace lists each launch in order, but labels every host call cudaLaunchKernel. This program issues that call more than a thousand times, so the report groups them in one row under one name.

The profiler knows the API names, but it does not know the phases of your algorithm. NVTX adds those phase names with one line at each range boundary. Later profiler lessons use the same ranges.

You will make two captures of one binary, one without NVTX data and one with named ranges. You will also measure the cost of the convergence test.

What a timeline shows that a timer cannot

Nsight Systems traces a whole process against one clock and draws it as rows. The rows you care about today are four:

  • CUDA API, what the host thread called and how long that call blocked it.
  • Kernels, what ran on the device and when.
  • Memory, every copy, with its direction and its size.
  • NVTX, your own phase names, once you add them.

Everything this course has measured so far came from CUDA events, which report device time between two records on a stream. That answers "how long did this kernel run" and is blind to the question underneath it: what was the host doing meanwhile.

A gap on the kernel row under a busy API row means the host did not submit device work in time. The cause may be launch overhead or other host work.

A long API bar over a busy kernel row means the host is blocked while the device runs. These two cases can look alike, so check both rows.

Hardware. nsys collects through CUPTI's activity and callback interfaces, not through hardware performance counters, so it traces CUDA and NVTX as an ordinary user on a local card and on a free Colab T4. Nsight Compute counter access is a separate permission question; it was blocked for an ordinary user on this project's node, while the current free-Colab case remains untested. Its CPU sampling is the part that wants /proc/sys/kernel/perf_event_paranoid at 2 or less, and NVIDIA states the fallback plainly: "If the Sampling Environment is not OK, you will still be able to run various trace operations." (https://docs.nvidia.com/nsight-systems/InstallationGuide/index.html , checked 2026-08-30). Run nsys status -e to see which half you have. If you cannot install nsys, this page ships its own reports and the exercise is doable from them.

Diagram: one PageRank iteration, three ways. Three horizontal bands, each with a CUDA API row above a GPU kernel row above an NVTX row, all on one time axis. Band 1, no annotation: five short API bars all labelled cudaLaunchKernel, five kernel bars under them, the NVTX row empty. Caption "5 launches, 1 name." Band 2, annotated: the same two rows, plus one iteration bar holding scale, spmv, dangling-mass and combine nested inside it. Caption "5 launches, 4 phase names, 1 iteration." Band 3, the convergence test added: the API row gains a cudaMemcpy bar that opens where the last kernel was queued and closes where it finishes, so it spans the whole kernel row, with copy-residual under it. Caption "4 bytes, and the bar covers everything queued before it." Alt text: "Five kernel launches share one API name until NVTX gives each a phase name. The four-byte convergence copy draws a bar that spans every kernel queued before it."

Why kernel names are not enough

Kernel names do not replace NVTX ranges. The kernel row lacks three kinds of context.

It has no nesting. Fifty iterations of five kernels give you 250 bars in five repeating colours and no way to say "show me iteration 34", which is the one that stalled.

It only covers kernels. A host-side graph transpose, a cudaMalloc, the counting sort day 40 runs before device work begins: none of those is a kernel, so none gets a kernel name. Host work can dominate the total time.

And it names the function, not the role. sumPartials runs twice per iteration in this program, once for the dangling mass and once for the residual. Same kernel, same name, two different jobs, and only one of them is on the critical path of the convergence test.

An NVTX range is a labelled interval on the calling thread. nvtxRangePushA opens one and nvtxRangePop closes it, they nest per thread, and NVTX v3 is header-only, so there is no library to link. The library "introduces close to zero overhead if no tool is attached" (https://github.com/NVIDIA/NVTX/blob/release-v3/c/include/nvtx3/nvToolsExt.h , checked 2026-08-30), which is why the ranges below stay in the shipped binary instead of hiding behind a build flag.

Two captures of one binary

Full program in code/day41-nsys/pagerank_nvtx.cu. Four decisions shape it.

One binary, two command lines. nsys profile -t cuda traces the CUDA API and the device and ignores the NVTX calls; -t cuda,nvtx picks them up. So the anonymous capture and the named capture come from the same executable and the same source, and nothing can drift between them. A -D flag and two builds would have given two programs to keep in step.

Only the convergence policy moves. Three timed loops, the same fifty iterations, the same kernels, differing only in how often the host reads the residual: never, every tenth iteration, every iteration. The program then requires the three rank vectors to be bit-identical, so "the copy policy changed the answer" is a failure rather than a footnote.

Every phase is a range, and the ranges nest. iteration wraps scale, spmv, dangling-mass and combine, so nsys stats --report nvtx_sum gives you the phase breakdown as a table without opening the GUI.

Events time it, the timeline explains it. The program prints a per-kernel table from CUDA events with a warm-up per kernel, because lazy module loading makes each kernel's first launch pay its own load. If that table and the profiler disagree, one of the two is measuring something else.

Here is the iteration, five launches inside four ranges:

static void iterate(const DeviceGraph& g, const Work& w, const float* x,
                    float* xNext) {
    const int blocks = gridFor(g.n);

    nvtxRangePushA("scale");
    scaleByOutDegree<<<blocks, kThreadsPerBlock>>>(x, g.invOut, w.contrib, g.n);
    nvtxRangePop();

    nvtxRangePushA("spmv");
    spmvCsrScalar<<<blocks, kThreadsPerBlock>>>(g.rowPtr, g.colIdx, w.contrib,
                                                w.y, g.n);
    nvtxRangePop();

    nvtxRangePushA("dangling-mass");
    sumDanglingPartial<<<kReduceBlocks, kThreadsPerBlock>>>(x, g.dangling,
                                                            w.partials, g.n);
    sumPartials<<<1, kThreadsPerBlock>>>(w.partials, w.scalars, kReduceBlocks);
    nvtxRangePop();

    nvtxRangePushA("combine");
    combineRanks<<<blocks, kThreadsPerBlock>>>(w.y, w.scalars, xNext, g.n,
                                               kDamping);
    nvtxRangePop();
}

spmv is the range to watch on the kernel row. It is day 37's thread-per-row kernel, so its gather through colIdx lands each lane in a 32-byte sector of its own instead of coalescing into four, exactly as day 11 measured. Naming it does not fix it. It

tells you whether it is what your program is waiting on.

And here is the convergence test, which is why this page needs a timeline rather than another table. Whether to stop is a host decision, so one float has to come back, and a synchronous device-to-host cudaMemcpy "returns only once the copy has completed" (https://docs.nvidia.com/cuda/cuda-runtime-api/api-sync-behavior.html , checked 2026-08-30). It is a device synchronize wearing four bytes as a disguise.

static void residual(const Work& w, const float* x, const float* xPrev,
                     size_t n) {
    nvtxRangePushA("residual");
    sumAbsDiffPartial<<<kReduceBlocks, kThreadsPerBlock>>>(x, xPrev, w.partials,
                                                           n);
    sumPartials<<<1, kThreadsPerBlock>>>(w.partials, w.scalars + 1,
                                         kReduceBlocks);
    nvtxRangePop();

    float value = 0.0f;
    nvtxRangePushA("copy-residual");
    CUDA_CHECK(cudaMemcpy(&value, w.scalars + 1, sizeof(float),
                          cudaMemcpyDeviceToHost));
    nvtxRangePop();
    (void)value;
}

Note. The graph is generated in the program, not read from day 40's OpenAlex files. Same node count and a similar edge count, so the timeline has the same shape, but the ranks are not day 40's ranks and this page never prints a top twenty. This keeps the exercise self-contained: the program runs without a data file. It also fits in L2, which is why the program reports milliseconds and no GB/s.

Results

Measured. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo and profiled with Nsight Systems 2024.3.2.

Captured 2026-09-01 on the project's verification node. The full transcript is in the page's evidence file, and both reports and their CSV exports ship at content/profiles/nsight-systems/.

Re-verified on 2026-09-02 with CUDA 13.0 (V13.0.88), driver 580.173.02, on the same Tesla T4. The program passed every correctness gate and the checked-every-step policy remained a 1.15x tax. Fresh plain and NVTX Nsight Systems captures are listed in evidence; the NVTX export measured 40805620 ns of copy-residual against 2415822 ns of spmv, so the lesson's timeline conclusion is unchanged.

one launch each, mean of 10 runs after 3 warm-ups
kernel                        ms
--------------------  ----------
scaleByOutDegree          0.0059
spmvCsrScalar             0.1786
sumDanglingPartial        0.0095
sumAbsDiffPartial         0.0088
sumPartials               0.0050
combineRanks              0.0041
one iteration, summed     0.2031

50 iterations, the same arithmetic. Only the convergence
test moves, and it reads four bytes when it runs.
policy                  total (ms)  per iter (ms)  D2H
--------------------  ------------  -------------  ---
never checked               9.9676         0.1994    0
checked every 10           10.1131         0.2023    5
checked every step         11.3456         0.2269   50
checked every step costs 1.14x never checked

And the top of the nvtx_sum export from the -t cuda,nvtx capture, which is the artifact the exercise reads:

 Time (%)  Total Time (ns)  Instances   Avg (ns)   Range
 --------  ---------------  ---------  ----------  ------------------
     40.9         98769395          1  98769395.0  :correctness
     22.3         53982410        359    150368.8  :iteration
     15.3         36972808        164    225444.0  :copy-residual
      2.1          5125947        359     14278.4  :dangling-mass
      1.3          3066298        359      8541.2  :scale
      1.1          2705884        359      7537.3  :spmv
      1.1          2687104        359      7485.0  :combine
      0.9          2145782        164     13084.0  :residual

What I expected, and what happened

  1. Matched. The widest range inside iteration is copy-residual: 36.97 ms across its 164 instances against 2.71 ms for all 359 spmv instances, 13.7x in total time for a range that moves four bytes against one that walks the whole edge list. The cost is not the bytes. It is the synchronisation: the host cannot decide whether to stop until the device answers, and the answer costs a full round trip.

  2. Did not match. Checked every step costs 1.14x never checked, under the 1.2 the prediction named as its floor. The sync is cheaper here than day 9's round trip suggested, because this copy rides on a warm context inside a tight loop.

    The prediction said under 1.2 would probably kill claim 1 too, and that half was wrong: copy-residual still dwarfs every compute range inside the iteration. Both things are true at once, a 14 percent tax on the whole loop and the single widest range inside it.

  3. Partly matched. The -t cuda capture's nvtx_sum is empty and the -t cuda,nvtx capture yields the thirteen named ranges. On call count, cudaLaunchKernel tops the API table at 2,219 calls. On total time it sits second at 16.9 percent, behind cudaMemcpy at 68.6, so "topped by" held for the count the exercise sorts by and not for time.

  4. Matched. The widest single API bar is the first CUDA call in the process, a cudaMemcpy at 18.0 ms against a 201 us median for the same API, because context creation hides inside it. Under the profiler the effect is smaller than day 9's bare 255.966 ms, but it is still 90 times the loop's whole per-iteration budget, and it is nowhere near the loop.

Read the kernel and API rows together. In the shipped report, the kernel row goes idle during copy-residual while the API row stays busy. The device has finished, and the host is deciding whether to stop.

Run it yourself

Run this on a CUDA system with Nsight Systems installed. Compile for your GPU's compute capability if it does not support sm_75.

nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o pagerank_nvtx pagerank_nvtx.cu
nsys profile -t cuda      -o day41-plain ./pagerank_nvtx
nsys profile -t cuda,nvtx -o day41-nvtx  ./pagerank_nvtx

nsys ships with the full CUDA Toolkit installer and as a standalone download at https://developer.nvidia.com/nsight-systems . A toolkit installed through conda gives you nvcc and not nsys, and it needs the cuda-nvtx-dev package before <nvtx3/nvToolsExt.h> resolves.

There is no Compiler Explorer embed: no profiler runs there, and the program's three timed loops and its double-precision reference go past the 20 second run cap. If you have no GPU, read /setup/learn-cuda-without-a-gpu and do the exercise from the shipped reports.

Exercise

Take your own pagerank.cu from day 40, wrap each of the five launches in iterateHand in an NVTX range plus one around the whole iteration, and profile a converging run with -t cuda,nvtx. Then find the widest bar in one iteration and say whether it is one of the ranges you wrote.

Time: 30 to 45 minutes. Submit: your nvtx_sum table, the name of the widest bar, and one sentence on what it is measuring.

Check: use a ratio so results remain comparable across GPUs. Take the widest host-side bar in an iteration and divide its total time by the sum of your five kernel ranges.

Under 1.1 means every bar you have is a launch and you have not annotated the thing that blocks. Over 1.1 means you found it, and the excess is what the host spent standing still.

nsys stats --report nvtx_sum --report cuda_api_sum --format csv --output . day41-nvtx.nsys-rep prints both tables, and the shipped report gives the same two if you cannot run nsys.

Hint 1

A range measures the calling thread, not the device. Which call inside one iteration returns immediately, and which one refuses to return until the device has caught up?

Hint 2

Look at cuda_api_sum rather than your own ranges, and sort by total time instead of by call count. Then put a range around whatever is at the top and compare it against the five you already have.

Solution

The five ranges you wrote are all launches, and a launch returns as soon as the work is queued, so all five are narrow whatever the kernels cost. The widest bar in the iteration belongs to a call you did not annotate: the cudaMemcpy that brings the convergence residual back. It cannot start until the queue ahead of it has drained, so its bar spans the kernels rather than sitting beside them.

Use the ratio, not the raw time. The shipped figure comes from content/profiles/nsight-systems/day41-nvtx_nvtx_sum.csv; compare the range layout as well as the values.

A blocking API bar includes the device work that it waits for. Compare the bar with the kernel row, then measure the part left after that work.

Checking convergence every tenth iteration removes nine of ten host waits without changing the answer.

Pitfalls

Every bar on your timeline says cudaLaunchKernel. Either the program has no NVTX ranges, or the capture was taken without them: -t cuda alone drops NVTX even when the calls are compiled in. Add nvtx to the trace list and re-run.

A push with no pop. nvtxRangePushA and nvtxRangePop nest per thread, so one missing pop places everything after it inside the range you meant to close. The timeline then shows one long bar that resembles a stall. An early return between a push and its pop is the usual cause.

You read the widest bar as the slowest thing. A blocking API bar spans everything it waits for. Before concluding that a copy is expensive, check whether the kernel row under it is busy for the same interval. If it is, the API duration includes the kernel time.

Startup dominates the full-program trace. Context creation, allocation and any host-side setup all land on the timeline, and on a short program they dominate it. Wrap the part you care about in its own NVTX range and read that range, or use --capture-range to start collection later. Day 9 prices the startup itself.

nsys status -e says the sampling environment failed. CPU sampling and CUDA tracing are separate. NVIDIA's own installation guide says that with sampling unavailable "you will still be able to run various trace operations", and CUDA plus NVTX tracing is one of them.

You profiled a timing harness instead of a program. This program's per-kernel table launches each kernel thirteen times to get a mean. That helps timing but makes the trace hard to read. Profile the loop you care about and leave the repetition to the event timers.

Go deeper

Next

Day 42 profiles one kernel with hardware performance counters and explains ERR_NVGPUCTRPERM. It includes a shipped .ncu-rep, just as this page includes a .nsys-rep.

Day 43 optimizes the slowest range from this trace. Day 48 reduces the launch count.