英伟达-CUDA-入门笔记-全-

英伟达 CUDA 入门笔记（全）

001：CUDA C++基础

概述

在本节课中，我们将开始学习如何编写语法正确的CUDA程序，并确保程序能给出正确的结果。这是任何编程学习的第一步，之后我们才会关注性能优化。

课程介绍

我是来自NVIDIA的Bob Corvello，是一名解决方案架构师。我的同事Robbie Srs（来自橡树岭国家实验室）和Max Cats（来自劳伦斯伯克利国家实验室）也参与了本次课程。

感谢橡树岭和伯克利实验室主办本次活动，并提供所有必要的技术支持。本系列课程计划包含九个模块，每个模块时长约一小时，并配有课后练习来巩固所学知识。

本系列课程的目标是为您提供一个广泛而基础的CUDA编程入门。CUDA编程是我们访问GPU进行通用计算或加速工作流程的方法之一。

需要说明的是，仅凭今天这一小时的课程，您可能无法完全掌握CUDA。完整的入门至少需要前三个模块的内容。

前三个模块将涵盖必要的语法和理解，使您能够编写语法正确且结果准确的CUDA程序。第三个模块将同时向您介绍GPU架构，并从中引出两个最重要的优化概念，帮助您编写能在GPU上高效运行的代码。

其余模块将更具专题性，涵盖特定主题以扩展您作为CUDA程序员的能力。我们强烈建议您参加前三个模块，以获得最连贯的CUDA编程入门。所有课程材料将被录制并提供，方便您补课或复习。

课后练习

每个模块都配有大约三个课后练习，旨在花费约一小时完成。这些练习可以在橡树岭或伯克利实验室的GPU加速集群上完成。练习的目的首先是巩固课堂所学知识，其次也会引入和扩展新的概念。

例如，今天的练习将巩固今天学到的概念，同时也会引入一些新概念，如CUDA错误检查。错误检查是学习CUDA编程旅程中非常重要但常被忽视的一部分，CUDA运行时提供的错误线索能帮助您快速定位问题。另一个例子是，今天的课程主要讨论一维处理，而练习会将部分概念扩展到二维，从而扩展您的知识。

我们鼓励您完成课后练习，相关材料已在GitHub上提供。

提问与答疑

在课程中如有疑问，您可以使用现场麦克风提问，或通过聊天功能留言。我会在讲解过程中适时暂停，查看并回答聊天区的问题。如果未能及时回答，我的同事Robert Crs和Max Cats也会提供帮助。

进入正题：CUDA C++ 基础

明确了目标后，让我们开始今天的学习。我们的目标是启动教学过程，教会您如何编写语法正确并能给出正确答案的CUDA程序。我认为，在理解如何确保代码正确之前，过分关注性能是没有太大意义的。

因此，今天我们将本着相对简单的原则介绍CUDA编程，重点学习语法和如何编写语法正确的程序。

什么是CUDA？

“CUDA”是一个含义丰富的术语。在NVIDIA，我们在多种不同场景下使用它。

架构：CUDA代表“统一计算设备架构”。这个术语是在2006-2007年间，我们向世界推出可编程通用GPU计算能力时创造的。我们当时做出的一个重要决定是标准化此架构，要求此后我们开发和制造的所有GPU都使用完全相同的可编程架构来提供GPU加速环境下的通用计算能力。
编程环境：我们也用CUDA来指代编程环境本身。
工具包：CUDA工具包是一个软件集合，您可以下载它来设置开发环境或运行CUDA程序。

CUDA C++

今天我们将重点学习CUDA C++。虽然名义上是C++，但众所周知C++包含了C的许多特性，因此这两种语言有相当大的重叠。CUDA正式宣称符合某个有文档记录的C++标准。

在您的GPU学习之旅中，您可能还会遇到其他内容，例如CUDA Fortran、CUDA Python等，我将其统称为语言绑定。CUDA C/C++并非进入GPU编程的唯一途径，您同样可以从Python或Fortran等其他语言绑定开始。

然而，第一个推出的语言绑定就是C/C++绑定，它至今仍是我们向世界展示GPU可编程能力的核心。许多其他语言绑定本质上都是构建在CUDA C++所建立的架构和底层机制之上的。

因此，您今天学到的概念是普适且相关的。即使您决定未来的学习主要使用Python、Fortran或其他语言，这些知识依然有用。

总结

本节课我们一起了解了CUDA入门系列课程的整体安排、课后练习的重要性，并初步认识了CUDA的多重含义以及CUDA C++作为核心语言绑定的地位。下一节，我们将开始学习具体的CUDA C++编程语法。

002：CUDA共享内存

概述

在本节课中，我们将学习CUDA编程中的线程协作与通信，特别是共享内存的使用。我们将通过一个名为“一维模板”的示例问题，来探索如何让线程在同一个线程块内高效地共享数据，并实现同步操作。

上一节我们介绍了CUDA的基本概念，如主机与设备、内核启动以及内存管理。本节中，我们将看看如何让线程之间进行协作。

课程内容

大家好，我是来自ULCF的Tom Papicuttor。今天，我们邀请到了英伟达的Bob Corbella来进行讲解。这是九部分CUDA培训系列的第二部分。

现场有来自橡树岭国家实验室的参与者，也有来自英伟达加州办公室的参与者，还有通过Webex远程加入的朋友。请大家在提问时说明自己所在的平台。

Bob在讲解过程中会定期暂停以回答问题。报告结束后，我们将进入实践环节。现场参与者可以获得帮助，远程参与者也可以通过Webex获得支持。实践练习和Bob的幻灯片都可以在课程注册页面找到。

由于第一节课的内容已在线发布，我将跳过大部分开场白。但我要感谢我的同事，以及橡树岭国家实验室和NurRS提供的所有支持和后勤保障。

如果你听了上个月的课程，就会知道前几节课的内容是紧密相连的。我建议尽可能跟上前三四节课，因为它们是CUDA编程的基础。随后的课程会更专题化。如果错过了前四节中的任何一节，可能会遗漏重要的基础知识。

希望你已经有机会学习了上个月的第一节课。在那节课中，我们通过一个“向量加法”的示例问题，介绍了多种概念。

我们学习了“主机”和“设备”的术语。我们学习了 __global__ 关键字。主机和设备的术语实际上存在于CUDA语法中，用于泛指在CPU（主机）或GPU（设备）上进行的活动。__global__ 关键字用于标识一个内核函数，这是我们作为CUDA程序员编写并启动在GPU上运行的代码的方式。

我们还学习了如何向运行的内核传递参数，特别是如何在GPU上进行内存管理。我们学习了如何在GPU内存中分配指针、释放这些指针，以及如何使用 cudaMemcpy API在主机内存和这些指针之间复制数据。我们还探索了启动并行内核，并学习了如何单独使用线程块和线程，以及如何将两者结合起来。

在上节课中，我提到还有更多内容需要学习。我们接下来要探讨的主题之一就是线程的协作与通信。我当时提到，我们需要一个新的示例来展开讲解。这就是我们现在要引入的内容。

