奇链 奇链
← 文章 / 编程开发
NVIDIA 开发者博客 1小时前 · 2026-09-03 05:07:38 · 2 阅读

实战现代 CUDA 工具箱：逐步优化指南

NVIDIA CUDA 仍是 GPU 加速计算的基石，驱动着从科学模拟到大规模 AI 训练的各类应用。 但编写正确、可维护且高性能的 CUDA 代码颇具挑战：内存 bug 往往隐蔽难察，若缺少合适的分析工具，性能瓶颈也极易被忽略，而手工编写的 GPU 算法也很难达到优化库的效率。幸运的是，现代 CUDA 工具链已日趋成熟，许多难题如今都有了简洁的解决方案。 本文我们将逐步梳理 NVIDIA 提供的调试、基准测试和性能优化工具。每次只需改动少量行，就能让示例代码更安全、更易维护、更快。 通过六个递进步骤，本文将介绍：
  1. 如何借助现代 CCCL API 与 Compute Sanitizer 快速定位索引 bug
  2. 如何利用 NVTX 增强 Nsight Systems 性能分析
  3. 如何在 block 级和 device 级使用 CUB 优化算法
  4. 如何通过池化容器管理 GPU 内存
  5. 如何通过 pinned 容器加速主机到设备的传输
  6. 如何通过为每个线程分配独立 stream 及异步传输来并行化 GPU 工作

作为本文的配套资源，我们提供了代码，并支持在Google Colab上运行。

起点：图像处理流水线示例

输入为红、绿、蓝三通道图像流，首先将数据从 CPU 传输至 GPU，随后将其从 RGB 转换为灰度图。

然后，对于图像中每个 32×32 像素的图块，通过排序像素并选取中间值来计算中位数。最后，将每个图块的中值复制回 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.
图 2. 图像处理流程的第二阶段：将每张灰度图划分为 32×32 像素的图块，对每个图块内的像素排序并选取中值，然后将中值从设备复制回主机

基础代码示例

下面是完整的起始代码。本文的每个优化步骤都建立在此基础上。

