100 DAYS OF GPU Course preview
Day 88Module 9
in-technical-review

Profiling a PyTorch model down to the kernel

This CUDA profile comes from PyTorch's profiler recipe and is sorted by total CUDA time:

                                                   Name     Self CUDA    CUDA total
                                        model_inference       0.000us      11.666ms
                                           aten::conv2d       0.000us      10.484ms
                                      aten::convolution       0.000us      10.484ms
                                     aten::_convolution       0.000us      10.484ms
                              aten::thnn_conv2d_forward      10.484ms      10.484ms

Four of the five rows show 10.484 ms of GPU time, but the whole run contains 11.666 ms (https://docs.pytorch.org/tutorials/recipes/recipes/profiler_recipe.html , checked 2026-09-01). Summing the column gives more time than the program used because these calls are nested. Only the last row ran a kernel.

This lesson shows how to read the table, map an aten:: row to a kernel name, and decide whether a custom operator can improve the measured work.

Self time, total time, and the rows that share one kernel

The PyTorch profiler records two time values. Each operator gets a total time, which includes its children, and a self time, which excludes work attributed to a child. aten::conv2d dispatches to aten::convolution, so both total values include the same subtree.

The self column locates the work within that call tree.

The sort key changes the ranking. Sorting by cuda_time_total ranks wrappers above the thing they wrap; sorting by self_cuda_time_total ranks the leaves, which are the rows worth reading.

Two kinds of row can be a leaf: a dispatched operator that launched a kernel and waited for nothing else, and the kernel itself, which the profiler lists under its device name. The recipe says so plainly: "Note the occurrence of on-device kernels in the output (e.g. sgemm_32x32x32_NN for CUDA ...)".

The aten:: row names the model operation, while the mangled C++ name under it identifies the GPU kernel. Nsight Systems uses the same kernel name.

torch.profiler traces through CUPTI's activity interface, the same one nsys uses, so it needs no hardware performance counters and no root. record_shapes=True adds each operator's input shapes to the record, which turns one aten::add row averaging five different calls into five rows you can tell apart.

Diagram: one profiler table, three shapes of row. A left column of operator rows with indentation showing nesting, and a right column giving each row's self CUDA time, with the launched kernel drawn as a leaf under the row that launched it. Band 1, the Linear: aten::linear, aten::addmm and one GEMM kernel leaf, with zeros in the self column on the two wrappers. Caption "3 rows, 1 kernel, 1 nonzero self time." Band 2, the eager chain: five sibling rows, mul, add, gelu, add, clamp, each with its own elementwise kernel leaf. Caption "5 rows, 5 kernels, 11 floats moved per element." Band 3, the fused operator: one row, day88::fused_chain, one kernel leaf. Caption "1 row, 1 kernel, 3 floats per element." Alt text: "A profiler table drawn as a tree. Wrapper rows carry zero self time and the leaf under them carries all of it. Five elementwise rows move eleven floats per element where one fused row moves three."

Choose a row you can improve

The slowest row in a model is often a GEMM implemented by cuBLAS. Day 44 spent a whole lesson getting a hand-written float4 SIMT matmul to 67.7 percent of cublasGemmEx on the measured GPU. Replacing that library kernel is unlikely to improve the model.

Several small elementwise rows can be a better target: a scale, a bias, an activation, a residual add, a clamp. Individually none of them is the bottleneck.

Together they read and write the same tensor over and over, which is the memory-bound shape day 48 measured: fusing four such stages into one kernel cut the bytes moved by 3.00x and the time by 3.08x. A model-specific chain may have no matching library call.

Use a library when one implements the operation. Write a custom operator when fusion removes intermediate memory traffic that the library interface cannot express. This follows day 39: use the library where it has an entry point.

The profile distinguishes these cases. The program below measures both.

One block, two spellings

Full program in code/day88-torch-profiler/profile_model.py, with the operator in fused_chain.cu beside it. One nn.Linear of 1024 by 1024 over a batch of 8192, then day 48's chain, written twice.

The two paths compute the same function. Both blocks share one set of weights through load_state_dict, the chain constants are the same literals in the Python and in the kernel, and both sides ask for the tanh approximation of GELU by name. Where the two differ is where the intermediates live, and nowhere else.

The elementwise tensor does not fit in cache. At 8192 by 1024 floats, the tensor is 32 MiB. The measured GPU has 4 MiB of L2, so each eager stage accesses DRAM.

A smaller chain could measure cache bandwidth instead.

Timing is CUDA events, profiling is separate. Collection adds uneven overhead, so the program times first and profiles second. It does not mix the two measurements.

class Block(nn.Module):
    """A Linear, then day 48's chain. `fused=True` swaps the five
    elementwise operators for one call into fused_chain.cu."""

    def __init__(self, dim, fused):
        super().__init__()
        self.fc = nn.Linear(dim, dim, bias=True)
        self.fused = fused

    def forward(self, x, residual):
        h = self.fc(x)
        if self.fused:
            return torch.ops.day88.fused_chain(h, residual)
        h = h * GAMMA + BETA
        h = F.gelu(h, approximate="tanh")
        h = h + residual
        return h.clamp(LO, HI)

Count the floats each spelling moves per element. The eager path pays two for the multiply, two for the bias add, two for the GELU, three for the residual add because it reads a second tensor, and two for the clamp: eleven. The fused path reads h, reads the residual and writes the output: three.

Day 48's hand-staged version of the same chain moved nine rather than eleven, because it did the scale and the bias in one kernel. Eager PyTorch cannot, since h * GAMMA + BETA is two operators, so it is one extra launch and one extra round trip. The eager path launches one kernel per operator, which creates that extra work.

The operator that replaces those five rows is one kernel with the whole chain in registers, and TORCH_LIBRARY in the same file puts it on torch.ops.day88:

__global__ void fusedChainKernel(const float* __restrict__ x,
                                 const float* __restrict__ r,
                                 float* __restrict__ out, size_t n) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (i < n) {
        float v = scaleBias(x[i]);
        v = gelu(v);
        v = v + r[i];
        out[i] = clampRange(v);
    }
}