我们将介绍今天课程要重点关注的下一问题。这个问题将帮助我们解析和探索一些新的CUDA语法和功能。

在深入探讨之前，我想说明，本月的课程可能会短很多。因此，我们预计会有更多时间用于提问。如果你在复习上个月课程或完成作业时产生了问题，今天无疑是提问的好时机。正如Tom提到的，你可以随时在聊天区提问。我会在讲解过程中暂停一两次查看聊天内容。我预计本小时的后半部分会有充足的时间回答听众的问题。所以，在我们学习本月材料时，请随时思考问题。

我们将引入的示例问题，我上个月提到过，叫做“一维模板”。模板操作的核心思想是，一个模板反映了一个数据窗口或范围，这个窗口被连续地应用于一个数据集以产生结果。我们称之为模板操作。我们将在讲解中看到一个一维模板的例子。

为了将一维模板应用于一维数组，我们首先需要理解模板本身的一些特性。

模板有一个宽度。宽度是窗口的整体宽度，或者说用于进行逐点计算的基础数据的宽度。在这个例子中，我们假设一个宽度为7的模板。我们还可以定义一个称为“半径”的术语。半径不过是模板中心点左侧或右侧的数据量。在本例中，半径为3，这给出了7的整体宽度。通常，在进行模板操作时，我们有模板中心的概念。因此，当我们讨论模板维度时，一维模板的宽度通常是奇数。如果我们有一个二维模板（也可以称为窗口操作），我们可能有两个维度都是奇数，这样会有一个明确定义的中心点，并且在中心点上下或左右的计算上具有对称性。这里我们有一个一维模板，所以我们只关注一维操作、一维模板、一维窗口上的计算。这些计算在本例中将聚焦于基础数据集的7个元素。

我们想象基础数据集比7个元素长得多。这里有一个可以被视为更长数据集的例子。我们可以把这些绿色块想象成像素（如果我这么说请见谅，但说“像素”比说“绿色块”更容易）。让我们想象它们是像素，或者是与一维数组相关联的数据。这些绿色数据被设想为同时存在于输入和输出数组中。因此，我们有一个一维操作，我们有一组数据将从输入转换为输出，输入和输出的大小相同或大致相同。我们将对此数据集应用模板操作以产生输出数据集，我们稍后会看到这具体意味着什么。

我们还可以想象，这张幻灯片上描绘的这组绿色像素或绿色块，实际上只是我整个数据集的一个子集。请记住，在CUDA中，我们有工作单元的分层分解概念，第一层层次是线程块。因此，我们可以想象一个内核启动包含大量线程，这些线程被分组到块中。正如我们稍后将看到的，在本节课中，我们将特别关注线程块级别的行为。因为我们将要引入的概念，即通信和同步，是在线程块级别工作的。这将迫使我们作为CUDA程序员，开始思考将问题分解为包含一组线程的块。

总结

本节课中，我们一起学习了CUDA线程协作与通信的基础，重点介绍了共享内存的概念。我们通过一维模板的示例，理解了如何在线程块内组织线程、共享数据并进行同步，这是构建更复杂并行算法的重要一步。

003：CUDA基础优化第一部分

在本节课中，我们将从关注编写语法正确的CUDA代码，转向探讨如何编写高性能的CUDA代码。我们将首先了解GPU架构的基本概念，并学习第一个核心优化目标：暴露足够的并行性。理解这一点将帮助你在编写代码时，就为GPU的高效运行打下基础。

GPU架构概述

上一节我们介绍了编写CUDA程序的基本语法。本节中，我们来看看为什么性能优化需要了解GPU架构。在计算机科学中，代码的性能往往与运行它的机器架构密切相关。为了理解如何编写高性能CUDA代码，我们需要对GPU架构有一个基本的认识。

我们将简要回顾Kepler、Maxwell、Pascal和Volta这几代架构。对于本次课程的参与者而言，Summit超算上使用的Volta架构最为重要。不过，我们今天要讨论的“暴露足够并行性”这一核心概念，在所有GPU架构中都是普遍适用的。

GPU如何执行：线程与线程块

为了理解“暴露足够并行性”的含义，我们需要了解GPU是如何组织和执行计算任务的。GPU的计算核心被组织成多个流式多处理器。

以下是一个简化的执行模型：

线程：是CUDA中最基本的执行单元。
线程块：一组线程被组织成一个线程块，在一个流式多处理器上执行。
网格：所有线程块组成一个网格，在GPU上启动。

GPU的高性能来自于其能够同时管理并执行成千上万个线程。如果线程数量不足，GPU的众多计算核心就会处于闲置状态，无法充分发挥其计算能力。

核心优化目标：使用大量线程

因此，我们从架构分析中得出的第一个、也是最重要的优化优先级，可以转化为一个直接的编程建议：使用大量线程。

对于刚学完前两节语法课程的开发者来说，这是一个简单易懂的目标。你的算法应该能够启动并利用海量线程来执行。这是实现高性能CUDA程序的基础，其重要性超过我们今天要讨论内容的50%。

为了实现这个目标，我们需要关注内核的启动配置。启动配置决定了你创建了多少个线程块，以及每个线程块包含多少个线程。

以下是设计启动配置时需要考虑的几个关键点：

总线程数：应远大于你的数据量或任务量，以确保GPU被充分利用。
线程块大小：通常是32的倍数（即一个Warp的大小），例如128、256、512。
网格大小：通过总线程数除以线程块大小来计算，确保能覆盖所有需要处理的数据。

通过合理配置，你可以确保GPU获得足够多可以并行执行的工作，从而隐藏内存访问延迟，并让所有计算单元保持忙碌。

总结

本节课中，我们一起学习了CUDA性能优化的第一个核心理念。我们了解到，GPU的架构设计使其擅长海量并行计算，因此编写高性能CUDA代码的首要任务是暴露足够的并行性，即在程序中启动并使用大量线程。

我们探讨了GPU通过流式多处理器、线程块和线程来组织执行的基本模型，并明白了如果线程数量不足，GPU的强大算力将无法被有效利用。最后，我们将这一架构知识转化为具体的编程实践——精心设计内核的启动配置，以确保网格和线程块能提供远超GPU物理核心数量的线程。

在下一节课中，我们将深入GPU的另一个关键子系统：内存层次结构。我们将学习如何高效地利用全局内存、共享内存等，这是CUDA编程中第二个至关重要的优化方向。无论你使用CUDA C++、CUDA Python还是CUDA Fortran，本节课所学的“使用大量线程”这一原则都同样适用。

004：CUDA基础优化第二部分

概述

在本节课中，我们将学习CUDA编程中第二个至关重要的优化概念：如何高效利用内存子系统。我们将重点关注全局内存的吞吐量以及共享内存的有效使用。理解这些概念将帮助您编写出性能更佳的CUDA代码，而无需过度依赖性能分析工具。

上一节我们介绍了通过暴露大量并行性（即运行大量线程）来优化执行行为。本节中，我们来看看如何通过优化内存访问模式来进一步提升性能。

内存层次结构回顾

