现代 CUDA 工具箱的实际应用：分步优化演练-优快云博客

NVIDIA CUDA 仍然是 GPU 加速计算的基础，为从科学模拟到大规模 AI 训练的一切提供支持。

但是，编写正确、可维护和高性能的 CUDA 代码可能具有挑战性：显存错误隐藏在显而易见的视野中，如果没有正确的仪器，性能瓶颈可能是不可见的，手动滚动的 GPU 算法几乎无法与优化库的效率相匹配。幸运的是，现代 CUDA 工具链已经显著成熟，其中许多挑战现在都有简单直接的解决方案。

在本博文中，我们将介绍 NVIDIA 提供的用于调试、基准测试和改进代码的工具。每次只需更改少量行，我们就能使示例代码更安全、更易于维护和更快。

本博文将介绍六个渐进步骤：

如何通过采用现代 CCCL API 和 Compute Sanitizer 轻松找到索引错误
如何使用 NVTX 改进 Nsight Systems 基准测试
如何在块和设备级别使用 CUB 的优化算法
如何通过池化容器管理 GPU 显存
如何使用固定容器加速主机到设备的传输
如何通过为每个线程分配自己的流和异步传输来并行化 GPU 工作

作为本博文的配套内容，我们提供了在 Google Colab 上运行的代码和选项。

起点：图像处理工作流示例

从红色、绿色和蓝色图像的输入流开始，将数据从 CPU 传输到 GPU。然后将 RGB 转换为灰度。

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.

图 1. 图像处理工作流的第一部分。首先，将 RGB 图像从主机复制到设备，然后将其从 RGB 转换为灰度

然后，对于图像中的每个 32 x 32 像素图块，通过对像素进行排序并选择中间值来计算中值。最后，将每个图块的中位数复制回 CPU。

图 2. 图像处理工作流的第二部分：将每个灰度图像分为 32 x 32 像素的图块，对每个图块中的像素进行排序以选择其中间值，然后将中间值从设备复制到主机

基础代码示例

以下是完整的起始代码。本博文中的每个步骤都会改进。

#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;
}

此代码首先定义两个内核：

computeRGBToGray 加载红色、绿色和蓝色输入图像的值，以将其转换并写入灰度输出图像。
computeMedian 计算输入灰度图像每个图块的中位数。每个线程块将图块从全局内存加载到共享内存。然后，使用单个线程对数组进行排序，并写入在全局输出中值数组中的中间索引 (对应中值) 处找到的值。

在正文中，定义示例中使用的常量后，CPU 内存将分配给每张图像和中间值。

然后，使用 OpenMP 为这三张图像分别并行运行图像处理工作流。工作流首先在 GPU 上分配所需显存，然后再将数据从 CPU 传输到 GPU。之后启动两个核函数，将 RGB 转换为灰度，并计算中值。最后，在释放内存之前，将中间值结果复制回 CPU。

此代码有几个问题需要逐步解决。

1. Compute Sanitizer 和 CCCL API：轻松发现错误并编写更安全的代码

我们先来运行代码。

code_steps$ ./build/0_base_error_example
CUDA error in 0_base_error_example.cu at line 105: an illegal memory access was encountered

虽然代码中有一些错误检查，但当您收到类似“非法内存访问”的错误消息时，您应该首先使用 computer – sanitizer 进行进一步调查。

使用 NVIDIA 功能正确性检查套件 Compute Sanitizer，我们可以直接识别难以发现的错误：

$ 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

运行上述操作直接表明，代码在 0_base_error_example.cu 的第 55 行存在超出限制的共享写入问题。

tile[index] = d_image_gray[index];

该行使用全局索引错误地加载共享内存中的数据。由于共享内存是在线程块级别定义的，因此我们需要更改索引。为避免索引错误，我们在 CCCL 中引入了新的 API，以区分全局索引和块级索引。要使用它，您首先需要使用新的 cuda::launch API 启动内核：

auto config = cuda::make_config(cuda::block_dims(...), cuda::grid_dims(...));
cuda::launch(stream, config, kernel_name<decltype(config)>, input)