```cpp #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) // 图像像素的类型别名 using pixel_t = uint8_t; // 将 RGB 三通道图像合并为单通道灰度图像的 kernel __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) { // 计算线程在网格中的全局索引 const int x = threadIdx.x + blockIdx.x * blockDim.x; const int y = threadIdx.y + blockIdx.y * blockDim.y; // 边界检查，仅处理图像范围内的线程 if (x < width && y < height) { // 计算线程在图像中的索引 const int i = x + y * width; // 将 RGB 转换为灰度并存入全局内存 d_image_gray[i] = static_cast(0.299f * d_image_r[i] + 0.587f * d_image_g[i] + 0.114f * d_image_b[i]); } } // 计算灰度图像每个 tile 中位数的 kernel template __global__ void computeMedian(pixel_t *d_image_gray, pixel_t *d_median, int width, int height) { // 计算线程在网格中的全局索引 const int x = threadIdx.x + blockIdx.x * blockDim.x; const int y = threadIdx.y + blockIdx.y * blockDim.y; // 边界检查，仅处理图像范围内的线程 if (!(x < width && y < height)) return; // 分配用于存储 tile 的共享内存 __shared__ pixel_t tile[TILE_WIDTH * TILE_WIDTH]; // 计算线程在图像中的索引 const int index = x + y * width; // 将 tile 的灰度值从全局内存加载到共享内存 tile[index] = d_image_gray[index]; // 同步，确保所有线程都已完成数据加载 __syncthreads(); // 使用单线程冒泡排序对 tile 数组排序 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]); // 每个线程块将排序后位于中间位置的中位数存入全局中位数数组 const int medianIndex = (TILE_WIDTH * TILE_WIDTH) / 2; d_median[blockIdx.x + blockIdx.y * gridDim.x] = tile[medianIndex]; } } int main() { // 定义示例常量 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; // 分配 CPU 内存，用于存储各图像 tile 的中位数以及 RGB 和灰度图像 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)); // 并行处理每张图像的图像处理流水线 #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; // 为每个容器分配 GPU 内存 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))); // 将各容器的内存从 CPU 复制到 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)); // 启动 GPU kernel，将 RGB 图像转为灰度 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()); // 启动 GPU kernel，计算图像中每个 tile 的中位数 computeMedian<<>>(d_image_gray, d_median, IMAGE_LENGTH, IMAGE_LENGTH); CUDA_CHECK_ERROR(cudaGetLastError()); // 将 GPU 上的中位数内存复制回 CPU CUDA_CHECK_ERROR(cudaMemcpy(h_medians[i].data(), d_median, (NB_TILE_X * NB_TILE_Y) * sizeof(pixel_t), cudaMemcpyDeviceToHost)); // 释放 GPU 内存 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; } ``` ```cpp

这段代码首先定义了两个内核：

  • computeRGBToGray 从红、绿、蓝三张输入图像读取像素值，转换为灰度并写入输出图像。
  • computeMedian 计算输入灰度图像每个图块的中值。每个线程块将图块从全局内存加载到共享内存，随后由一个线程对数组排序，将中值（位于中间索引处）写入全局输出中值数组。

在 main 函数中，定义了示例所用的常量后，分别为每幅图像和中值数组分配了 CPU 内存。

接着使用 OpenMP 并行对三张图像依次运行图像处理流水线。流水线首先在 GPU 上分配所需内存，然后将数据从 CPU 传输到 GPU，之后启动 RGB 转灰度和计算中值的两个内核。最后将中值结果复制回 CPU 并释放内存。

这段代码存在若干缺陷，接下来将逐步修复。

1. Compute Sanitizer 与 CCCL API：轻松发现 bug，写出更安全的代码

我们先来运行这段代码。

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

虽然代码中包含了一些错误检查，但遇到"illegal memory access"这类错误信息时，应首先使用 compute-sanitizer 进行深入排查。

借助 Compute Sanitizer（NVIDIA 功能正确性检查套件），我们可以直接定位到一个不易察觉的 bug：

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

代码在 shared memory 中错误地使用了全局索引来加载数据。由于 shared memory 是在 thread block 层面定义的，我们需要修改索引方式。为了避免索引错误，CCCL 中引入了新的 API 来区分全局索引和 block 级索引。使用它的第一步是用新的 cuda::launch API 启动 kernel：

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

然后在 kernel 内部使用新的索引 API：

template <typename Configuration>
__global__ void kernel_name(Configuration config, ...) {
    // 获取并展开每个全局索引
    const auto [x, y, z] = cuda::gpu_thread.index(cuda::grid, config);

    // 获取 block 索引结构（包含 block_idx.x、.y、.z）
    const auto block_idx = cuda::gpu_thread.index(cuda::block, config);
}

如果不使用 compute-sanitizer 或新 API，也可以用 cuda::std::span 或其多维变体 cuda::std::mdspan 替代裸指针来直接发现这类错误。cuda::std::span 和 cuda::std::mdspan 是对连续内存的非所有权视图，可用于抽象掉具体容器类型。通过 span 访问数据比裸指针更安全，因为在调试模式下越界访问会触发断言。

kernel 应更新为：

// 二维 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.

shared memory 中的内存访问也应通过 cuda::shared_memory_mdspan 加以保护，如下所示：

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

现在运行代码，一切应该能正常执行而不出错。

借助新的启动API及其索引机制、span（包裹原始指针），以及compute-sanitizer，越界访问要么根本不会发生，要么会被立即捕获。更多compute-sanitizer的信息，请参阅高效CUDA调试：如何使用NVIDIA Compute Sanitizer追踪Bug。

2. Nsight Systems 与 NVTX：对代码进行规范基准测试

现在代码已无Bug，可以借助NVIDIA Nsight Systems进行基准测试。它能可视化程序时间线，让你清楚每个函数何时被调用、执行了多久。

为了让时间线可视化更清晰，我们用NVTX将各个关键代码段包裹起来：

void image_compute(...)
{
  // 整个函数的NVTX作用域
  nvtx3::scoped_range fun_scope("Image compute");

 // 对特定代码段进行push/pop的NVTX范围
  nvtxRangePushA("Kernel median");

  // 启动GPU内核计算图像每个瓦片的像素中值
  ...

  // 在特定代码段结束时pop该范围
  nvtxRangePop();
}

得到的结果如下：

Nsight Systems时间线：上方为GPU硬件活动，下方依次是CPU线程和NVTX行。三段较长的computeMedian内核区间几乎占满整个GPU时间线，而RGB转灰度内核和内存传输则表现为狭窄区间。外层Image Compute NVTX范围跨度约6.8秒，表明中值计算是主要瓶颈。
图3. 带有NVTX标注的初始Nsight Systems时间线。中值计算内核占据主导，几乎贡献了全部GPU活动量和6.8秒图像处理时长的绝大部分

从图3的GPU硬件（CUDA HW）分析结果可以看到，GPU主要忙于内核计算（占GPU时间的98.5%），而内存操作仅占1.5%。

其中，中值计算内核占用了绝大部分运行时间，每个图像耗时2.1秒（见右侧黄色框，显示computeMedian的统计信息，耗时2.142秒）。

从CPU（线程）部分可见，图像处理总共耗时6.8秒，大部分时间花费在计算三个灰度图像的中值上。

至此，我们已经明确了最值得优先优化的操作。如需了解更多关于Nsight Systems的信息，请参阅使用NVIDIA Nsight Systems优化CUDA内存传输。关于NVTX的更多信息，请参阅CUDA技巧：使用NVTX生成自定义应用分析时间线。

3. CUB：在GPU上直接表达算法

处理常见算法时，手写自定义内核既容易出错，也很难得到高效实现。因此无论是对设备端模式还是核内原语，都应优先使用 CUB。 CUB 是 NVIDIA 并行算法库，随 CCCL 一同提供。它在多个粒度上暴露了高度优化的例程：设备级（cub::Device*）、线程块级（cub::Block*）和 warp 级（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 进行基准测试：
使用 CUB 优化后，Nsight Systems 时间线上的 GPU 内核活动被压缩到显示范围末尾的短暂区间内，而 CPU 与 CUDA API 行中则出现更宽泛的内存分配区间。Image Compute NVTX 区间覆盖约 635 毫秒，表明在内核加速之后，与 cudaMalloc 相关的工作已成为主要开销。
图 4. 用 CUB 替换自定义算法后的 Nsight Systems 时间线。中位数计算时间降至微秒级，完整流水线耗时约 635 毫秒，内存分配成为主要开销

中位数的计算时间现已降至仅 773 微秒（再次参考 computeMedian 黄色弹出框中的经过时间），速度提升 2717 倍。计算全部三张图片的总耗时现在为 635 毫秒，整体提速 10 倍。

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

这可以得到很大改善。

4. 池化内存容器：便捷且高效的内存管理

使用 cudaMalloc 分配 GPU 内存可能带来意想不到的负面影响：例如忘记调用 cudaFree 导致内存泄漏，或在代码关键部分产生高昂的内存操作开销。

为此，我们建议使用 CCCL 的异步内存容器 cuda::device_buffer。与 C++ 的 std::vector 类似，容器一旦超出作用域，其内存将被自动释放。

此外，该缓冲区由内存池支撑，因此重复的分配与释放操作无需每次都承担完整的 cudaMalloc / cudaFree 开销。

为使用这些 GPU 内存容器，我们对代码进行如下更新：

// 用于管理 GPU 内存分配的资源
cuda::device_memory_pool_ref device_resource = cuda::device_default_memory_pool(cuda::device_ref{0});

// 稍后解释，当前不重要
cuda::stream stream{cuda::device_ref{0}};

// 使用未初始化的容器分配 GPU 内存
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 时间现在主要由内存传输主导。计算所有图像的耗时几乎全花在第一步——将三张图像的 RGB 通道数据从 CPU 复制到 GPU。

可以通过优化显著提升主机到设备的内存传输速度。

5. Pinned 内存：加速主机到设备传输

CPU 数据分配默认是页面可移动的（pageable），GPU 无法直接访问。CUDA 驱动必须先分配一块临时的页锁定（pinned，即固定）主机数组，把数据复制到 pinned 数组，再从 pinned 数组传输到设备内存。

若提前知道 CPU 内存将复制到 GPU，建议直接使用 pinned 内存进行分配。

CCCL 提供了一个固定内存主机容器 `cuda::host_buffer`，可通过 `cuda::make_pinned_buffer` 工厂函数创建： ```cpp // 分配 CPU 内存以存储图像瓦片、中值以及红、绿、蓝和灰度图像 // 这些 CPU 容器与 std::vector 不同，使用固定内存分配 std::vector> h_images_r(NB_IMAGES, cuda::make_pinned_buffer(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 倍。 有关固定内存的更多信息，请参阅 How to Optimize Data Transfers in CUDA C/C++。 细心的读者可能从一开始就注意到一个值得关注的现象： 虽然我们使用了不同的 CPU 线程，但所有操作（内存传输和内核）都在 GPU 上顺序执行。 让我们来解决这个问题。