为了理解如何优化内存访问，我们首先需要回顾GPU的内存层次结构。GPU的内存系统从最靠近处理单元的高速存储到容量较大但速度较慢的存储，形成了一个层次结构。

以下是GPU内存层次结构的主要组成部分：

寄存器：这是每个线程私有的、速度最快的内存资源。几乎所有底层机器指令都从寄存器读取数据并将结果写回寄存器。寄存器使用主要由编译器管理，程序员通常无需直接干预。
共享内存与L1缓存：这两者都是位于GPU芯片上的资源，具有高带宽和低延迟的特点。
- 共享内存：是一个可由程序员显式分配和使用的本地内存数组，用于数据的临时存储。每个线程块通常至少有48KB的共享内存可用。
- L1缓存：是一种硬件管理的缓存，旨在通过保留最近使用的数据来提升访问速度，对程序员基本透明。
L2缓存：这是一个设备范围的资源，所有进出全局内存的数据都需要经过L2缓存。与L1类似，它通过缓存频繁访问的数据来提升性能。
全局内存：这是GPU中容量最大，但访问延迟最高、带宽相对较低的内存。优化对全局内存的访问是本课的重点。

全局内存吞吐量优化

现在，让我们深入探讨如何优化全局内存的访问。全局内存的访问模式对性能有决定性影响。为了获得高带宽，关键是要实现合并访问。

合并访问原则

当GPU中的一个线程束（Warp，通常是32个线程）访问全局内存时，如果所有线程访问的内存地址是连续的，并且对齐到特定的边界（例如128字节），那么这些访问可以被合并成一个或少数几个内存事务。这种高效的访问模式称为“合并访问”。

反之，如果线程束中的线程访问分散的、不连续的内存地址，则会导致多个内存事务，严重降低有效带宽。这种低效的访问模式称为“非合并访问”或“分散访问”。

核心优化原则可以总结为以下公式：
高效全局内存访问 ≈ 实现合并访问

为了实现合并访问，在编写内核时应注意：

确保同一线程束内的线程访问连续的全局内存地址。
尽量使访问的起始地址对齐到缓存行边界。

共享内存的高效使用

在理解了全局内存优化后，我们接下来看看如何利用共享内存来进一步提升性能。共享内存的访问速度远高于全局内存，因此可以作为一种软件管理的缓存。

使用共享内存的模式

一个常见的使用模式是“平铺”（Tiling）算法。其基本思想是：

将数据从全局内存分块（Tile）加载到共享内存中。
让线程块内的线程协作处理共享内存中的这块数据。
将结果写回全局内存。

这种模式能显著减少对全局内存的重复访问，尤其适用于存在数据重用的算法（如矩阵乘法、卷积等）。

以下是使用共享内存的一个简化代码框架示意：

__global__ void kernel(float* input, float* output) {
    // 1. 声明共享内存数组
    __shared__ float tile[TILE_SIZE];

    // 2. 协作将全局内存数据加载到共享内存
    int idx = ...; // 计算全局索引
    tile[threadIdx.x] = input[idx];
    __syncthreads(); // 确保所有数据加载完成

    // 3. 从共享内存读取数据进行计算
    float data = tile[some_index];
    // ... 执行计算 ...

    // 4. 将结果写回全局内存
    output[idx] = result;
}

使用共享内存时，必须注意使用 __syncthreads() 来同步线程块内的线程，确保数据在共享内存中准备就绪后再被读取。

总结

本节课中我们一起学习了CUDA基础优化的第二个核心支柱：高效利用内存子系统。

我们首先回顾了GPU的内存层次结构，理解了从高速的寄存器、共享内存/L1缓存，到L2缓存，最后到全局内存的访问速度差异。接着，我们重点探讨了优化全局内存访问的关键——实现合并访问，以确保获得高内存带宽。最后，我们介绍了如何利用更快的共享内存作为软件管理的缓存，通过“平铺”等模式来减少对全局内存的访问，从而提升整体性能。

掌握这两个概念——优化全局内存访问模式与有效利用共享内存——是编写高性能CUDA代码的基石。在后续课程中，我们将学习更多基于性能分析驱动的优化技巧。

005：Atomics, Reductions, Warp Shuffle 🧮

在本节课中，我们将学习CUDA编程中的三个核心概念：原子操作、归约操作和Warp洗牌操作。这些技术对于解决需要线程间协作和数据共享的并行计算问题至关重要。

概述

大家好，我是来自NVIDIA的Bob Carbella。在今天的第五次CUDA培训课程中，我们将探讨原子操作、归约操作和Warp洗牌操作。虽然听起来像是三个独立的主题，但我们会发现它们之间存在紧密的联系，并且常常被用来解决同一类编程问题。

与之前的基础性课程不同，从本节课开始，我们将进入更具专题性的学习阶段。课程结构将分为两部分：第一部分是约一小时的讲解，第二部分是问答环节。此外，由于本次作业内容更为复杂，我们将在后续安排一次专门的作业讲解会。

从“变换”到“归约”

在之前的课程和作业中，我们处理的大多数问题都属于“变换”类问题。这类问题的特点是：输入数据集的大小与输出数据集的大小大致相同。例如，对数组中的每个元素进行某种运算，并生成一个大小相近的新数组。

然而，我们接下来要处理的问题有所不同。考虑一个简单的C语言示例：计算一个包含100，000个元素的数组的总和。

int a[100000];
int sum = 0;
for (int i = 0; i < 100000; i++) {
    sum += a[i];
}

在这个例子中，输入是一个庞大的数组，但输出只是一个单一的值（sum）。这类问题被称为“归约”问题。将这类问题并行化，需要我们思考新的“线程策略”。

什么是线程策略？

线程策略是CUDA编程中的一个核心设计问题，它回答的是：“我应该让每个线程做什么？” 当我们编写一个CUDA核函数时，本质上是在编写单个线程要执行的代码。CUDA的并行化机制（如网格和线程块维度）会将这段代码复制到成千上万个线程上并行执行。

因此，在开始编写任何CUDA内核之前，我们必须首先确定线程策略。对于“变换”类问题，策略通常很直观：让每个线程处理一个（或几个）输入元素，并产生一个（或几个）输出元素。

但对于“归约”问题，情况就变得复杂了。所有线程都需要协作，将大量数据“归约”成一个或几个值。这引出了我们本节课要学习的三个关键技术。

原子操作

原子操作是解决多线程同时读写共享内存时数据竞争问题的关键工具。它确保了对某个内存位置的操作是“不可分割”的，即在该操作完成前，其他线程无法访问该位置。

以下是原子加操作的示例：

// 所有线程都向全局内存中的同一个地址累加自己的值
atomicAdd(&global_sum, thread_value);

核心概念：atomicAdd 等原子函数保证了即使在成千上万个线程同时执行的情况下，对 global_sum 的累加操作也能正确、顺序地完成，而不会发生数据覆盖或丢失。

归约操作

归约操作是指将一组数据通过某种二元运算（如加法、求最大值）合并成单个值的过程。在CUDA中实现高效的归约需要巧妙的线程协作模式。

