The Modern CUDA Toolbox in Practice: A Step-by-Step Optimization Walkthrough

This NVIDIA Developer post is a vendor walkthrough of the modern CUDA toolbox on a toy pipeline: RGB images to GPU, grayscale, 32×32 tile medians, copy back. The baseline hits an illegal memory access. Compute Sanitizer finds an out-of-bounds shared write: a global index used as a shared-memory index. The fix path is CCCL's cuda::launch indexing (grid vs block coordinates), plus cuda::std::mdspan and shared_memory_mdspan that assert on overflow.

Nsight Systems with NVTX shows median kernels at 2.1 seconds per image and 6.8 seconds end-to-end, 98.5% kernel time. Replacing the RGB kernel with cub::DeviceTransform and the bubble-sort median with cub::BlockRadixSort drops median compute to 773 microseconds — 2,717× — and the pipeline to 635 milliseconds. Pooled cuda::device_buffer cuts remaining time 2.6× by killing cudaMalloc churn. Pinned cuda::host_buffer makes host-to-device copies about 10× faster, to ~25 milliseconds for three images. Giving each OpenMP thread its own cuda::stream and asynchronous copy_bytes overlaps the work to a final 23 milliseconds, about 300× versus 6.8 seconds. Colab, code, and a YouTube class are linked.

Vendor tutorial, and unusually useful because it starts from a crash. The 2,717× number is real for this example and not a promise about your model. The example is rigged to make hand-rolled bubble sort look as bad as it is. That is still the lesson: CUB exists; your single-threaded shared-memory sort does not.

The sequence of bottlenecks is the actual curriculum. Bugs, then the algorithm, then allocators, then pageable copies, then the default stream serializing OpenMP threads that thought they were parallel. Each step is a one-API change. That is the CCCL pitch: stop being a CUDA cowboy.

Signal Desk's read: this post is Nvidia teaching the 2026 stack — launch API, spans, pools, pinned host buffers, per-thread streams — because the old cudaMalloc/cudaMemcpy folklore still ships in production. The 300× is a commercial for CCCL. It is also a true description of how much performance was hiding in not being careful. Teams training or serving on Blackwell who still bubble-sort a tile should be embarrassed, not inspired.

Compute Sanitizer as step zero should be policy. 'Illegal memory access' without it is superstition.

CUDA C++ has grown a standard library. CCCL (Thrust, CUB, libcudacxx) is that library. This walkthrough is how Nvidia wants new GPU code to look: fewer raw pointers, more containers, library algorithms, and a profiler open.