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

NVIDIA published a step-by-step tutorial demonstrating how to debug, benchmark, and optimize CUDA code using modern tools such as Compute Sanitizer, Nsight Systems, NVTX, CUB, and CCCL, applied to an RGB-to-grayscale image processing pipeline with median filtering. The tutorial covers six incremental improvements, including adopting the CCCL API, using CUB's optimized algorithms, managing GPU memory with pooled containers, and parallelizing work with per-thread streams, with code available on GitHub and Google Colab.

NVIDIA CUDA remains the foundation of GPU-accelerated computing, powering everything from scientific simulations to large-scale AI training. But writing correct, maintainable, and performant CUDA code can be challenging: memory bugs hide in plain sight, performance bottlenecks can be invisible without the right instrumentation, and hand-rolled GPU algorithms rarely match the efficiency of optimized libraries. Fortunately, the modern CUDA toolchain has matured significantly, and many of these challenges now have straightforward solutions. In this blog post, we will walk through the tools NVIDIA offers to debug, benchmark, and improve your code. With only small line changes each time, we are going to make the example code safer, easier to maintain, and faster. Across six incremental steps, this post will cover: - How to easily find indexing bugs by adopting the modern CCCL API and Compute Sanitizer - How to improve Nsight Systems benchmarks with NVTX - How to use CUB’s optimized algorithms at the block and device level - How to manage GPU memory through pooled containers - How to speed up host-to-device transfers with pinned containers - How to parallelize GPU work by giving each thread its own stream and asynchronous transfers As a companion to this blog post, we provide the code https://github.com/NVIDIA/accelerated-computing-hub/tree/main/resources/blogs/modern cuda cpp blogpost/code steps and the option to run on Google Colab https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/resources/blogs/modern cuda cpp blogpost/code steps/modern cuda toolbox tutorial.ipynb . Starting point: An image processing pipeline example From an input stream of red, green, and blue images, start by transferring the data from the CPU to the GPU. Then convert those from RGB to grayscale. Then, for each 32 by 32 pixel tile in the image, compute the median by sorting the pixels and selecting the middle value. Finally, copy the median of each tile back to the CPU. The base code example Below is the full starting code. Each step in this post improves on it. define CUDA CHECK ERROR call do { \ cudaError t err = call; \ if err = cudaSuccess { \ std::cerr << "CUDA error in " << FILE << " at line " << LINE << ": " \ << cudaGetErrorString err << std::endl; \ std::exit EXIT FAILURE ; \ } \ } while 0 // Alias for an image pixel using pixel t = uint8 t; // Kernel converting the red, green and blue images into a single gray image global void computeRGBToGray const pixel t d image r, const pixel t d image g, const pixel t d image b, pixel t d image gray, int width, int height { // Compute the thread global index in the grid const int x = threadIdx.x + blockIdx.x blockDim.x; const int y = threadIdx.y + blockIdx.y blockDim.y; // Boundary check selecting only threads within the image boundary if x < width && y < height { // Compute the thread index in the image const int i = x + y width; // Convert from rgb to grayscale and store the result in global memory d image gray i = static cast<pixel t 0.299f d image r i + 0.587f d image g i + 0.114f d image b i ; } } // Kernel computing the median of each tile in the grayscale image template <int TILE WIDTH, int HISTO SIZE global void computeMedian pixel t d image gray, pixel t d median, int width, int height { // Compute the thread global index in the grid const int x = threadIdx.x + blockIdx.x blockDim.x; const int y = threadIdx.y + blockIdx.y blockDim.y; // Boundary check selecting only threads within the image boundary if x < width && y < height return; // Allocate the shared memory in which we will store the tile shared pixel t tile TILE WIDTH TILE WIDTH ; // Compute the thread index in the image const int index = x + y width; // Load the tile's grayscale value from global memory into shared memory tile index = d image gray index ; // Synchronize to make sure all threads have loaded their data syncthreads ; // Sort the tile array using a single threaded bubble sort if threadIdx.x == 0 && threadIdx.y == 0 { for int i = 0; i < TILE WIDTH TILE WIDTH; ++i for int j = i + 1; j < TILE WIDTH TILE WIDTH; ++j if tile i tile j cuda::std::swap tile i , tile j ; // Each thread block stores the median, found in the middle index after sorting, in the global median array const int medianIndex = TILE WIDTH TILE WIDTH / 2; d median blockIdx.x + blockIdx.y gridDim.x = tile medianIndex ; } } int main { // Define all the example constants constexpr auto TILE WIDTH = 32; constexpr auto HISTO SIZE = 256; constexpr auto NB TILE X = 250; constexpr auto NB TILE Y = NB TILE X; constexpr auto IMAGE LENGTH = TILE WIDTH NB TILE X; constexpr auto IMAGE SIZE = IMAGE LENGTH IMAGE LENGTH; constexpr auto NB IMAGES = 3; constexpr auto INIT VALUE = 4; // Allocate the CPU memory to store the images tiles medians and for the red, green, blue and grayscale images std::vector<std::vector<pixel t h images r NB IMAGES, std::vector<pixel t IMAGE SIZE, 4 ; std::vector<std::vector<pixel t h images g NB IMAGES, std::vector<pixel t IMAGE SIZE, 4 ; std::vector<std::vector<pixel t h images b NB IMAGES, std::vector<pixel t IMAGE SIZE, 4 ; std::vector<std::vector<pixel t h images gray NB IMAGES, std::vector<pixel t IMAGE SIZE, 0 ; std::vector<std::vector<pixel t h medians NB IMAGES, std::vector<pixel t NB TILE X NB TILE Y ; // Run the image processing pipeline for each image, in parallel pragma omp parallel for for int i = 0; i < NB IMAGES; ++i { pixel t d image r, d image g, d image b, d image gray, d median; // Allocate the GPU memory for each container CUDA CHECK ERROR cudaMalloc &d image r, IMAGE SIZE sizeof pixel t ; CUDA CHECK ERROR cudaMalloc &d image g, IMAGE SIZE sizeof pixel t ; CUDA CHECK ERROR cudaMalloc &d image b, IMAGE SIZE sizeof pixel t ; CUDA CHECK ERROR cudaMalloc &d image gray, IMAGE SIZE sizeof pixel t ; CUDA CHECK ERROR cudaMalloc &d median, NB TILE X NB TILE Y sizeof pixel t ; // Copy the memory of each container from CPU to GPU CUDA CHECK ERROR cudaMemcpy d image r, h images r i .data , IMAGE SIZE sizeof pixel t , cudaMemcpyHostToDevice ; CUDA CHECK ERROR cudaMemcpy d image g, h images g i .data , IMAGE SIZE sizeof pixel t , cudaMemcpyHostToDevice ; CUDA CHECK ERROR cudaMemcpy d image b, h images b i .data , IMAGE SIZE sizeof pixel t , cudaMemcpyHostToDevice ; // Launch a GPU kernel to convert the RGB images to grayscale dim3 blockSize TILE WIDTH, TILE WIDTH ; dim3 gridSize cuda::ceil div IMAGE LENGTH, blockSize.x , cuda::ceil div IMAGE LENGTH, blockSize.y ; computeRGBToGray<<<gridSize, blockSize d image r, d image g, d image b, d image gray, IMAGE LENGTH, IMAGE LENGTH ; CUDA CHECK ERROR cudaGetLastError ; // Launch the GPU kernel to compute the median of every tile in the image computeMedian<TILE WIDTH, HISTO SIZE <<<gridSize, blockSize d image gray, d median, IMAGE LENGTH, IMAGE LENGTH ; CUDA CHECK ERROR cudaGetLastError ; // Copy the GPU median memory back to the CPU CUDA CHECK ERROR cudaMemcpy h medians i .data , d median, NB TILE X NB TILE Y sizeof pixel t , cudaMemcpyDeviceToHost ; // Free the GPU memory CUDA CHECK ERROR cudaFree d image r ; CUDA CHECK ERROR cudaFree d image g ; CUDA CHECK ERROR cudaFree d image b ; CUDA CHECK ERROR cudaFree d image gray ; CUDA CHECK ERROR cudaFree d median ; } return 0; } This code starts by defining two kernels: computeRGBToGray loads the values of the red, green, and blue input images to convert and write them to the grayscale output image. computeMedian computes the median of each tile of an input grayscale image. Each thread block loads the tile from global memory to shared memory. Then, a single thread is used to sort the array and write the value found at the middle index, corresponding to the median, in the global output median array. In the main, after defining the constants used for the example, the CPU memory is allocated for each image and for the medians. The image processing pipeline is then run for each of the three images, in parallel, using OpenMP. The pipeline starts by allocating the required memory on the GPU before transferring the data from the CPU to the GPU. The two kernels to convert from RGB to grayscale and to compute the median are launched afterward. Finally, we copy back the median results to the CPU before freeing the memory. This code has several flaws that will be addressed, step by step. 1. Compute Sanitizer and the CCCL API: Find bugs easily and write safer code Let’s start by running the code. code steps$ ./build/0 base error example CUDA error in 0 base error example.cu at line 105: an illegal memory access was encountered While the code has some error checking in it, when you get an error message like this “illegal memory access,” you should start by using compute-sanitizer to investigate further. Using Compute Sanitizer https://developer.nvidia.com/compute-sanitizer , the NVIDIA functional correctness checking suite, we can directly identify a bug that can be hard to spot: bash $ compute-sanitizer ./build/0 base error example ========= COMPUTE-SANITIZER ========= Invalid shared write of size 1 bytes ========= at void computeMedian< int 32, int 256 unsigned char , unsigned char , int, int +0x170 in 0 base error example.cu:55 ========= by thread 0,3,0 in block 20,0,0 ========= Access at 0x6440 is out of bounds Running the above directly shows that the code suffers from an out-of-bound shared write at line 55 of 0 base error example.cu . tile index = d image gray index ; The line incorrectly loads data in shared memory, using a global index. Since shared memory is defined at the thread block level, we need to change the indexing. To avoid indexing mistakes, a new API was introduced in CCCL to distinguish global and block-level indexing. To use it, you first need to launch your kernel using the new cuda::launch API: auto config = cuda::make config cuda::block dims ... , cuda::grid dims ... ; cuda::launch stream, config, kernel name<decltype config , input Then, using the new indexing API inside the kernel: template <typename Configuration global void kernel name Configuration config, ... { // Retrieve and expand each global index const auto x, y, z = cuda::gpu thread.index cuda::grid, config ; // Retrieve the block index structure containing block idx.x, .y, .z const auto block idx = cuda::gpu thread.index cuda::block, config ; } Without using compute-sanitizer or the new API, this mistake could also have been directly spotted by using cuda::std::span or its n-dimensional variant cuda::std::mdspan instead of raw pointers. cuda::std::span and cuda::std::mdspan are non-owning views over contiguous memory and are useful to abstract the exact container away. Accessing data through spans is safer than through raw pointers in part because in debug mode, out-of-bounds accesses will trigger an assertion. The kernel should be updated as: // Alias for a 2-dimensional mdspan template <typename T using span 2d = cuda::std::mdspan<T, cuda::std::dims<2 ; template <typename Configuration global void computeMedian ..., span 2d<const pixel t d image gray, ... If you run the code with these changes you’ll get something like the following: $ ./build/1 span libcudacxx/include/cuda/std/ mdspan/mdspan.h:436: operator : block: 16,0,0 , thread: 0,30,0 Assertion mdspan: operator out of bounds access failed. Memory accesses in shared memory should also be protected by using a cuda::shared memory mdspan as in the following snippet: shared pixel t shared TILE WIDTH TILE WIDTH ; cuda::shared memory mdspan tile 2d shared, TILE WIDTH, TILE WIDTH ; Now you can run and everything should execute properly without errors. Using the new launch API and its indexing mechanism, spans over raw pointers, and compute-sanitizer, out-of-bounds accesses either do not happen or are caught right away. For more information on compute-sanitizer, see Efficient CUDA Debugging: How to Hunt Bugs with NVIDIA Compute Sanitizer https://developer.nvidia.com/blog/debugging-cuda-more-efficiently-with-nvidia-compute-sanitizer/ . 2. Nsight Systems and NVTX: Benchmark your code properly Now the code is bug-free, it is ready to be benchmarked using NVIDIA Nsight Systems. It allows you to visualize the program timeline: know when each function is called and for how long. To make the timeline visualization easier, we wrap every interesting code section using NVTX: void image compute ... { // NVTX range for the scope of the whole function nvtx3::scoped range fun scope "Image compute" ; // NVTX range that is pushed and then popped for a specific code section nvtxRangePushA "Kernel median" ; // Launch the GPU kernel to compute the median of every tile in the image ... // Pop the range at the end of the specific code section nvtxRangePop ; } Which yields the following result: In the GPU hardware CUDA HW section of the profiler output in Figure 3, above, it is reported that the GPU is mainly busy with kernels 98.5% of the GPU time while the memory operations only take 1.5% of the GPU time. Of the two kernels, the one computing the medians is taking the majority of the runtime with 2.1 seconds for each image see the yellow box on the right, where the stats for computeMedian are shown, and the elapsed time is 2.142s . From the CPU thread section, we see the image computation takes 6.8 seconds in total, with most of the time being spent on computing the medians for the three grayscale images. We now know the first operation to optimize in order to have the greatest impact. For more information on Nsight Systems, see Optimizing CUDA Memory Transfers with NVIDIA Nsight Systems https://developer.nvidia.com/blog/optimizing-cuda-memory-transfers-with-nsight-systems/ . For more information on NVTX, see CUDA Pro Tip: Generate Custom Application Profile Timelines with NVTX https://developer.nvidia.com/blog/cuda-pro-tip-generate-custom-application-profile-timelines-nvtx/ . 3. CUB: Express algorithms directly on the GPU When dealing with common algorithms, writing custom kernels is error-prone and will most likely result in an inefficient implementation. Whenever possible, both for device-side patterns and for in-kernel primitives, it is recommended to use CUB. CUB is the NVIDIA parallel algorithm library shipped through CCCL. It exposes highly optimized routines at several granularities: device-wide cub::Device , block-level cub::Block , and warp-level cub::Warp . For the RGB to grayscale step, we can replace the custom kernel with cub::DeviceTransform::Transform . It applies a user-provided function to a tuple of input iterators and writes the result to an output iterator, executing on the GPU: // Use CUB to convert the RGB images to grayscale cub::DeviceTransform::Transform cuda::std::make tuple d image r, d image g, d image b , // inputs d image gray, // output IMAGE SIZE, // size host device pixel t r, pixel t g, pixel t b // functor { return static cast<pixel t 0.299f r + 0.587f g + 0.114f b ; }, stream ; For the median, programming a parallel block-level sort by hand is complex and slow. Instead, we directly leverage CUB’s block-level radix sort inside the kernel: // Declare and allocate the storage for CUB BlockRadixSort using BlockRadixSort = cub::BlockRadixSort<... ; shared typename BlockRadixSort::TempStorage temp storage; // Load the tile's grayscale value from global memory pixel t thread keys 1 ; thread keys 0 = d image gray y, x ; // Perform the thread-block-level radix sort BlockRadixSort temp storage .Sort thread keys ; // Select the thread found at the middle index // Write its value which is, after sorting, the median, in the global median array if block idx.x == TILE WIDTH / 2 && block idx.y == TILE WIDTH / 2 d median grid block idx.y, grid block idx.x = thread keys 0 ; Following this change, we benchmark again using Nsight Systems: The time required to compute the median is now only 773 microseconds again, look at the elapsed time in the computeMedian yellow pop-out image , 2717x faster. The overall time to compute all three images is now 635 milliseconds, 10x faster. If we reassess the current bottleneck: the time spent on memory allocations represents around 83% of the total image compute runtime. This can be greatly improved. 4. Pooled Memory Containers: Convenient and faster memory management Allocating GPU memory using cudaMalloc can have unexpected negative effects: leaks by forgetting to call cudaFree and costly memory operations in critical parts of your code. Instead, we recommend using CCCL’s asynchronous memory containers, cuda::device buffer . Like C++ std::vector , the memory is automatically deallocated once the container goes out of scope. Additionally, a memory pool backs the buffer, so repeated allocations and deallocations do not pay the full cost of cudaMalloc / cudaFree every time. To use the GPU memory containers, we update the code accordingly: // Resource to handle the GPU memory allocations cuda::device memory pool ref device resource = cuda::device default memory pool cuda::device ref{0} ; // Explained at a later stage, unimportant for now cuda::stream stream{cuda::device ref{0}}; // Allocate the GPU memory using uninitialized containers cuda::device buffer<pixel t d image r = cuda::make buffer<pixel t stream, device resource, IMAGE SIZE, cuda::no init ; ... After this change, we analyze the timeline again: The time spent on memory allocation is now almost nonexistent; the time it takes to compute an image has improved by 2.6x. The GPU time is now memory-dominated. Almost all the time to compute all images is spent on first copying the red, green and blue, for the three images, from CPU to GPU. It is possible to greatly speed up the host-to-device memory transfers. 5. Pinned memory: Faster host-to-device memory transfers CPU data allocations are pageable by default, which the GPU cannot access directly. The CUDA driver must first allocate a temporary page-locked, or pinned, host array, copy the host data to the pinned array, and then transfer the data from the pinned array to device memory. When it is known in advance that CPU memory will be copied to the GPU, it is advised to directly allocate using pinned memory. CCCL exposes a pinned-memory host container, cuda::host buffer , that can be built through the cuda::make pinned buffer factory: // Allocate the CPU memory to store the image tiles, medians, and for the red, green, blue, and grayscale images // Those CPU containers, contrary to std::vector, are allocated using pinned memory std::vector<cuda::host buffer<pixel t h images r NB IMAGES, cuda::make pinned buffer<pixel t stream, IMAGE SIZE, ... ; ... Following those changes, we can benchmark again: The time it takes to do the host-to-device memory transfers has been significantly reduced; it now only takes 25 ms to compute all images, 10x faster. For more information on pinned memory, see How to Optimize Data Transfers in CUDA C/C++ https://developer.nvidia.com/blog/how-optimize-data-transfers-cuda-cc/ . One surprising behavior might have caught the eye of some readers since the beginning: Though we are using different CPU threads, all operations memory and kernels are being executed sequentially on the GPU. Let’s fix it. 6. Streams: Parallelize operations on the GPU By default, all operations kernels, memory allocations, or transfers are launched on what we call the default stream: it can be viewed as a queue of tasks the GPU needs to execute in order. In this example, we need a stream for each image/thread. CCCL provides cuda::stream , an owning self-managed version of CUDA streams. It can simply be constructed inside the parallel for loop, so each OpenMP thread gets its own queue of GPU work. To efficiently leverage streams, we also need to use the asynchronous API: each GPU operation launched by the CPU should not be waited upon until completion. To saturate the GPU, each CPU thread should launch as many operations as possible, as fast as possible, without waiting for them to first complete. Kernels and CUB device calls are already asynchronous by default and are launched on the passed stream. To launch asynchronous copies between the host and the device, we use the new CCCL cuda::copy bytes API. It is advised in any modern CUDA code never to rely on the default stream and to always rely on streams. We update the code accordingly: A dedicated init stream is used for the initial allocations of the pinned host buffers. Each iteration of the parallel for loop now owns its own cuda::stream for the computation pipeline: // Stream used for initial host buffer allocations cuda::stream init stream{cuda::device ref{0}}; ... // Allocate the host pinned buffers on init stream: std::vector<cuda::host buffer<pixel t h images r NB IMAGES, cuda::make pinned buffer<pixel t init stream, IMAGE SIZE, ... ; ... // Sync before launching operations on another stream: init stream.sync ; pragma omp parallel for for int i = 0; i < NB IMAGES; ++i { // One different stream per thread cuda::stream stream{cuda::device ref{0}}; ... // GPU buffer allocations using the stream owned by each thread cuda::device buffer<pixel t d image r = cuda::make buffer<pixel t stream, device resource, IMAGE SIZE, cuda::no init ; ... // Copy the memory of each container from CPU to GPU asynchronously using the stream owned by each thread cuda::copy bytes stream, h images r i , d image r ; ... // Use CUB to convert the RGB images to grayscale asynchronously using the per thread stream cub::DeviceTransform::Transform ..., stream.get ; // Launch the GPU kernel to compute the median of every tile in the image using the per thread stream cuda::launch stream, ... ; // Copy the GPU median memory back to the CPU cuda::copy bytes stream, d median, h medians i ; ... // To make sure the copy bytes is finished before accessing results on the host stream.sync ; } Following those changes, we can take a final look at the timeline: We now have a complete overlap between our kernels and memory copies. The final duration to compute all three images following all our improvements is 23 milliseconds, starting from 6.8 seconds. Your turn Using the CUDA Developer’s Toolbox, we made the code safer, easier to maintain and faster. No low-level optimizations were used, yet the code is 300x faster. Try out this code https://github.com/NVIDIA/accelerated-computing-hub/tree/main/resources/blogs/modern cuda cpp blogpost/code steps yourself, and run on Google Colab https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/resources/blogs/modern cuda cpp blogpost/code steps/modern cuda toolbox tutorial.ipynb if you like. We have also built a full class to learn how to use those tools in detail. It is freely available on YouTube https://www.youtube.com/playlist?list=PL5B692fm6--vWLhYPqLcEu6RF3hXjEyJr alongside links to practice on Google Colab https://github.com/NVIDIA/accelerated-computing-hub/tree/main/tutorials/cuda-cpp .