一个经典的并行归约策略是树形归约：

每个线程块将其内部的数据归约成一个部分和。
然后，这些部分和被进一步归约（可能通过另一个内核或原子操作）以得到最终结果。

这种策略极大地减少了直接使用原子操作带来的全局内存访问冲突和性能开销。

Warp洗牌操作

Warp洗牌是CUDA提供的一种在同一个Warp（通常是32个线程）内进行线程间数据交换的机制。它比通过共享内存进行数据交换速度更快，延迟更低。

以下是使用Warp洗牌进行归约的示例：

// 假设每个线程都有一个值 `val`
for (int offset = 16; offset > 0; offset /= 2) {
    val += __shfl_down_sync(0xffffffff, val, offset);
}
// 现在，Warp内第一个线程的 `val` 包含了该Warp所有线程值的和

核心概念：__shfl_down_sync 内在函数允许Warp内的线程直接从其他线程的寄存器中获取数据，无需经过共享内存，从而实现了极低开销的数据共享和归约。

总结

在本节课中，我们一起学习了CUDA中用于解决归约和线程协作问题的三个关键技术：

原子操作：用于安全地在多线程环境下更新共享变量，是解决数据竞争的基础，但可能成为性能瓶颈。
归约操作：通过树形结构等并行算法，高效地将大量数据合并为少数值，是许多科学计算的核心。
Warp洗牌操作：在Warp级别实现高速、低延迟的线程间数据交换，是优化归约和特定计算模式的强大工具。

理解这些概念及其相互关系，是编写高效、正确CUDA并行程序，特别是处理非变换类问题的重要一步。在接下来的作业中，你将有机会亲自实践这些技术。

006：托管内存

在本节课中，我们将学习CUDA托管内存（也称为统一内存）的概念。这是一种旨在简化GPU编程中内存管理的技术。我们将探讨其基本原理、特性、典型用例，并重点讨论其性能表现。请注意，托管内存的主要目标是简化编程，而非直接提升性能。

什么是托管内存？

上一节我们介绍了GPU计算通常包含的三个步骤：数据复制到GPU、执行内核、结果复制回CPU。托管内存旨在简化这一流程。

托管内存的核心思想是消除显式内存复制的“样板代码”。它通过创建一个单一指针来实现，该指针在主机（CPU）代码和设备（GPU）代码中均可使用。从程序员视角看，数据似乎只有一份副本，无需手动管理主机和设备上的两个独立副本。

这项功能大约在CUDA 6时代（2012-2013年，与Kepler GPU同期）被引入CUDA编程模型。

托管内存的基本用法

要使用托管内存，您只需使用 cudaMallocManaged 函数来分配内存，而不是分别使用 malloc（主机）和 cudaMalloc（设备）。

以下是分配托管内存的基本代码示例：

// 分配托管内存
int *data;
cudaMallocManaged(&data, N * sizeof(int));

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-intro/img/163d4829ac6bf48680a5870cc82ae046_2.png)

// 现在，`data` 指针既可以在主机代码中使用，也可以在设备内核中使用
// ... 在主机上初始化数据 ...
// ... 启动内核处理 `data` ...
// ... 在主机上访问结果 ...

// 释放内存
cudaFree(data);

使用托管内存后，数据迁移（在主机和设备之间移动）由CUDA运行时系统自动管理，程序员无需编写显式的 cudaMemcpy 调用。

托管内存的关键特性

以下是托管内存的一些重要特性：

按需分页：数据并非在分配时立即全部传输。只有当CPU或GPU实际访问（读取或写入）内存的某个页面时，该页面才会被迁移到访问它的处理器上。
内存超额订阅：在支持它的系统上（如带有GPU的服务器），应用程序可以分配比GPU物理显存更大的托管内存。CUDA运行时会在需要时在主机内存和显存之间交换数据页面。
一致性模型：在计算能力6.0及以上的GPU上，托管内存提供系统范围的原子内存操作一致性。这意味着主机和设备可以同时对同一内存位置进行原子操作。
简化代码移植：对于原本为CPU编写的代码，将其指针改为指向托管内存，并添加内核启动，可能就足以让部分计算在GPU上运行，从而简化移植过程。

性能考量

现在，我们来看看托管内存的性能。重申一遍，托管内存的主要目标是编程简化，而非性能优化。

基础性能：由于存在按需分页和数据迁移的开销，使用托管内存的代码性能可能低于精心手动管理内存复制的代码。首次访问数据时可能会产生页面错误和迁移延迟。
性能优化API：CUDA提供了一些API来帮助管理托管内存的性能：
- cudaMemPrefetchAsync：此函数允许您将数据预取到特定设备（例如GPU）的内存中，以减少内核执行时的页面错误延迟。您可以在启动内核前，将所需数据预取到GPU。
- cudaMemAdvise：此函数允许您向运行时提供关于数据使用模式的建议（例如，cudaMemAdviseSetPreferredLocation 可以设置数据的首选存放位置），以帮助运行时做出更优的数据迁移决策。

总结

本节课我们一起学习了CUDA托管内存（统一内存）。我们了解到，它通过提供单一指针和自动数据迁移，显著简化了GPU编程中的内存管理。我们讨论了其按需分页、超额订阅等关键特性，并审视了其性能特点。重要的是要记住，托管内存是一种生产力工具，用于简化开发，通常需要配合 cudaMemPrefetchAsync 等API来优化性能，而非直接提供最高性能的解决方案。在决定是否使用托管内存时，应在编程便利性和性能要求之间进行权衡。

007：并发性

在本节课中，我们将要学习CUDA编程中的并发性概念。并发性是优化CUDA程序性能的关键技术之一，它允许我们同时执行多个操作，从而缩短程序的总体运行时间。我们将从动机开始，探讨为何需要并发性，然后介绍实现并发所需的基础知识，特别是“固定内存”的概念。最后，我们将学习实现并发的具体机制和不同使用场景。

动机：为何需要并发？

上一节我们介绍了CUDA的基本编程模型。现在，我们来看看如何通过并发性来优化这个模型。

在CUDA编程中，一个典型的三步序列是：将数据从主机复制到设备，执行内核，然后将结果从设备复制回主机。如果我们按照目前所学的方式编写代码，这些操作通常是串行执行的。这意味着内核执行必须等待数据复制完成，而结果复制也必须等待内核执行完成。

这种串行执行方式导致程序的总执行时间是各个操作时长的简单相加。然而，如果我们能够重叠这些操作，让数据复制和内核执行同时进行，就有可能显著缩短程序的总体运行时间。这正是并发性优化的核心目标——通过同时利用多个处理器来最大化系统性能。

固定内存：实现并发的基础

为了实现上述的并发操作，我们首先需要理解一个关键概念：固定内存。

现代操作系统（如Linux或Windows）的内存子系统具有一定的复杂性，可能对程序员并不完全透明。其核心思想是，操作系统将虚拟内存空间与物理内存（即系统RAM）分离开来。

