The GPU toolchain — profilers & kernel DSLs.

Two halves of the toolchain. Profilers tell you what is slow: Nsight Systems for the timeline (is the GPU even busy? are you launch-bound? CPU-bound?), Nsight Compute for one kernel's guts (Roofline, occupancy, stall reasons). Kernel producers let you write fast code without raw PTX: CUTLASS (templated CUDA), TensorRT (inference compiler), Triton (a Python DSL), and Helion (a higher-level DSL that compiles down to Triton). Pick the highest-level tool that gives you the control you actually need.

Article 2.1 made one demand over and over: measure. This article is about the instruments. It splits cleanly into two parts. First, the profilers — because you cannot optimize what you cannot see, and GPU performance is famously counter-intuitive. Second, the kernel-authoring tools — because hand-writing optimized CUDA is hard, and a stack of increasingly high-level abstractions now lets you get most of the performance for a fraction of the effort.

Part 1 — Profiling: the two Nsight tools

NVIDIA's profiling story is two complementary tools, and using the right one for the right question saves enormous time. The mental model: Nsight Systems zooms out to the whole application; Nsight Compute zooms in to one kernel.

Nsight Systems — the timeline

Nsight Systems is a low-overhead, system-wide profiler. It shows a timeline of CPU threads, CUDA API calls, memory transfers, and kernel executions all aligned on one clock. Its job is to answer the structural questions that dwarf any single-kernel tuning:

• Is the GPU actually busy? A surprising amount of "slow GPU code" is a GPU sitting idle while the CPU prepares work or a synchronous copy blocks. The timeline shows gaps immediately.
• Are you launch-bound? Thousands of tiny kernels with gaps between them means launch overhead dominates — the fix is fusion or CUDA Graphs, not kernel tuning.
• Are host↔device copies serializing with compute? If memcpys and kernels run back-to-back instead of overlapping, you need asynchronous copies on separate streams.
• Are your streams actually concurrent? The timeline reveals false serialization from an accidental synchronization point.

# capture a timeline, then open report.nsys-rep in the GUI
nsys profile -o report --stats=true ./my_app
# quick text summary of where time went
nsys stats report.nsys-rep

Nsight Compute — the kernel microscope

Once Nsight Systems tells you which kernel matters, Nsight Compute dissects it. It re-runs the kernel with heavy instrumentation and produces the exact metrics article 2.1 asked for:

• Roofline chart — your kernel plotted against the memory and compute ceilings, so you instantly see memory-bound vs compute-bound and how much headroom remains.
• Achieved occupancy vs theoretical, plus the limiting resource (registers / shared memory / block size).
• Memory-transaction efficiency — the coalescing health check; low values point straight at strided access.
• Warp-stall reasons — a breakdown of why warps are not issuing (waiting on memory, on a barrier, on a long-latency instruction). This is the single most actionable view: it names your bottleneck.
• Source-to-SASS correlation — maps assembly hot-spots back to your source lines.

# profile one kernel with the full "detailed" metric set
ncu --set full -o kernel_report ./my_app
# or target a specific kernel by name, limit to first launch
ncu -k myKernel -c 1 --set full ./my_app

Nsight Compute is heavyweight — it can run a kernel many times to gather all counters — so you point it at the one or two kernels Nsight Systems flagged, not the whole app. The combination is the entire measured-optimization loop from article 2.1, made real.

Part 2 — Producing fast kernels

Now the other half: writing the kernels. There is a ladder of abstraction here, from raw CUDA C++ at the bottom to Python DSLs at the top. Higher on the ladder means less code and less hardware detail to manage, usually at some cost in peak control. The skill is choosing the lowest-effort rung that still meets your performance and flexibility needs.

CUTLASS — templated CUDA for matrix kernels