6. 流：在 GPU 上并行化操作

默认情况下，所有操作（kernel、内存分配或数据搬运）都会在 default stream 上启动，可视为 GPU 需按序执行的任务队列。 本例中每个图像（或线程）均需独立流。CCCL 提供了 cuda::stream，这是一种自管理且带所有权声明的 CUDA stream。只需在并行 for 循环内直接构造，每个 OpenMP 线程便各拥一条独立的 GPU 任务队列。 要高效利用 stream，必须搭配异步 API：CPU 发出 GPU 操作后不应阻塞等待。为跑满 GPU，每个 CPU 线程应尽可能快、尽可能多地并发发射操作，无需等待前序任务完成。Kernel 和 CUB device 调用默认即异步，会直接挂载到传入的 stream 上执行；而 host 与 device 间的异步拷贝，则通过 CCCL 新增的 cuda::copy_bytes API 完成。 现代 CUDA 编程的最佳实践是：摒弃对默认 stream 的依赖，全程显式使用自定义 stream。 据此更新代码： 针对 pinned host buffer 的初始分配，使用独立的 init_stream；并行 for 循环的每次迭代现均绑定一条专属的 cuda::stream 用于计算 pipeline： ```cpp // 用于初始主机缓冲区分配的流 cuda::stream init_stream{cuda::device_ref{0}}; ... // 在 init_stream 上分配主机锁定缓冲区： std::vector> h_images_r(NB_IMAGES, cuda::make_pinned_buffer(init_stream, IMAGE_SIZE, ...)); ... // 在另一个流上启动操作前进行同步： init_stream.sync(); #pragma omp parallel for for (int i = 0; i < NB_IMAGES; ++i) { // 每个线程使用不同的流 cuda::stream stream{cuda::device_ref{0}}; ... // 使用每个线程拥有的流进行 GPU 缓冲区分配 cuda::device_buffer d_image_r = cuda::make_buffer(stream, device_resource, IMAGE_SIZE, cuda::no_init); ... // 使用每个线程拥有的流异步地将每个容器的内存从 CPU 复制到 GPU cuda::copy_bytes(stream, h_images_r[i], d_image_r); ... // 使用 CUB 配合每个线程的流异步地将 RGB 图像转换为灰度图像 cub::DeviceTransform::Transform(..., stream.get()); // 使用每个线程的流启动 GPU 内核以计算图像中每个图块的均值 cuda::launch(stream, ...); // 将 GPU 上的均值内存复制回 CPU cuda::copy_bytes(stream, d_median, h_medians[i]); ... // 确保在访问主机上的结果之前完成字节复制 stream.sync(); } ``` 经过这些改动后，我们最后来看一下时间线：
放大版 Nsight Systems GPU 时间线，展示多个 CUDA Stream 行。主机到设备的拷贝、内核执行和设备到主机的拷贝按行交错排列并在时间上重叠。高亮时段约为 23 毫秒，展示了三个图像处理管线的并发执行。
图 7. 为每张图片分配独立 CUDA Stream 后的最终 Nsight Systems 时间线：三条流上的内存传输与内核执行相互重叠，总处理时间从 6.8 秒降至约 23 毫秒

现在，内核执行与内存拷贝已实现完全重叠。

经过所有优化后，计算全部三张图片的最终耗时为 23 毫秒，相比最初的 6.8 秒大幅缩短。

你来试试

借助 CUDA Developer's Toolbox，我们让代码更安全、更易维护、运行更快。未使用任何底层优化手段，但代码性能已提升 300 倍。

你可以亲自尝试这份 代码，也可在 Google Colab 上运行。

我们还制作了一套完整教程课程，帮助你深入学习这些工具的使用方法。课程免费发布于 YouTube，并附有 Google Colab 练习链接。

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.
原始来源： NVIDIA 开发者博客

评论 (0)