固定内存是指被“锁定”在物理RAM中的主机内存页，操作系统不会将其交换到磁盘，并且其物理地址是固定的。这对于CUDA的异步内存传输至关重要，因为GPU的DMA引擎需要知道稳定的物理地址才能高效地直接访问主机内存。

以下是使用固定内存的代码示例：

cudaMallocHost(&pinned_host_ptr, size); // 分配固定内存
// ... 使用 pinned_host_ptr 进行异步操作
cudaFreeHost(pinned_host_ptr); // 释放固定内存

并发机制与使用场景

理解了固定内存后，我们现在可以探讨CUDA中实现并发的具体机制。

CUDA提供了多种机制来实现不同层面的并发。以下是几种主要的并发使用场景：

内核执行与内存拷贝的重叠：这是最常见的并发优化。通过使用流和异步内存拷贝，可以让一个流中的内核执行与另一个流中的数据拷贝同时进行。
多个内核的并发执行：如果设备有足够的资源，可以同时启动多个内核，让它们并行执行。
主机与设备的并发执行：主机CPU可以在GPU执行内核的同时，进行自己的计算任务，实现CPU-GPU的协同工作。

为了实现这些并发场景，我们需要使用CUDA流。流是一系列按顺序执行的命令序列。不同流中的命令可以并发执行。以下是创建和使用流的示例：

cudaStream_t stream;
cudaStreamCreate(&stream); // 创建流
cudaMemcpyAsync(dev_ptr, host_ptr, size, cudaMemcpyHostToDevice, stream); // 在指定流中进行异步拷贝
myKernel<<<grid, block, 0, stream>>>(...); // 在指定流中启动内核
cudaStreamSynchronize(stream); // 等待流中所有操作完成
cudaStreamDestroy(stream); // 销毁流

总结

本节课中我们一起学习了CUDA中的并发性。我们从优化程序执行时间的动机出发，认识到重叠数据拷贝和内核执行的重要性。为了实现这种并发，我们首先学习了固定内存的概念及其作用。最后，我们探讨了使用CUDA流来实现内核执行与内存拷贝重叠等并发场景的具体方法。掌握这些并发技术是进行CUDA程序高级性能优化的关键步骤。

008：GPU性能分析 🚀

在本节课中，我们将学习GPU性能分析的核心方法论——分析驱动优化。我们将了解如何系统性地使用性能分析工具来识别代码瓶颈，并基于数据而非猜测来指导优化工作。课程将涵盖关键的性能限制因素，如内存带宽、计算吞吐量和延迟，并介绍如何使用NVIDIA Nsight工具进行实践分析。

分析驱动优化方法论 💡

上一节我们概述了课程目标，本节中我们来深入探讨分析驱动优化的核心思想。

分析驱动优化是一种系统性的性能提升方法。其核心在于：让性能分析工具的数据指导优化方向，而非依赖直觉或未经证实的假设。

许多开发者常会陷入针对特定概念（如“消除线程束分化”）进行盲目优化的误区。然而，在没有数据支持的情况下，这类优化往往收效甚微，甚至可能完全无关紧要。分析驱动优化要求我们遵循一个可重复、有逻辑的流程：

使用性能分析工具收集数据。
识别最主要的性能瓶颈。
针对该瓶颈进行代码重构或优化。
验证优化效果，并重复此过程。

这种方法论是长期实践积累的成果，具有普适性和可靠性。无论工具或代码如何变化，这套方法论都能持续生效。

关键性能分析路径 🛣️

理解了方法论后，我们来看看分析过程中可能遇到的几种主要性能限制类型。根据初始分析结果，我们通常会沿着以下几条关键路径进行深入探索：

内存带宽限制：当数据在GPU显存与计算单元之间的传输速度成为瓶颈时。
计算吞吐量限制：当GPU的计算单元（如CUDA核心）利用率不足，无法跟上指令发射速度时。
延迟限制：当线程因等待数据（如全局内存访问）或同步而长时间空闲时。

沿着不同路径前进，我们需要关注不同的性能指标和分析重点。接下来，我们将引入性能指标的概念，用于量化和确定具体是哪种因素在限制性能。

现代分析工具：Nsight Systems & Nsight Compute 🛠️

上一节我们介绍了性能限制的类型，本节中我们来看看帮助我们识别这些限制的现代工具。

一个理想的智能工具能够封装专家经验，为用户提供最直接、最有用的优化指南。NVIDIA的Nsight Systems和Nsight Compute正是这样的工具。它们在近几年推出，并集成了开发者社区多年积累的智能分析能力。

Nsight Systems提供系统级的性能分析，适用于观察应用程序的整体行为、CPU-GPU交互以及多GPU/多流活动。Nsight Compute则专注于GPU内核的微观性能分析，提供详细的硬件计数器数据和优化建议。

在本次演示的后半部分，我们将主要跟随Nsight Compute工具提供的“线索”进行分析。该工具内置的专家系统，其底层逻辑正是基于我们前面讨论的分析驱动优化方法论。

本节课中，我们一起学习了GPU性能分析的核心——分析驱动优化方法论，了解了内存带宽、计算吞吐量和延迟等关键性能限制路径，并认识了Nsight Systems和Nsight Compute这两款强大的现代性能分析工具。记住，让数据而非猜测来指导你的优化工作，是持续提升代码性能的关键。

009：Cooperative Groups 🚀

在本节课中，我们将学习CUDA编程模型中的一个重要扩展——Cooperative Groups（协作组）。我们将探讨如何创建和管理不同规模的线程组，以实现更灵活的线程间协作与同步。课程内容涵盖线程块级、网格级、多设备级协作组以及集合操作。通过本课的学习，您将掌握使用协作组编写高效、可组合CUDA代码的方法。

课程概述 📋

协作组功能旨在为线程组之间的协作提供支持。在CUDA编程中，线程协作主要涉及两个基本概念：同步（如执行屏障）和线程间通信（数据共享）。协作组扩展了这些概念，允许我们创建灵活大小的线程组，并设计能够处理这些可变大小组的算法。

上一节我们介绍了CUDA编程的基础概念，本节中我们来看看如何利用协作组实现更高级的线程协作。

协作组的动机 💡

协作组子系统主要促进线程组之间的合作。在之前的课程中，我们学习了线程协作的含义，包括 __syncthreads() 和共享内存的使用。协作组在此基础上，允许我们：

创建灵活的并行分解，将大组细分为小组。
设计能够处理可变大小线程组的算法。
编写可跨软件边界组合的代码。

其中一个最引人注目的概念是网格级同步（grid-wide sync）。在传统的CUDA编程模型中，实现整个网格的同步是困难且受限的。协作组引入的网格级同步功能，为许多算法问题提供了强大的解决方案。

传统CUDA的协作机制 ⚙️

在CUDA 9引入协作组之前，程序员可用的线程协作机制相对有限。主要依赖于线程块内的固有协作原语。

以下是传统CUDA中线程协作的核心机制：

线程块内同步：使用 __syncthreads() 函数。
线程间通信：通过共享内存（__shared__ 变量）实现。

这些机制虽然强大，但缺乏在不同规模线程组（尤其是跨线程块或整个网格）上进行灵活协作和同步的直接支持。