然后，在内核中使用新的索引 API：

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);
}

如果不使用 computing – sanitizer 或新 API，也可以使用 cuda::std::span 或其 n 维变体 cuda::std::mdspan 而不是原始指针来直接发现此错误。cuda::std::span 和 cuda::std::mdspan 是连续内存的非拥有视图，有助于抽象出确切的容器。通过 span 访问数据比通过原始指针更安全，部分原因是在调试模式下，越界访问将触发断言。

内核应更新为：

// 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, ...)

如果您在运行代码时进行了这些更改，您将获得以下内容：

$ ./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.

还应使用 cuda::shared_memory_mdspan 保护共享内存中的内存访问，如下文代码段所示：

__shared__ pixel_t shared[TILE_WIDTH * TILE_WIDTH];
cuda::shared_memory_mdspan tile_2d(shared, TILE_WIDTH, TILE_WIDTH);

现在，您可以运行，一切都应正常执行，而不会出现错误。

使用新的启动 API 及其索引机制，跨原始指针，而计算清理器、越界访问要么不发生，要么立即被捕获。有关 compute-sanitizer 的更多信息，请参阅Efficient CUDA Debugging：How to Hunt Bug with NVIDIA Compute Sanitizer (高效 CUDA 调试：如何使用 NVIDIA Compute Sanitizer 追踪错误) 。

2. Nsight Systems 和 NVTX：正确地对代码进行基准测试

现在，该代码没有错误，可以使用 NVIDIA Nsight Systems 进行基准测试。它允许您可视化程序时间轴：了解每个函数何时被调用以及调用时长。

为简化时间轴可视化，我们使用 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();
}

这会生成以下结果：

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.

图 3. 包含 NVTX 标注的初始 Nsight Systems 时间轴。中值计算内核在配置文件中占主导地位，几乎占据所有 GPU 活动，并且占据了 6.8 秒图像处理持续时间的大部分

在以上图 3 分析器输出的 GPU 硬件 (CUDA HW) 部分中，据悉 GPU 主要忙于处理内核 (占 GPU 时间的 98.5%) ，而显存运算仅占 GPU 时间的 1.5%。

在两个内核中，计算中值的内核会占用大部分运行时间，且每张图像需要 2.1 秒 (请参阅右侧的黄色框，其中显示了 computeMedian 的统计信息，运行时间为 2.142 秒) 。

在 CPU (线程) 部分，我们发现图像计算总共需要 6.8 秒，其中大部分时间用于计算三个灰度图像的中间值。

我们现在知道要优化的第一个运算，以便产生最大的影响。有关 Nsight Systems 的更多信息，请参阅使用 NVIDIA Nsight Systems 优化 CUDA 内存传输。有关 NVTX 的更多信息，请参阅CUDA Pro 提示：使用 NVTX 生成自定义应用程序配置文件时间轴。

3. CUB：直接在 GPU 上使用 Express 算法

在处理常见算法时，编写自定义内核容易出错，并且极有可能导致实现效率低下。无论对于设备端模式还是内核内基元，建议尽可能使用 CUB。

CUB 是通过 CCCL 提供的 NVIDIA 并行算法库。它以多种粒度呈现高度优化的例程：设备级 (cub::Device*) 、块级 (cub::Block*) 和线程束级 (cub::Warp*) 。

对于 RGB 到灰度的步骤，我们可以将自定义内核替换为 cub::DeviceTransform::Transform。它将用户提供的函数应用于输入迭代器的元组，并将结果写入输出迭代器，在 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);

对于中间值，手动对并行块级排序进行编程既复杂又缓慢。相反，我们直接在内核中利用 CUB 的块级基数排序：

// 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];

作出此更改后，我们再次使用 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.

图 4. 将自定义算法替换为 CUB 后的 Nsight Systems 时间轴。中值计算被缩短到微秒，整个工作流大约需要 635 毫秒，而内存分配是主要成本

