--- 格式版本: 2 标题: "The Modern CUDA Toolbox in Practice: A Step-by-Step Optimization Walkthrough" 原文链接: "https://developer.nvidia.com/blog/the-modern-cuda-toolbox-in-practice-a-step-by-step-optimization-walkthrough/" 发布日期: "2026-09-02" 发布时间校准状态: "found" 发布时间需复核: "否" 发布时间来源: "rule:configured_publication_date_rule" 发布时间证据: "nvidia-developer-blog-published-meta text:scrape: Published: 2026-09-02" 发布时间校准原因: "信源发布日期识别规则直接确认发布时间" 发布时间校准置信度: "high" 发布时间候选数量: 1 发布时间严格候选数量: 1 发布时间原页读取状态: "source template page reused from URL open" 发布时间未找到原因: "" 发布时间校准时间: "2026-09-04T14:49:37+08:00" 发布时间仲裁状态: "skipped" 发布时间仲裁尝试次数: 0 发布时间仲裁耗时毫秒: 0 发现时间: "2026-09-04T14:48:38+08:00" 入库时间: "2026-09-04T06:49:38.220Z" 来源平台: "固定入口" 搜索渠道: "fixed_url" 搜索词: "https://developer.nvidia.com/blog/" 匹配关键词: - "GPU" - "AI" - "performance" 相关厂家: - "NVIDIA" - "Google" 相关专家: [] 内容类型: "网页" 抓取工具: "Free Fetch + Defuddle" 清洗工具: "Defuddle Markdown + Defuddle/Readability 正文提取" 原始附件: [] AI优质: "否" AI打分: 38 AI分档: "非优质" AI质检状态: "不通过" AI打分理由: "正文主线是基于图像处理示例的CUDA代码调试与性能优化教程，介绍Compute Sanitizer、CCCL/CUB、内存池、锁页内存和多流并发，并给出6.8秒降至23毫秒的示例结果；并非超节点、AI Rack或机架级基础设施。来源为NVIDIA官方开发者博客，正文完整且代码与测试数据充分。固定知识库未显示同篇内容，但未命中不能证明首次发布；本文也未正式发布新的机架级硬件、协议标准、生产级基础设施基准或规模部署，且没有客户、量产或交付信号。命中“教程与运维选型”硬否决项，当前页面不具备超节点情报库准入价值。" AI质检模型: "gpt-5.6-sol" AI质检时间: "2026-09-08T16:42:39+08:00" AI主题相关性: 2 AI来源权威性: 15 AI新颖性: 5 AI技术细节: 6 AI商业部署信号: 0 AI完整性: 10 AI评分提示词版本: "v17-精简生产版" AI评分提示词SHA256: "48fb9777f386026761b4873eaff30807694fb11e9b352d7c69bf2dfde750cc7d" AI评分知识库版本: "knowledge_base_v1-20260819+runtime.91" AI评分知识库SHA256: "63412978a77e0c992a0c417bf6a25f0ef83df5c342c2a67607a122c0f8bd250e" AI评分知识库检索词: "[\"NVIDIA\",\"Google\",\"GPU\",\"performance\",\"https://developer.nvidia.com/blog/\",\"NPU\",\"CUDA\",\"GPU-accelerated\",\"CCCL\",\"API\",\"NVTX\",\"CUB\"]" AI评分知识库命中: "[{\"id\":\"runtime-9f43a79e9c23b1f329d8a71a\",\"title\":\"AWS and NVIDIA to Deliver 2 Million Additional GPUs and Next-Generation Infrastructure for Agentic and Physical AI\",\"sourceType\":\"ai_excellent_article\",\"time\":\"2026-08-26\",\"matchedTerms\":[\"NVIDIA\",\"GPU\",\"performance\",\"CUDA\",\"GPU-accelerated\",\"API\"],\"rank\":-14.50697457573956},{\"id\":\"runtime-35f284698f34b54313c05e16\",\"title\":\"7 Restore LLM Inference Capacity in Seconds with Shadow Engine Recovery in NVIDIA Dynamo | August 2026 NVIDIA Dynamo's shadow engine recovery feature, presented by Maksim Khadkevich, Vikram Sharma Mailthody, Mohammed Abdulwahhab, and Schwinn Saereesitthipitak, reduces LLM inference recovery time from 283 seconds to 7.3 seconds by utilizing a preinitialized shadow engine on the same GPUs as the active engine. It leverages NVIDIA CUDA, GPU Memory Service (GMS), and Dynamic Resource Allocation (DRA\",\"sourceType\":\"ai_excellent_article\",\"time\":\"\",\"matchedTerms\":[\"NVIDIA\",\"GPU\",\"NPU\",\"CUDA\",\"API\"],\"rank\":-12.924988060001422},{\"id\":\"july-correct-0089\",\"title\":\"Setting a World Record for MoE Pre-Training on NVIDIA GB300 NVL72\",\"sourceType\":\"labeled_article\",\"time\":\"2026-07\",\"matchedTerms\":[\"NVIDIA\",\"GPU\",\"performance\",\"https://developer.nvidia.com/blog/\",\"CUDA\",\"API\"],\"rank\":-11.852388987267318},{\"id\":\"runtime-3dcabc270e938e0c3b6d5245\",\"title\":\"Nvidia wants to bypass the CPU with an open-source AI storage overhaul\",\"sourceType\":\"ai_excellent_article\",\"time\":\"2026-08-06\",\"matchedTerms\":[\"NVIDIA\",\"Google\",\"GPU\",\"performance\",\"API\"],\"rank\":-11.522608023852857},{\"id\":\"runtime-bf4039a0342d37545e9459a2\",\"title\":\"Most Neoclouds Suck At Security\",\"sourceType\":\"ai_excellent_article\",\"time\":\"2026-08-30\",\"matchedTerms\":[\"NVIDIA\",\"Google\",\"GPU\",\"performance\",\"https://developer.nvidia.com/blog/\",\"CUDA\",\"API\"],\"rank\":-10.907963548939733}]" 采集批次: "2026年9月4日14点35分46秒" 采集批次ID: "20260904-143546-06b8cfc0" 去重键: "https://developer.nvidia.com/blog/the-modern-cuda-toolbox-in-practice-a-step-by-step-optimization-walkthrough" --- 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: 1. How to easily find indexing bugs by adopting the modern CCCL API and Compute Sanitizer 2. How to improve Nsight Systems benchmarks with NVTX 3. How to use CUB’s optimized algorithms at the block and device level 4. How to manage GPU memory through pooled containers 5. How to speed up host-to-device transfers with pinned containers 6. 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. ![Diagram divided into a blue host section on the left and a green device section on the right. Three stacks of red, green, and blue channel images appear above a CPU icon. Colored lines connect them to three overlapping grayscale images above a GPU icon, representing the host-to-device transfers and RGB-to-grayscale conversion.](https://developer-blogs.nvidia.com/wp-content/uploads/2026/08/image1-7.webp) Figure 1. The first part of the image-processing pipeline. First, copy the RGB images from host to device, then convert them 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. ![Diagram with a blue host section containing an RGB image stack and CPU icon on the left, and a green device section containing a GPU icon on the right. Above the GPU, a grayscale image is divided by red grid lines into tiles. An example sequence of pixel values is shown being sorted into ascending order, with the middle value highlighted, and the resulting median represented in a smaller output image. A line back to the CPU represents copying the medians to host memory.](https://developer-blogs.nvidia.com/wp-content/uploads/2026/08/image4-6.webp) Figure 2. The second part of the image-processing pipeline: Divide each grayscale image into 32-by-32-pixel tiles, sort the pixels in each tile to select its median, then copy the median values from device to host ### 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(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 __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> h_images_r(NB_IMAGES, std::vector(IMAGE_SIZE, 4)); std::vector> h_images_g(NB_IMAGES, std::vector(IMAGE_SIZE, 4)); std::vector> h_images_b(NB_IMAGES, std::vector(IMAGE_SIZE, 4)); std::vector> h_images_gray(NB_IMAGES, std::vector(IMAGE_SIZE, 0)); std::vector> h_medians(NB_IMAGES, std::vector(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<<>>(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<<>>(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: ``` $ 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, input) ``` Then, using the new indexing API inside the kernel: ``` template __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 using span_2d = cuda::std::mdspan>; template __global__ void computeMedian(..., span_2d 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: ![Nsight Systems timeline showing GPU hardware activity above CPU thread and NVTX rows. Three long computeMedian kernel intervals occupy almost the entire GPU timeline, while RGB-to-grayscale kernels and memory transfers appear as narrow intervals. The enclosing Image Compute NVTX range spans approximately 6.8 seconds, revealing median computation as the primary bottleneck.](https://developer-blogs.nvidia.com/wp-content/uploads/2026/08/image3-7.webp) Figure 3. The initial Nsight Systems timeline with NVTX annotations. Median-computation kernels dominate the profile, accounting for nearly all GPU activity and most of the 6.8-second image-processing duration 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(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: ![Nsight Systems timeline after the CUB optimization. GPU kernel activity is compressed into short intervals near the end of the displayed range, while the CPU and CUDA API rows contain a much wider memory-allocation interval. The Image Compute NVTX range covers roughly 635 milliseconds, showing that cudaMalloc-related work now dominates after the kernels were accelerated.](https://developer-blogs.nvidia.com/wp-content/uploads/2026/08/image5-5.webp) Figure 4. The Nsight Systems timeline after replacing the custom algorithms with CUB. Median computation is reduced to microseconds and the complete pipeline takes about 635 milliseconds, leaving memory allocation as the dominant cost 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 d_image_r = cuda::make_buffer(stream, device_resource, IMAGE_SIZE, cuda::no_init); ... ``` After this change, we analyze the timeline again: ![Nsight Systems timeline containing repeated groups of memory-copy and short kernel intervals for three images. The previously dominant allocation region is absent. Wide cudaMemcpy intervals and matching Memory Copy In NVTX ranges occupy most of the remaining pipeline, while the RGB-to-grayscale and median kernels are comparatively brief.](https://developer-blogs.nvidia.com/wp-content/uploads/2026/08/image6-5.webp) Figure 5. The Nsight Systems timeline after adopting pooled device-memory containers. Allocation overhead is nearly eliminated, making host-to-device memory transfers the next visible bottleneck 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> h_images_r(NB_IMAGES, cuda::make_pinned_buffer(stream, IMAGE_SIZE, ...)); ... ``` Following those changes, we can benchmark again: ![Nsight Systems timeline showing three consecutive image-processing groups on one GPU stream. Each group contains a shortened Memory Copy In range followed by RGB-to-grayscale and median kernels and a Memory Copy Out range. The Image Compute NVTX range lasts approximately 25 milliseconds, but the three groups remain serialized rather than overlapping.](https://developer-blogs.nvidia.com/wp-content/uploads/2026/08/image7-4.webp) Figure 6. The Nsight Systems timeline after allocating host data in pinned memory; host-to-device transfers are much shorter and all three images complete in about 25 milliseconds, although their operations still execute sequentially on the default stream 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> h_images_r(NB_IMAGES, cuda::make_pinned_buffer(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 d_image_r = cuda::make_buffer(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: ![Enlarged Nsight Systems GPU timeline with multiple CUDA stream rows. Three sequences of host-to-device copies, kernels, and device-to-host copies are staggered across the stream rows and overlap in time. The highlighted duration is approximately 23 milliseconds, demonstrating concurrent execution of the three image pipelines.](https://developer-blogs.nvidia.com/wp-content/uploads/2026/08/image2-9.webp) Figure 7. The final Nsight Systems timeline after assigning a CUDA stream to each image; memory transfers and kernels overlap across three streams, reducing the total processing time from 6.8 seconds to about 23 milliseconds 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).