协作组的核心概念 🧠

协作组扩展了CUDA模型，引入了“组”作为一等公民。一个组可以包含任意数量的线程，从几个线程到整个网格的所有线程。

核心操作包括：

创建组：从现有线程集合中定义一个新的协作组。
组内同步：确保组内所有线程到达同步点后再继续执行。
组间通信：在组内或特定模式的线程间交换数据。

不同级别的协作组 📊

协作组功能支持多个级别，每种服务于不同的目的。

以下是主要的协作组类型：

线程块级协作组 (Thread Block Level)
- 用于线程块内部的协作。
- 可以细分为更小的小组（如经线组Warp、线程块瓦片Tile）。
网格级协作组 (Grid Level)
- 允许跨线程块的线程进行协作。
- 实现了之前难以达成的网格级同步。
多设备协作组 (Multi-Device Groups)
- 扩展协作范围到多个GPU设备。
- 用于复杂的多GPU算法。
集合操作组 (Collectives Groups)
- 专为集合通信模式设计（如规约、扫描）。
- 提供优化的原语。

网格级同步的重要性 ⚡

网格级同步是协作组带来的一个关键能力。在早期CUDA模型中，除了在内核启动之间隐式同步外，没有安全的方法让网格中所有线程在同一个执行点汇合。

网格级同步使得以下操作成为可能：

实现需要全局屏障的算法。
在网格范围内协调对全局内存的访问。
编写更清晰、更易于推理的多阶段内核。

其基本用法如下：

// 创建包含网格中所有线程的组
auto grid_group = cooperative_groups::this_grid();
// 执行网格级同步
grid_group.sync();

示例与应用场景 🛠️

协作组适用于多种并行模式。

以下是一些典型的应用场景：

并行规约 (Parallel Reduction)：可以动态地将参与规约的线程组织成不同大小的组，实现更灵活的规约树。
扫描/前缀和 (Scan/Prefix Sum)：协作组有助于管理扫描操作中的依赖关系。
生产者-消费者模式 (Producer-Consumer Patterns)：网格中的线程组可以协作处理任务队列。
迭代求解器 (Iterative Solvers)：需要在每次迭代后进行全局同步的算法。

总结与回顾 🎯

本节课中我们一起学习了CUDA Cooperative Groups的核心概念与应用。我们了解到，协作组功能填补了传统CUDA模型在灵活线程协作方面的空白，特别是引入了强大的网格级同步能力。通过创建和管理不同规模的线程组，我们可以设计出更高效、更模块化且可组合的CUDA内核。

关键要点包括：

协作组将“线程组”作为编程模型中的显式实体。
它支持从经线到整个网格乃至多设备级别的协作。
网格级同步解锁了新的算法设计可能性。
使用协作组可以编写出更清晰、更易于维护的并行代码。

建议您通过官方文档和练习进一步巩固对这些概念的理解，并将其应用到实际的CUDA项目中去。

010：多线程与CUDA并发 🚀

在本节课中，我们将要学习如何在多线程环境中使用CUDA流来实现并发执行。我们将探讨CUDA流的基本概念、默认流的行为，以及如何通过流来重叠计算和数据传输，从而提升GPU的利用率和应用程序的性能。

上一节我们介绍了CUDA并发的基本概念，本节中我们来看看CUDA流的具体语义和行为。

CUDA流是CUDA中用于管理操作执行顺序的抽象。以下是关于CUDA流的一些基本语义：

在同一个流中发出的两个操作，将按照发出顺序执行。这意味着，如果操作A和操作B被依次提交到同一个流中，操作B必须等待操作A完成后才能开始执行。
在不同流中发出的操作，CUDA不保证它们之间的执行顺序。这意味着，操作A可能在操作B之前、期间或之后执行。

通常，一个“操作”指的是内存拷贝（如cudaMemcpyAsync）或内核函数调用。更广泛地说，任何接受流参数的CUDA API调用（如流回调）都属于此类操作。对于您自己编写的内核，调用时也可以指定一个可选的流参数。

理解了流的基本语义后，我们来看一个具体的代码示例，它展示了如何创建流并实现计算与拷贝的重叠。

// 创建两个CUDA流
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 在流1上异步拷贝数据到设备
cudaMemcpyAsync(dst, src, size, cudaMemcpyHostToDevice, stream1);

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-intro/img/4bfc3307933eb52ba61dd372c6495da3_5.png)

// 在流2上启动一个内核
myKernel<<<grid, block, 0, stream2>>>(...);

// 让CPU线程等待流2上的所有操作完成
cudaStreamSynchronize(stream2);

在这个例子中，内存拷贝操作在stream1上执行，而内核启动在stream2上。由于这两个操作在不同的流中，它们可能会并发执行。cudaStreamSynchronize(stream2)调用会阻塞CPU线程，直到stream2上的内核执行完毕。需要注意的是，如果内存拷贝的数据是内核计算所必需的，那么这种安排就是不安全的，因为CUDA不保证跨流操作的顺序。

为了更清晰地说明如何利用流进行任务分块和重叠，我们来看一个向量运算的例子。

这个例子类似于我们作业中的代码，它将一个大型数组分成多个块，并使用多个流来重叠主机到设备的数据拷贝、内核计算以及设备到主机的回拷。

以下是其执行逻辑的简化描述：

在stream0上：执行第一块数据的 H2D拷贝 -> 内核计算 -> D2H回拷。
在stream1上：执行第二块数据的 H2D拷贝 -> 内核计算 -> D2H回拷。

在每个流内部，操作严格按照提交顺序执行。但是，stream0和stream1之间的操作没有顺序约束，因此它们可以并发执行。关键在于，必须确保在每个流内部，数据拷贝在内核启动之前完成。这种模式不仅适用于分块处理单一数组，也适用于准备和处理多个独立的数据集。

在深入多线程之前，我们必须理解默认流的行为，因为它常常是并发编程中的陷阱。

如果您从未指定过流，那么您的所有操作都使用了默认流。传统默认流的行为是：它会阻塞整个GPU设备。这意味着，即使您在其他流中提交了任务，一旦向默认流提交工作，它会等待所有其他流上的任务完成后才开始执行。在此期间及默认流执行过程中，其他流上的任务也无法开始。

所有主机线程共享同一个默认流。在复杂的并发场景中，应避免使用传统默认流。一个常见的误解是，先启动一个任务（使用默认流），再为第二个任务创建一个新流，期望它们并发。实际上，第一个任务仍在默认流中，会阻塞第二个任务，无法实现并发。

您可以通过在编译时使用--default-stream per-thread标志来修改这一行为。启用后，每个主机线程将拥有自己独立的“普通”默认流，它不再阻塞整个设备。

掌握了流的基础知识后，我们现在可以探讨CUDA并发的几种典型用例。

CUDA并发主要有以下两种应用场景：