CUTLASS (CUDA Templates for Linear Algebra Subroutines) is NVIDIA's open-source library of C++ template building blocks for high-performance GEMM and convolution. It implements all the techniques from this series — hierarchical tiling across thread blocks, warps, and threads; double-buffered async copies; Tensor-Core mma instructions — as composable, well-tuned components. Its killer feature is the fusible epilogue: you can append your own element-wise operations (bias, activation, scaling) directly onto a GEMM so the result never round-trips through HBM. When cuBLAS's fixed set of operations is not enough but you do not want to write a GEMM from scratch, CUTLASS is the answer. The newer CuTe abstraction (tensors + layouts) makes expressing these tilings considerably cleaner.

TensorRT — the inference compiler

TensorRT is a different kind of tool: an inference optimizer and runtime. You hand it a trained model (via ONNX or a framework export) and it produces a heavily optimized "engine" for a specific GPU. Its main optimizations:

• Layer/operator fusion — automatically fuses conv+bias+activation and similar chains into single kernels (exactly the fusion principle from article 2.1).
• Precision calibration — runs the model in FP16 or INT8 (with calibration to preserve accuracy), or FP8 on Hopper.
• Kernel auto-tuning — benchmarks multiple kernel implementations for each layer on your actual GPU and picks the fastest.
• Memory planning — reuses activation buffers to shrink the memory footprint.

TensorRT is the standard last-mile step for deploying a model on NVIDIA hardware when latency and throughput matter. You write no kernels; you accept a longer build step in exchange for a runtime engine that is often 2–5× faster than the eager framework. TensorRT-LLM is the LLM-specialized variant, adding paged KV cache, in-flight batching, and the attention/decoding optimizations from the DeepSeek and KOG articles.

Triton — write kernels in Python

Triton (originally from OpenAI) is the abstraction that changed who can write GPU kernels. It is a Python DSL where you express computation in terms of blocks/tiles rather than individual threads. You write the tile logic; the Triton compiler handles the painful parts automatically — memory coalescing, shared-memory allocation, and much of the occupancy tuning. A Triton matmul or fused softmax is a few dozen readable lines and routinely matches or beats hand-tuned CUDA.

import triton, triton.language as tl

@triton.jit
def add_kernel(x_ptr, y_ptr, out_ptr, n, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)   # this block's tile of indices
    mask = offs < n                            # guard the tail
    x = tl.load(x_ptr + offs, mask=mask)       # coalescing handled for you
    y = tl.load(y_ptr + offs, mask=mask)
    tl.store(out_ptr + offs, x + y, mask=mask)

Notice what is absent: no threadIdx, no explicit shared memory, no __syncthreads(). You think in tiles; the compiler lowers tiles to warps and threads. This is why Triton is now the backend that PyTorch's torch.compile generates fused kernels into — it hits the sweet spot of control and productivity. The mental model from Phase 1 still fully applies: you choose tile sizes for coalescing and reuse, and Triton's autotuner sweeps configurations to find a good occupancy. You understand the why; Triton handles the how.

Helion — a higher-level DSL above Triton

Helion is a newer Python-embedded DSL (from the PyTorch ecosystem) that sits one rung above Triton. Where Triton still asks you to manage tiles and program IDs explicitly, Helion lets you write kernels in a more PyTorch-like, loop-oriented style and compiles down to Triton, autotuning the tiling and launch configuration for you. The pitch is "ML kernels with minimal boilerplate": you express the math close to how you would in NumPy/PyTorch, and Helion handles the decomposition into tiles, the autotuning search, and the lowering — giving much of Triton's performance with even less of the manual tile bookkeeping. It is the current frontier of the "write less, get fast kernels" trend, and a sign of where kernel authoring is heading: the human specifies intent and data layout, the compiler searches the implementation space.

Choosing a tool: a decision guide

Situation	Reach for	Why
Standard dense matmul / conv	cuBLAS / cuDNN	Already optimal; do not reinvent.
GEMM with custom fused epilogue	CUTLASS	Templated GEMM + your own epilogue, no HBM round-trip.
Deploy a trained model fast	TensorRT / TensorRT-LLM	Auto fusion, precision, kernel auto-tuning.
Custom fused op, want productivity	Triton	Python tiles, compiler handles coalescing/occupancy.
ML kernel, minimal boilerplate	Helion	PyTorch-like, autotunes, compiles to Triton.
Exotic shape, need every last %	raw CUDA / CuTe	Full control when nothing higher suffices.
Any of the above, but slow	Nsight Systems → Compute	Measure before and after every change.