计算中值所需的时间现在只有 773 微秒 (同样，在 computeMedian yellow 弹出式图像中查看消耗时间) ，速度提高了 2717 倍。现在，计算全部三张图像的总时间为 635 毫秒，速度提高了 10 倍。

如果我们重新评估当前瓶颈：在内存分配上花费的时间约占总图像计算运行时的 83%。

这可以得到很大的改进。

4. 池化内存容器：便捷、快速的内存管理

使用 cudaMalloc 分配 GPU 显存可能会产生意想不到的负面影响：忘记调用 cudaFree 而造成泄露，以及在代码的关键部分进行昂贵的内存操作。

相反，我们建议使用 CCCL 的异步内存容器 cuda::device_buffer。与 C++ std::vector 类似，一旦容器超出范围，内存将自动取消分配。

此外，内存池支持缓冲区，因此重复分配和取消分配无法支付每次 cudaMalloc/ cudaFree 的全部成本。

为使用 GPU 显存容器，我们会相应地更新代码：

// 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);
...

更改后，我们将再次分析时间轴：

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.

图 5. 采用池化设备内存容器后的 Nsight Systems 时间轴。分配用度基本消除，主机到设备的内存传输成为下一个明显瓶颈

现在，分配内存所花费的时间几乎不复存在；计算图像所需的时间缩短了 2.6 倍。

GPU 时间现在以内存为主。在计算所有图像时，几乎所有时间都花在从 CPU 到 GPU 的三个图像的红色、绿色和蓝色上。

可以大幅加速主机到设备的内存传输。

5. 固定显存：更快的主机到设备显存传输

CPU 数据分配默认可分页，GPU 无法直接访问。CUDA 驱动程序必须首先分配一个临时的锁页或固定主机阵列，将主机数据复制到固定阵列，然后将数据从固定阵列传输到设备内存。

当预先知道 CPU 内存将复制到 GPU 时，建议使用固定内存直接分配。

CCCL 提供了一个固定内存主机容器 cuda::host_buffer，可通过 cuda::make_pinned_buffer 工厂进行构建：

// 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, ...));
...

作出这些更改后，我们可以再次进行基准测试：

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.

图 6. 在固定内存中分配主机数据后的 Nsight Systems 时间轴；主机到设备的传输更短，所有三个镜像都在大约 25 毫秒内完成，尽管它们的操作仍然在默认流上依次执行

主机到设备内存传输所需的时间已显著缩短；现在计算所有图像只需 25 毫秒，速度提高了 10 倍。

有关固定内存的更多信息，请参阅如何在 CUDA C/ C++ 中优化数据传输。

从一开始，一些读者就注意到了一种惊人的行为：

虽然我们使用不同的 CPU 线程，但所有操作 (内存和内核) 都是在 GPU 上依次执行的。

我们来解决这个问题。

6. 流：在 GPU 上并行处理运算

默认情况下，所有操作 (核函数、内存分配或传输) 均在我们所说的默认流上启动：默认流可被视为 GPU 按顺序执行的任务队列。

在本示例中，我们需要为每个图像/ 线程创建流。CCCL 提供 cuda::stream，这是一个拥有自管理版本的 CUDA 流。它可以在并行的 for 循环中简单构建，因此每个 OpenMP 线程都会获得自己的 GPU 工作队列。

为了有效利用流，我们还需要使用异步 API：CPU 启动的每个 GPU 操作不应等待完成。为了使 GPU 饱和，每个 CPU 线程应尽可能快地启动尽可能多的操作，而无需等待它们首次完成。默认情况下，核函数和 CUB 设备调用已经是异步的，并且会在通过的流上启动。要在主机和设备之间启动异步副本，我们使用新的 CCCL cuda::copy_bytes API。

建议在任何现代 CUDA 代码中永远不要依赖默认流，并始终依赖流。

我们会相应地更新代码：

专用的 init_stream 用于固定主机缓冲区的初始分配。现在，并行 for 循环的每个迭代都为计算管道拥有自己的 cuda::stream：

// 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();
}

作出这些更改后，我们可以最终查看时间轴：

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.