主机-设备执行并发：启动一个内核后，如果CPU端暂时不需要其结果，我们可以不等待（不调用cudaStreamSynchronize），而是让CPU代码继续执行其他任务。
并发内核执行：在不同的流中启动多个内核，它们有可能并发执行。这里需要强调“有可能”，因为实际并发取决于GPU硬件资源的利用率。如果两个内核各自都能高效利用大部分GPU资源（如SM、寄存器等），那么当它们被同时提交时，很可能一个会先占用所有资源执行，另一个则需等待资源释放，导致只有很少的重叠执行时间。要观察到真正的内核并发，通常需要那些资源占用率低但执行时间长的内核。

本节课中我们一起学习了CUDA流在多线程并发编程中的核心作用。我们明确了流的语义，即同流顺序、异流无序。通过代码示例，我们学习了如何创建和使用流来重叠计算与数据传输。我们特别指出了传统默认流对设备全局的阻塞特性，并给出了解决方案。最后，我们探讨了主机-设备并发及内核并发两种模式，并理解了实现真正内核并发所需的资源条件。合理利用CUDA流是优化GPU应用程序性能的关键技术之一。

011：CUDA多进程服务 (MPS) 🚀

在本节课中，我们将学习CUDA多进程服务。我们将探讨当多个进程共享同一GPU时会发生什么，理解默认行为，并介绍MPS如何帮助提升此类场景的性能。

概述

在开始之前，我们先思考一个高性能计算中的常见问题：假设我有一个固定量的工作需要完成，我可以选择将其分配到N个MPI进程中。那么，如何确定最优的进程数N？以及，应该将这些进程分配到多少个GPU上？通常，我们会为每个进程分配一个GPU，并尽可能多地增加GPU数量以缩短求解时间，或者尽可能大地利用每个GPU的内存来处理问题。

一个简单的用例

为了具体分析，我们来看一个简单的CUDA C内核。这个内核的功能非常简单：它接收一个数组并将其中的每个元素加倍。这是一个典型的内存密集型工作负载。

__global__ void double_array(double *array, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        array[idx] = 2.0 * array[idx];
    }
}

这个工作负载的大小与数组长度n成正比。我们假设n为1024的三次方，即大约10亿个元素。对于一个双精度数组（每个元素8字节），这大约占用8GB内存，这已经接近现代GPU的显存容量。

在单个GPU上运行此内核（连续启动1000次）时，每次内核执行大约需要19.7毫秒。那么，一个自然的问题是：我能否通过使用更多的MPI进程来加快速度？

多进程与单GPU

上一节我们介绍了在单个GPU上运行单个进程的情况。本节中我们来看看，如果我们在同一个GPU上运行多个MPI进程，会发生什么。

NVIDIA GPU有几种计算模式，这些模式决定了GPU如何与同时运行的多个进程交互。

默认模式：这是GPU出厂时的设置。在此模式下，多个进程可以同时在一个GPU上运行。
独占进程模式：在此模式下，一个GPU一次只能被一个进程使用。这在共享工作站环境中很有用，可以防止用户相互干扰。

需要明确的是，在现代NVIDIA GPU上，多个进程共享一个GPU在默认模式下是可行的。MPS（多进程服务）旨在增强这种场景的性能，但并非引入了全新的功能。

不同的计算中心可能配置不同的默认模式。例如，在Summit超算上，默认分配的是独占进程模式，用户需要通过特定的作业标志来启用多进程共享。

性能分析实验

现在，让我们回到最初的问题：将同一个问题分解到多个MPI进程中，并在单个GPU上运行，性能会如何变化？

我进行了一个实验：将总问题大小（1024^3）平均分配给每个MPI进程，然后在单个GPU上运行不同数量的进程。

以下是测量结果（运行时间相对于单进程情况的比值）：

1个进程：基准，相对运行时间为1.0。
2个进程：相对运行时间约为1.10-1.12。
3到约20个进程：相对运行时间保持在大约1.10-1.12，没有显著恶化。

实验表明，当使用超过1个进程时，会产生大约10-12%的额外运行时间开销。但超过2个进程后，继续增加进程数（直到约20个），总运行时间基本保持稳定。

为什么需要考虑多进程？

既然没有让问题解决得更快，为什么我们还要考虑这种多进程共享GPU的场景呢？

原因在于实际的应用程序往往很复杂。许多遗留应用并未完全移植到GPU上运行，可能只有部分计算在GPU上执行，其余部分仍在CPU上。在这种情况下，你可能希望将CPU部分的工作分配到尽可能多的MPI进程中，以充分利用CPU资源，同时仍然让所有进程都能访问GPU进行计算。

默认的CUDA多进程支持使得这种混合计算模式成为可能，尽管会带来一些性能开销。

总结

本节课中我们一起学习了CUDA多进程服务的基础。我们了解到，默认情况下多个进程可以共享一个GPU，但这会引入约10%的性能开销。MPS服务的目的是优化这种多进程场景，减少开销。我们探讨了这种模式的应用场景，特别是在处理部分GPU化、部分CPU化的混合应用时，它提供了灵活的资源配置方式。

012：CUDA调试

概述

在本节课中，我们将学习如何调试CUDA程序。我们将从基础的CUDA错误管理开始，然后介绍两个重要的调试工具：Compute Sanitizer和CUDA-GDB。通过掌握这些知识和工具，你将能够更有效地识别和解决CUDA程序中的问题。

CUDA错误管理

上一节我们介绍了课程的整体目标，本节中我们来看看调试的第一步：CUDA错误管理。

所有CUDA运行时API调用都会返回一个错误码。例如，cudaSetDevice就是一个CUDA运行时API调用。任何以“CUDA”开头、且非用户自定义的函数，几乎都属于CUDA运行时API。

这些API调用返回的错误码类型是一个枚举类型，称为cudaError_t。然而，在开发中，我们经常看到开发者忽略这个错误码，例如直接调用cudaSetDevice而不检查其返回值。这是一种不严谨的做法。

如果你在CUDA代码中遇到问题，首要建议就是确保你进行了严谨的错误检查。严谨意味着捕获每一个CUDA运行时API调用的错误码，并进行适当的处理。如果返回的错误码不是cudaSuccess这个枚举值，至少应该打印出错误信息。

CUDA运行时API还提供了一个内置函数cudaGetErrorString。这个函数接收一个cudaError_t类型的值，并将其转换为人类可读的英文文本字符串。因此，最佳实践是始终检查这些错误码，并通过cudaGetErrorString打印出可读的错误信息。

内核错误检查

上一节我们介绍了如何检查CUDA运行时API的错误，本节中我们来看看如何检查内核启动的错误。

首先需要明确，CUDA内核启动是异步的。这意味着内核从主机线程（CPU代码）启动后，会与CPU线程并发执行。CPU线程不会自动获得内核何时完成执行的明确指示。

更重要的是，如果内核代码在执行过程中遇到错误，这个错误会在CPU代码后续某个不确定的时间点被报告。这可能导致一个看似无关的CUDA运行时API调用（例如cudaSetDevice）返回一个奇怪的错误，例如“misaligned address”或“illegal address”。这实际上是之前内核执行失败所遗留的错误状态。