The last pass in the program adds nothing to the arithmetic and everything to the timeline:

def nvtx_pass(eager, fused, x, residual):
    """The same two blocks again, under NVTX ranges and nothing else.

    torch.profiler already knows the operator names. Nsight Systems does not:
    without these ranges an nsys capture of a model is one long row of
    anonymous cudaLaunchKernel bars, which is exactly the problem day 41
    opened with. The other route is
    torch.autograd.profiler.emit_nvtx(), which puts a range around every
    autograd operation automatically; it is finer and noisier than this.
    """
    with torch.no_grad():
        for _ in range(RUNS):
            with torch.cuda.nvtx.range("eager-block"):
                eager(x, residual)
            with torch.cuda.nvtx.range("fused-block"):
                fused(x, residual)
    torch.cuda.synchronize()

Day 41 used four NVTX ranges around a PageRank iteration. torch.cuda.nvtx.range wraps the same nvtxRangePush and nvtxRangePop pair with a Python context manager on top, and it is what makes an nsys capture of a model readable. The profiler tells you which operator; the timeline tells you what the host was doing while the device waited.

Results

Verified on a Tesla T4, driver 580.173.02, CUDA 12.6 and torch 2.13.0+cu126. The clean standalone run passed correctness, produced both torch profiler tables and Chrome traces, and exited 0. The fused result had no elements outside rtol 1e-5, atol 1e-6; its worst absolute difference was 4.768e-07, and 46167 outputs reached the clamp.

measurement eager fused
whole block, CUDA events, mean of ten 7.7794 ms 6.6554 ms
elementwise CUDA time in clean profiler 1.5118 ms across five kernels 0.4104 ms in one kernel

The elementwise eager figure is the sum of the two aten::add rows (685.724 us), aten::clamp (276.447 us), aten::gelu (275.710 us) and aten::mul (273.887 us). Against the fused kernel's 410.397 us, that is 3.68x, almost exactly the 11-to-3 byte ratio. The whole block improved only 1.1689x because its GEMM was unchanged.

The NVTX CSV reports a shorter average host range for the fused block than for the eager block. The Nsight-wrapped rerun made torch profiler warn that CUPTI already had a subscriber.

Its torch tables are therefore not evidence. Nsight Systems did produce the .nsys-rep, NVTX summary and CUDA kernel summary, while the clean standalone run produced the torch profiler tables above.

The NVTX PushPop durations measure host submission between range boundaries, not GPU completion, because the ranges contain asynchronous launches and synchronization happens after the loop.