How the tools fit the optimization loop

Tying Phase 2 together: the workflow from article 2.1 maps directly onto these tools. You write the kernel at the highest comfortable rung (Triton/Helion, or call CUTLASS/TensorRT). You profile the application with Nsight Systems to confirm the GPU is busy and the kernel matters. You analyze that kernel with Nsight Compute to get its Roofline position, occupancy, and stall reasons. You apply the matching technique family — often just by changing a tile size in Triton or enabling a precision in TensorRT — and re-measure. The tools turn the abstract method into a tight, repeatable loop.

FAQ

Do I need Nsight Compute if I only write Triton?

Yes — and Triton integrates with it. Triton can dump the generated PTX/SASS, and Nsight Compute profiles the resulting kernel like any other. Tile-size choices in Triton have exactly the occupancy and coalescing consequences Nsight Compute measures, so the profiler is how you tune Triton autotuning intelligently rather than blindly.

Is Triton NVIDIA-only?

Triton started CUDA-focused but has growing support for AMD (ROCm) and other backends. That portability is part of its appeal — the same tile-level kernel can target multiple vendors, whereas raw CUDA is locked to NVIDIA. The abstraction layer is also a portability layer.

When does TensorRT not help?

For very dynamic models (highly variable shapes, lots of control flow) the engine-build assumptions break down, and for training (TensorRT is inference-only). It also locks you to a specific GPU and TensorRT version per engine, so it adds deployment rigidity. For stable inference shapes on fixed hardware, it is excellent; for research iteration, it is overhead.

Helion vs Triton — should I just always use Helion?

Helion is younger and more abstract, so it is great for standard ML kernels with minimal effort but gives you less direct control than Triton when you need to hand-tune an unusual pattern. A reasonable default: start in Helion (or torch.compile), drop to Triton when you need to control the tiling yourself, drop to CUTLASS/CUDA only when even that is not enough.

What's the difference between cuBLAS and cuDNN?

cuBLAS is for dense linear algebra (matrix multiply, etc.); cuDNN is for deep-learning primitives specifically (convolutions, pooling, normalization, attention). Both are NVIDIA's pre-tuned closed libraries and both are hard to beat on their home turf.

Takeaways

• Profile top-down: Nsight Systems (whole-app timeline, is the GPU busy?) before Nsight Compute (one kernel's Roofline, occupancy, stalls).
• No kernel tuning fixes an idle GPU — fix application structure (overlap, batching, fewer launches) first.
• CUTLASS for fusible GEMMs, TensorRT for deploying trained models, cuBLAS/cuDNN for standard ops.
• Triton lets you write tile-level kernels in Python; the compiler handles coalescing and occupancy mechanics.
• Helion sits above Triton — PyTorch-like, autotuned, compiles down to Triton.
• Higher-level DSLs remove boilerplate, not the need to understand the hardware. Phase 1 knowledge is what makes you good at them.
• Pick the highest rung that meets your needs; drop down only when it leaves performance on the table.

Phase 2 is complete: you can measure a kernel and write a fast one. Phase 3 changes scale entirely — from one GPU to many. Article 3.1 covers multi-GPU programming and model parallelism (Megatron-LM and the synchronization/access overheads that dominate at scale), and article 3.2 covers the datacenter: scheduling, virtualization, and multi-tenant sharing.

References & further reading

• NVIDIA Nsight Systems — the system-wide timeline profiler.
• NVIDIA Nsight Compute — kernel-level analysis with built-in Roofline and occupancy.
• CUTLASS — CUDA templates for linear algebra; CuTe layouts.
• TensorRT & TensorRT-LLM — inference compiler and runtime.
• OpenAI Triton — the Python kernel DSL.
• Helion — higher-level Python DSL that compiles to Triton.