因此，我们需要一种方法来检查内核执行是否出错。由于内核启动本身不返回错误码，我们需要使用cudaGetLastError函数。这个函数会获取上一次内核启动或CUDA运行时API调用产生的错误状态。

以下是检查内核错误的推荐方法：

在内核启动后，立即调用cudaGetLastError来捕获启动阶段的错误（例如参数配置错误）。
在内核执行完成后（通常通过cudaDeviceSynchronize或流同步来确保），再次调用cudaGetLastError来捕获内核执行过程中的错误。

通过这种方式，我们可以将错误定位到具体的操作步骤。

调试工具：Compute Sanitizer

上一节我们学习了如何通过代码进行错误检查，本节中我们来看看第一个自动化调试工具：Compute Sanitizer。

Compute Sanitizer 是一个功能强大的工具集，用于检测CUDA应用程序中的多种内存访问和同步错误。它类似于CPU上的Valgrind或AddressSanitizer工具。

以下是Compute Sanitizer提供的主要检查功能：

内存访问错误：检测越界访问、使用未初始化内存等。
竞争条件：检测多个线程对共享内存的不安全访问。
初始化错误：检测设备全局内存的初始化问题。
同步错误：检测死锁等同步问题。

使用Compute Sanitizer的基本方法是在编译时添加-lineinfo选项以保留行号信息，然后在运行时通过环境变量或命令行参数来启用特定的检查器。例如，使用compute-sanitizer --tool memcheck ./my_cuda_app来运行内存检查。

该工具会生成详细的报告，指出错误类型、发生位置（源代码文件和行号）以及相关的调用栈信息，极大地简化了内存和并发错误的调试过程。

调试工具：CUDA-GDB

上一节我们介绍了用于自动化错误检测的Compute Sanitizer，本节中我们来看看另一个强大的交互式调试工具：CUDA-GDB。

CUDA-GDB是GNU调试器（GDB）的扩展版本，专门用于调试CUDA应用程序。它允许开发者像调试CPU代码一样，交互式地调试运行在GPU上的内核代码。

以下是CUDA-GDB的一些核心功能：

设置断点：可以在主机代码或设备内核代码中设置断点。
单步执行：可以逐行执行内核代码。
检查变量：可以查看主机和设备上的变量值。
线程检查：可以检查和切换当前聚焦的CUDA线程（blockIdx, threadIdx）。
检查GPU状态：可以查询GPU的寄存器、共享内存、本地内存状态。

使用CUDA-GDB进行调试通常遵循以下流程：

使用-g -G标志编译CUDA代码（-g用于主机代码调试信息，-G用于设备代码调试信息）。
在终端启动CUDA-GDB：cuda-gdb ./my_cuda_app。
在GDB提示符下，使用标准的GDB命令（如break, run, next, print）以及CUDA扩展命令（如cuda thread, cuda block）进行调试。

通过CUDA-GDB，开发者可以深入内核内部，观察每一步的执行状态，这对于理解复杂的数据流和定位逻辑错误至关重要。

总结

本节课中我们一起学习了CUDA程序调试的核心知识。我们从基础的CUDA错误管理开始，强调了检查所有CUDA运行时API调用和内核执行状态的重要性。接着，我们介绍了两个强大的工具：Compute Sanitizer用于自动化检测内存和并发错误，CUDA-GDB用于交互式地调试内核代码逻辑。掌握这些方法和工具，将帮助你更高效地开发和排除CUDA应用程序中的故障。

013：CUDA Graphs 🚀

在本节课中，我们将要学习CUDA Graphs，这是CUDA中一项用于提升性能的特性。我们将专注于其核心概念，帮助初学者理解如何利用CUDA Graphs来优化程序执行。

概述

CUDA Graphs是一种将一系列CUDA操作（如内核启动、内存拷贝）及其依赖关系定义为独立于执行过程的结构。通过一次性定义并重复执行这个“图”，可以显著减少CPU的启动开销，从而提升整体性能，尤其是在执行大量短时内核时。

什么是CUDA Graphs？

CUDA Graphs是一系列由依赖关系连接的操作节点，其定义与执行过程是分离的。所有传统的CUDA工作流本质上都隐式地形成了一个执行图。

例如，一个包含三个内核启动的基本CUDA代码，在默认流中会形成一个线性的、由流顺序隐式定义依赖关系的图。

从左侧的传统流式执行，到右侧的图表示，我们可以看到CUDA Graphs是如何将工作流抽象为节点和依赖关系的。

图的节点与依赖

构成CUDA Graph的节点可以是任何异步CUDA操作。

以下是节点类型的示例：

内核启动（Kernel Launch）
内存拷贝（Memory Copy）
事件（Events）
CUDA回调函数（Callback Functions）

节点之间的依赖关系可以通过两种方式建立：

隐式依赖：由流中操作的顺序执行决定。
显式依赖：通过cudaStreamWaitEvent和事件信号来明确指定。

此外，图还可以包含子图，这增加了构建复杂工作流的灵活性。

图的优势：定义一次，重复执行

CUDA Graphs的核心优势在于，它可以被定义一次，然后重复执行多次。

这意味着我们只需支付一次图的构建成本，之后仅需一个简单的执行操作即可启动整个可能非常复杂的图。相比之下，传统的流式工作流需要逐个启动操作，并且每次想重新运行整个工作流时都需要重新提交所有命令。

当需要多次重复执行相同的工作流时，这种模式能有效回收CPU时间，实现延迟隐藏。

性能对比：流 vs. 图

下面的图示对比了两种执行模式。

在流式执行（上方）中，CPU需要忙碌地逐个启动内核，而GPU可能在等待下一个启动指令时处于空闲状态。

在图执行（下方）中，通过一次cudaGraphLaunch调用，整个图被隐式地提交执行，从而释放了CPU去处理其他任务，减少了GPU的空闲等待时间。

当内核执行时间非常短（接近甚至小于启动开销）时，使用图来减少累积的启动延迟尤为重要。启动开销通常在几微秒量级。

需要明确的是，使用CUDA Graphs并不会改变内核在GPU上的纯运行性能，其主要目标是降低程序的综合启动成本。

CUDA Graphs的三阶段执行模型

CUDA Graphs的执行遵循一个清晰的三阶段模型：

定义（Define）：在此阶段，我们构建图的结构。我们指定节点（操作）、参数、依赖关系、目标GPU和启动流等信息。
实例化（Instantiate）：将定义好的图转换为一个可执行的“图实例”。此过程会固定图的参数，为执行做准备。
执行（Execute）：将实例化的图启动到指定的CUDA流中运行。

总结

本节课我们一起学习了CUDA Graphs的基础知识。我们了解到，CUDA Graphs通过将一系列操作及其依赖预定义为图，实现了“定义一次，重复执行”的模式。这能有效减少CPU的启动开销，尤其有利于优化由大量短时内核构成的工作流，从而提升应用程序的整体性能。其核心的三阶段模型（定义、实例化、执行）为管理复杂的GPU操作提供了清晰的框架。

posted @ 2026-03-26 13:20 布客飞龙V 阅读(41) 评论(0) 收藏举报

刷新页面返回顶部

龙哥盟