Four things this should show, and what would kill each

  1. Refuted as written. The device kernel was volta_sgemm_128x128_tn, but torch 2.13 also attributed 6.243 ms of self CUDA time to aten::addmm and an Unrecognized row. The three rows describe the same GEMM activity, so this version's table does not preserve the predicted zero-self wrapper shape.
  2. Held. Nsight Systems ranked volta_sgemm_128x128_tn first by a wide margin over the five eager elementwise kernel classes. The clean profiler likewise measured the GEMM at 6.242 ms against 1.5118 ms for the eager chain.
  3. Held. The chain improved 3.68x against a 3.67x byte ratio, while the whole block improved only 1.1689x.
  4. Held with one tooling caveat. The clean torch profiler run and the Nsight Systems capture both worked without root or performance-counter permission. They did not collect concurrently: the nsys-wrapped rerun's second CUPTI subscriber warning is why the two profiler results come from separate runs.

A profile ranks rows, not proposed fixes. Focus on rows whose arithmetic or composition you control.

Run it yourself

Use a GPU with compute capability 7.5 or newer. PyTorch's CUDA support matrix lists Turing (7.5) for the 12.6, 13.0, and 13.2 builds of 2.13.0 (https://github.com/pytorch/pytorch/blob/main/RELEASE.md , checked 2026-09-01).

torch.utils.cpp_extension compiles fused_chain.cu with the nvcc on PATH, so its CUDA major version must match the PyTorch wheel. The extension uses that nvcc for its JIT build. Compiler Explorer cannot host this run because PyTorch is a 2.5 GB install and the runner has a 20-second compile limit.

Exercise

Profile a small model of your own, at least a Linear and three elementwise operators in a row. Find the slowest chain of memory-bound operators in the self_cuda_time_total table, replace it with one fused operator built the way day 87 built yours, and profile it again.

Time: 40 to 60 minutes. Submit: the two key_averages tables, the chain ratio, the whole-model ratio, and one sentence on why they differ.

Check: report two ratios rather than comparing raw milliseconds across GPUs. Count the floats per element your chain moves before and after, then divide the two summed self CUDA times. A chain ratio well under the byte ratio means something else bounds the kernel and day 49's roofline is the next tool; a chain ratio at the byte ratio with a model ratio near one means you fused something that was never the problem.

The harness checks your output against the eager chain at rtol 1e-5, atol 1e-6 and prints the first disagreeing index, so a faster wrong answer fails.

Hint 1

Two small rows can still repeat the same memory traffic. Identify what they share that neither row shows by itself.

Hint 2

Count loads and stores per element, not flops. Then count the same thing once the intermediates never leave a register. The ratio of those two counts is what your measurement has to be compared against.

Solution

The chain speeds up by roughly its byte ratio and the model speeds up by much less because the unchanged GEMM still dominates part of the run. Eleven floats per element against three gives the chain a factor of about three. The model gain depends on the fraction of its original time spent in that chain, as described by Amdahl's law.

Libraries provide common operators, but a model-specific scale_bias_gelu_residual_clamp chain may have no fused entry point. The small profiler rows can therefore identify useful custom-kernel work.

Pitfalls

You added up the CUDA total column and got more time than the run contains. Total is inclusive of children, so a dispatch chain counts its kernel once per level. Sort and sum self_cuda_time_total instead; the hook at the top of this page is what the other column looks like.

Your top row is cudaMalloc, or the first kernel is ten times slower than the rest. You profiled the first iteration, so the table holds allocator growth, autotuning and lazy module loading rather than your model. Run the model once before opening the profiler, as the program here does, with the warm-up discipline day 9 set.

Five different aten::add calls collapsed into one row and you cannot tell which one is expensive. The profiler averages by operator name unless you give it more to key on. Pass record_shapes=True and read key_averages(group_by_input_shape=True).

nsys shows your model as one anonymous row of launches. Either the program has no NVTX ranges or the capture dropped them: -t cuda alone ignores NVTX even when the calls are compiled in, and the trace list has to say -t cuda,nvtx. Day 41 has the same pitfall for the same reason.

You assumed the profiler needs root because ncu does. It does not.

Nsight Compute reads hardware performance counters and fails for a normal user with The user running <tool_name/application_name> does not have permission to access NVIDIA GPU Performance Counters or the Hardware Event System on the target device (https://developer.nvidia.com/nvidia-development-tools-solutions-err_nvgpuctrperm-permission-issue-performance-counters , checked 2026-09-01). torch.profiler and nsys both trace through CUPTI's activity interface, which uses a different permission.

Your operator launched on the wrong stream. A kernel on the default stream inside a model running on another one is a race that appears only under load. Pass at::cuda::getCurrentCUDAStream() as the fourth launch argument, as fused_chain.cu does.

Go deeper

Next

Day 89 compares CUDA, Triton, cuTile, CUTLASS, and other GPU tools. The choice starts with the cost shown by the profile. Day 90 then checks whether each capstone should have used a library call.