LLM生成的CUDA C++ Programming Guide 摘要

CUDA C++编程指南介绍了CUDA模型和接口，强调了在版本12.4中的更新。关键知识点包括：

1. GPU的优势：与CPU相比，GPU在相同价格和功耗范围内提供更高的指令吞吐量和内存带宽。这使得许多应用在GPU上运行速度更快，特别是在高度并行计算场景下。

2. GPU与CPU设计差异：CPU设计用于快速执行单一线程操作，并行处理少数线程；而GPU设计用于同时高效执行数千个线程，通过大量并行计算来抵消单线程性能较慢的问题，从而实现更高的吞吐量。

3. CUDA介绍：自2006年NVIDIA推出CUDA以来，它作为一种利用GPU并行计算引擎的通用并行计算平台和编程模型，使开发者能够以C++等高级语言编写程序，解决复杂计算问题，效率高于CPU。

4. 可扩展的编程模型：CUDA编程模型通过线程组层次结构、共享内存和屏障同步三个核心抽象，简化了并行编程，支持细粒度和粗粒度的数据及任务并行性，实现了自动可扩展性和跨不同数量多核处理器的透明执行。

5. 编程模型概述：文档结构中包含对编程模型的深入介绍，特别是如何在C++中表达CUDA概念，如定义并行执行的内核函数（kernels），使用特定语法指定执行配置，以及如何通过内建变量识别每个独立线程。
   综上所述，GPU优化的关键在于理解并利用其高度并行架构，通过CUDA这样的平台，采用适合并行处理的算法和数据结构，以及有效管理内存访问和线程协作，以达到性能最大化。

6. 线程层次结构：线程通过一个三维向量threadIdx来标识，这允许使用一维、二维或三维线程索引来形成线程块（thread block），自然地跨向量、矩阵或体积等域执行计算。

7. 线程块限制：每个线程块中的线程数量受限，因为它们期望驻留在同一个流式多处理器核心上，并且必须共享该核心有限的内存资源。当前GPU中，线程块最多可包含1024个线程。

8. 网格（Grid）组织：线程块组成一维、二维或三维的网格。网格中的线程块数量通常由待处理数据的大小决定，这通常超过系统中的处理器数量。

9. 独立执行的线程块：要求线程块独立执行，意味着可以以任意顺序并行或串行执行，便于跨任何数量的核心调度，从而编写可随核心数扩展的代码。

10. 线程间合作：通过共享内存和同步执行协调内存访问，线程块内的线程可以协同工作。__syncthreads()函数作为屏障同步块内所有线程，而Cooperative Groups API提供了丰富的线程同步原语。

11. 线程块集群（Thread Block Clusters）：从NVIDIA Compute Capability 9.0起，引入了线程块集群的概念，保证集群内的线程块在GPU处理集群（GPC）上协同调度。集群支持动态和编译时定义的大小，并且集群内的线程块可以使用分布式共享内存进行协作。

12. 内存层次结构：CUDA线程可以访问多种内存空间，包括私有本地内存、块内共享内存、全局内存、常量内存和纹理内存。不同内存空间针对不同类型的内存访问进行了优化，纹理内存还支持特定数据格式的地址模式和数据过滤。
    这些知识点对于理解如何在GPU上设计高效的并行算法和优化内存访问至关重要，有助于开发者充分利用GPU的并行计算能力。

异构编程（Heterogeneous Programming）

CUDA编程模型将CUDA线程视为在物理上与运行C++程序的主机分离的设备上执行，例如，在GPU上执行内核而其余C++程序在CPU上执行。该模型假设主机和设备各自在DRAM中维护独立的内存空间，即主机内存和设备内存。程序通过调用CUDA运行时来管理对内核可见的全局、常量和纹理内存空间，包括设备内存的分配与回收以及主机与设备内存间的数据传输。统一内存提供托管内存以桥接主机和设备内存空间，实现所有CPU和GPU对单一、连贯内存映像的访问，并支持设备内存的超订阅，简化应用移植。

异步SIMT编程模型（Asynchronous SIMT Programming Model）

CUDA编程模型允许从NVIDIA Ampere架构起始的设备通过异步编程模型加速内存操作。该模型定义了CUDA线程间的异步屏障同步行为，以及如何使用cuda::memcpy_async在GPU计算同时异步移动全局内存数据。

异步操作（Asynchronous Operations）

异步操作由CUDA线程发起并异步执行，如同由另一线程执行。良好的程序设计中，一个或多个CUDA线程会与此异步操作同步。异步操作使用同步对象来同步操作完成，这些对象可以由用户显式管理（如cuda::memcpy_async）或库内部隐式管理（如cooperative_groups::memcpy_async）。同步对象如cuda::barrier和cuda::pipeline可在不同线程作用域使用，用于定义可使用同步对象进行同步的线程集。

计算能力（Compute Capability）

设备的计算能力由版本号表示，标识GPU硬件支持的功能。版本号由主要修订号X和次要修订号Y组成，表示为X.Y。相同主要修订号的设备属于同一核心架构。从CUDA 7.0和CUDA 9.0起，Tesla和Fermi架构不再受支持。

编程接口（Programming Interface）

CUDA C++通过C++语言的最小扩展集，以及运行时库，为熟悉C++的用户提供编写设备执行程序的简单途径。核心语言扩展允许定义内核函数、指定网格和块维度等。运行时库提供了在主机上执行的内存管理、数据传输等功能的C/C++函数。此外，还介绍了基于更低级CUDA驱动API的编译流程，使用nvcc编译器将内核代码编译为设备上可执行的二进制代码。

1. 即时编译（Just-in-Time Compilation）：应用程序运行时加载的PTX代码会被设备驱动进一步编译为二进制代码，这称为即时编译。即时编译虽增加了应用加载时间，但能让应用受益于设备驱动内置的新编译器改进。它也是应用在编译时尚不存在的设备上运行的唯一途径。
2. 二进制兼容性：二进制代码针对特定架构生成，确保了从小版本到下一个相邻小版本的兼容性，但不保证跨大版本或向后兼容。例如，为计算能力X.y生成的cubin对象仅能在计算能力X.z（z≥y）的设备上执行。
3. PTX兼容性：某些PTX指令仅支持较高计算能力的设备。PTX代码针对特定计算能力生成后，可编译为不小于该计算能力的二进制代码，但基于较早PTX版本的二进制可能无法利用新硬件特性，如Tensor Core指令。
4. 应用兼容性：应用需加载与目标设备计算能力兼容的二进制或PTX代码。为了能在未来更高计算能力的架构上执行代码，应用需加载将被即时编译的PTX代码。
5. C++兼容性：编译器前端根据C++语法规则处理CUDA源文件，主机代码支持完整C++，而设备代码仅支持C++的子集。
6. 64位兼容性：nvcc的64位版本以64位模式编译设备代码，要求主机代码也以64位模式编译。
7. CUDA运行时：运行时库(cudart)提供管理设备内存、共享内存、页面锁定的主机内存、异步并发执行、多设备系统操作等函数，所有入口点前缀为cuda，且强调了错误检查和调用堆栈管理的重要性。
   这些内容覆盖了从代码编译、兼容性管理到运行时功能使用的关键方面，对进行GPU优化的开发者来说是基础且重要的知识。

初始化（Initialization）

自CUDA 12.0起，cudaInitDevice()和cudaSetDevice()调用负责初始化运行时系统及与指定设备关联的主上下文。若未进行这些调用，运行时将默认使用设备0并在处理其他API请求时按需自我初始化。这在计时运行时函数调用和解释首次进入运行时的错误代码时需注意。在12.0之前，cudaSetDevice()不初始化运行时，应用通常使用无操作调用cudaFree(0)来隔离运行时初始化与其他API活动，以实现更精确的计时和错误处理。

设备内存（Device Memory）

CUDA编程模型假设系统由主机和设备组成，每部分拥有独立的内存。内核在设备内存上执行，因此运行时提供了分配、释放设备内存以及在主机和设备内存间传输数据的功能。设备内存可作为线性内存或CUDA数组分配。CUDA数组是针对纹理获取优化的不透明内存布局。线性内存则在一个统一地址空间中分配，便于指针引用，其大小取决于主机系统和GPU的计算能力。

设备内存L2访问管理（Device Memory L2 Access Management）

从CUDA 11.0开始，计算能力8.0及以上的设备能够影响全局内存中数据在L2缓存中的持久性，从而可能提供更高的带宽和更低的延迟访问。L2缓存的一部分可以被预留用于持久化访问，这类访问对这部分缓存有优先使用权。L2预留大小可在一定范围内调整，并且在多实例GPU(MIG)模式下禁用。当使用多进程服务(MPS)时，L2预留大小不能通过cudaDeviceSetLimit改变，而是在MPS服务器启动时通过环境变量设置。

L2策略与访问属性（L2 Policy and Access Properties）

访问策略窗口定义了一个全局内存的连续区域及其在L2缓存中的持久性属性。通过CUDA流或CUDA图节点可以设置L2持久访问窗口，hitRatio参数用于指定获得持久化属性的访问比例，有助于避免缓存行抖动并减少L2缓存的数据移动量。

L2持久性示例（L2 Persistence Example）

示例展示了如何为持久化访问预留L2缓存，在CUDA流中利用预留的L2缓存，以及之后重置L2缓存的过程，详细说明了如何配置访问策略以优化内存访问性能。

1. 重置L2缓存访问为正常状态：之前的CUDA内核在L2缓存中的持久化缓存行可能长时间保留，即使已不再使用。因此，重置L2缓存到正常状态对于流式或普通内存访问利用具有正常优先级的L2缓存至关重要。存在三种方法可以将持久化访问重置为正常状态。
2. 管理L2预留缓存的利用率：多个并发在不同CUDA流上的CUDA内核可能被分配了不同的访问策略窗口，但它们共享L2预留缓存部分。因此，该预留缓存部分的总利用率是所有并发内核个别使用的总和。随着持久化访问量超过预留L2缓存容量，指定内存访问为持久化的益处会减少。应用需考虑如何有效管理这一预留缓存的使用。
3. 查询L2缓存属性：L2缓存相关的属性是cudaDeviceProp结构体的一部分，可以通过CUDA运行时API cudaGetDeviceProperties查询。
4. 控制用于持久化内存访问的L2缓存预留大小：通过CUDA运行时API cudaDeviceGetLimit查询和 cudaDeviceSetLimit设置用于持久化内存访问的L2预留缓存大小，最大值由cudaDeviceProp::persistingL2CacheMaxSize给出。
5. 共享内存：作为线程层次结构的一部分，共享内存通过__shared__内存空间指定符分配。它比全局内存快得多，可作为加速计算的暂存区，以减少全局内存访问。示例代码展示了不使用和使用共享内存的矩阵乘法实现，后者显著减少了对全局内存的访问次数。
6. 分布式共享内存：计算能力9.0及以后的设备引入了线程块集群，使集群内的线程能够访问参与集群的所有线程块的共享内存。这称为分布式共享内存，提供了额外的内存访问机制，可以基于线程块集群的需求动态调整，适用于超出单个线程块共享内存限制的场景，如大规模直方图计算。
7. 页面锁定的主机内存：通过CUDA运行时提供的函数，可以使用页面锁定（或固定）的主机内存，相比常规分页内存，它有更低的CPU-GPU数据传输延迟，支持零拷贝操作，并允许在系统中的任何设备上使用。为了跨所有设备提供这些优势，需要通过特定标志(cudaHostAllocPortable 或 cudaHostRegisterPortable)分配或注册页面锁定内存。
   ####在GPU优化领域，有几个关键知识点值得关注：
8. 写组合内存（Write-Combining Memory）：通过在cudaHostAlloc()调用中传递cudaHostAllocWriteCombined标志，可以将默认为缓存可访问的锁页主机内存改为写组合类型。这释放了主机的L1和L2缓存资源，使应用的其他部分能利用更多缓存，并且在PCI Express总线上传输时不被窥探，从而可能提升高达40%的传输性能。但主机直接读取写组合内存的速度非常慢，因此应主要应用于主机仅写入的内存。避免在WC内存上使用CPU原子指令，因为不是所有CPU都保证此功能。
9. 映射内存（Mapped Memory）：通过在cudaHostAlloc()或cudaHostRegister()中传递cudaHostAllocMapped或cudaHostRegisterMapped标志，可以使锁页主机内存块映射到设备地址空间中。这样，该内存块通常具有两个地址：一个在主机内存中，另一个在设备内存中，后者可通过cudaHostGetDevicePointer()获取并在内核中使用。直接从内核访问主机内存虽不如访问设备内存带宽高，但有其优势。由于映射的锁页内存是主机和设备共享的，必须使用流或事件同步内存访问以避免冲突。要获取任何映射锁页内存的设备指针，必须先启用页面锁定内存映射。映射锁页主机内存的原子操作并非从主机或其他设备的角度来看是原子性的。
10. 内存同步域（Memory Synchronization Domains）：自Hopper架构GPU及CUDA 12.0起，内存同步域特性有助于减轻内存栅栏操作导致的干扰问题。每个内核启动被赋予一个域ID，写操作和栅栏操作都标记这个ID，使得栅栏仅对匹配其域的写操作进行排序。在不同域间需要系统级同步，而同一域内设备级同步仍然足够。这要求跨域通信提前到系统级刷新，以满足累积性。
    这些概念对于优化GPU应用程序中的内存管理和数据传输至关重要，特别是在需要高效内存使用和减少数据传输延迟的场景下。

3.2.7.3. CUDA中的域使用

• 域访问：通过新的启动属性cudaLaunchAttributeMemSyncDomain和cudaLaunchAttributeMemSyncDomainMap访问。
• 域选择：逻辑域包括cudaLaunchMemSyncDomainDefault（默认）和cudaLaunchMemSyncDomainRemote（远程），后者用于隔离执行远程内存访问的内核的内存流量。
• 域映射：cudaLaunchAttributeMemSyncDomainMap提供逻辑到物理域的映射，有助于应用程序架构的灵活性。
• Hopper架构特性：Hopper架构有4个域，而CUDA在Hopper之前的设备上报告的域计数为1，以支持可移植代码。
• 默认行为：未设置时，逻辑域默认为默认域；默认映射将默认域映射到0，远程域映射到1（多于1个域的GPU）。
• 库集成示例：NCCL 2.16及以上版本将在CUDA 12.0及以后版本中使用远程域标记启动。

3.2.8. 异步并发执行

• 并发操作：CUDA支持主机与设备、内核之间以及数据传输的异步并发执行。
• 并发限制：并发程度取决于设备的功能集和计算能力。
• 主机与设备并发：通过异步库函数实现，如异步内存拷贝，允许在设备操作完成前释放主机线程控制权。
• 并发内核执行：计算能力2.x及以上的某些设备支持。限制包括不同CUDA上下文内的内核不能同时执行，需启用多进程服务（MPS）来实现跨进程的SM并行。
• 数据传输与内核执行重叠：支持的设备可通过异步引擎进行内存拷贝与内核执行的同时进行，涉及主机内存时需锁定页面。
• 并发数据传输：计算能力2.x及以上设备可能支持数据传输间的重叠，同样要求涉及的主机内存页锁定。
• 流(Streams)：管理并发操作的机制，命令序列按顺序执行，不同流间可以交错或并行执行。流的创建和销毁通过cudaStreamCreate()和cudaStreamDestroy()进行，支持命令的依赖管理和同步。

应用启示

这些知识点对于GPU优化至关重要，特别是在设计高性能计算应用时，通过合理配置域映射、利用异步执行模式以及流管理，可以显著提升数据处理的效率和并发性，减少等待时间，提升整体应用性能。特别是对于需要大量数据交换和复杂计算的应用场景，精细控制内存同步域和并发执行策略是提高系统吞吐量的关键。

1. 默认流（Default Stream）：未指定流参数的内核启动和主机-设备内存拷贝，默认使用默认流执行，并保持顺序执行。通过编译选项--default-stream可设定默认流为每个线程独立的常规流或所有线程共用的特殊NULL流，后者会隐式同步。

2. 显式同步（Explicit Synchronization）：包括cudaDeviceSynchronize()等待所有流命令完成，cudaStreamSynchronize()等待特定流的命令完成，cudaStreamWaitEvent()使流命令等待事件完成，以及cudaStreamQuery()查询流中所有前序命令是否完成。

3. 隐式同步（Implicit Synchronization）：某些操作（如依赖检查、事件记录等）会导致不同流间的命令不能并发执行，应用需遵循指南以提升并发内核执行潜力。

4. 重叠行为（Overlapping Behavior）：两个流间的执行重叠程度取决于命令发出的顺序及设备对数据传输与内核执行重叠、并发内核执行和并发数据传输的支持情况。

5. 主机函数（Host Functions）：通过cudaLaunchHostFunc()可在流中插入CPU函数调用，该函数在流中所有之前命令完成后执行，且后续命令需等待该函数完成才开始。

6. 流优先级（Stream Priorities）：通过cudaStreamCreateWithPriority()创建时可指定流的优先级，高优先级流的工作将优先于低优先级流执行。

7. 程序化依赖启动与同步（Programmatic Dependent Launch and Synchronization）：针对计算能力9.0及以上设备，允许依赖于主内核结果的次内核在主内核完成前启动，利用内核执行中的非依赖部分实现并发，以提高性能，并引入新的API支持这一机制，减少启动延迟。

8. 程序化依赖启动（Programmatic Dependent Launch）：在CUDA编程中，允许一个主核函数（primary kernel）和一个次核函数（secondary kernel）在同一CUDA流中启动。主核函数需在所有线程块执行完毕后，通过调用cudaTriggerProgrammaticLaunchCompletion来表明准备就绪，以便次核函数启动。次核函数必须使用可扩展启动API，并可设置cudaLaunchAttributeProgrammaticStreamSerialization属性，使得CUDA驱动能够在不等待主核函数完成及其内存刷新的情况下提前启动次核函数。若主核函数未显式触发，该触发会在所有主核函数线程块退出后隐式发生。在这种模式下，为确保数据一致性，次核函数必须使用cudaGridDependencySynchronize或其他机制来同步来自主核函数的结果数据。

9. CUDA图中的应用：程序化依赖启动也可应用于CUDA图中，通过流捕获或直接利用边数据实现。在CUDA图的两个核函数节点之间使用cudaGraphDependencyTypeProgrammatic类型的边，可以使得上游核函数对下游核函数中的cudaGridDependencySynchronize()可见。此类型要求使用特定的输出端口，如cudaGraphKernelNodePortLaunchCompletion或cudaGraphKernelNodePortProgrammatic。

10. CUDA图概述：CUDA图作为一种新的工作提交模型，允许将操作（如核函数启动）及其间的依赖关系定义为独立于执行流程的序列，从而实现定义一次、多次执行的效率提升。图的定义与执行分离，能够减少CPU启动开销，并为CUDA提供整个工作流程视图以进行潜在的优化，尤其对于短时执行的GPU核函数，能够显著降低总体执行时间的开销部分。

11. 图结构与创建：CUDA图由操作节点和它们之间的依赖边组成，这些依赖关系限制了操作的执行顺序。图的创建可以通过显式API或流捕获机制实现，后者能够将现有基于流的API代码段转换为图。流捕获通过cudaStreamBeginCapture和cudaStreamEndCapture函数实现，期间向流中添加的工作被记录为图的一部分而非立即执行。

12. 边数据（Edge Data）：自CUDA 12.3起，图中的边可以携带数据，用于修改依赖关系的行为，包括指定输出端口（触发时机）、输入端口（依赖哪部分节点）以及依赖类型。这为图中的依赖关系提供了更多控制，例如通过cudaGraphDependencyTypeProgrammatic类型支持核函数间的程序化依赖启动。边数据可以在创建和查询图的API中使用，同时也适用于某些流捕获API。

13. 跨流依赖处理：流捕获能够管理通过cudaEventRecord()和cudaStreamWaitEvent()表达的跨流依赖，前提是等待的事件被记录在同一捕获图中。当事件在捕获模式的流中被记录时，它会生成一个捕获事件，该事件代表捕获图中的多个节点集。如果流等待一个被捕获的事件，则会将该流置于捕获模式，并对后续操作添加对该捕获事件节点的额外依赖。所有相关的流最终都会被合并回最初调用cudaStreamBeginCapture()的源头流。

14. 禁止和未处理的操作：在流捕获期间，同步或查询处于捕获状态的流或捕获事件的执行状态是无效的，因为它们不代表已调度执行的项目。此外，在任何关联流正在进行捕获的情况下，使用遗留流（默认流）也是非法的。尝试合并两个不同的捕获图，或者在没有指定cudaEventWaitExternal标志的情况下从一个被捕获的流等待非捕获事件也是不允许的。某些异步API在流捕获模式下不支持，如cudaStreamAttachMemAsync()。

15. 失效处理：若在流捕获过程中尝试执行非法操作，相关的捕获图会被失效。一旦捕获图失效，继续使用相关联的正在捕获的流或捕获事件会返回错误，直到通过cudaStreamEndCapture()结束捕获并使流脱离捕获模式。

16. CUDA用户对象：用于帮助管理CUDA异步工作使用的资源生命周期，特别是在CUDA Graphs和流捕获场景中。用户对象通过与内部引用计数关联的用户定义析构回调，类似于C++的shared_ptr。资源引用可以由CPU端的用户代码和CUDA图持有。CUDA自动管理与图关联的引用，包括克隆、实例化以及销毁过程中的引用管理。用户对象通过cudaUserObjectCreate创建，并提供了一种手动信号同步对象的方式，但不允许在析构函数中直接调用CUDA API以避免阻塞CUDA内部线程。
    这些知识点对于优化GPU程序中的异步执行流程、资源管理和避免潜在的执行错误至关重要。
    ####在使用图进行工作提交时，过程分为三个阶段：定义、实例化和执行。当工作流程不变时，通过多次执行来分摊定义和实例化的开销，此时图相比流提供了明显优势。图是对工作流程（包括内核、参数和依赖关系）的快照，以便快速高效地重放。若工作流程改变，则图过时需修改；结构重大变化（如拓扑或节点类型）需重新实例化源图，因为需要重新应用与拓扑相关的优化技术。
    重复实例化的成本会减少图执行的整体性能优势，但常见情况是仅节点参数（如内核参数和cudaMemcpy地址）改变而图拓扑保持不变。为此，CUDA提供了“图更新”的轻量级机制，允许在不重建整个图的情况下就地修改某些节点参数，这比重新实例化更高效。更新在下次图启动时生效，不影响之前的图启动，即使它们在更新时正在运行。图可反复更新和重新启动，因此多个更新/启动可以在一个流上排队。
    CUDA提供两种更新已实例化图参数的机制：整体图更新和单个节点更新。整体图更新允许用户提供一个拓扑相同但节点包含更新参数的cudaGraph_t对象。单个节点更新允许用户显式更新单个节点的参数。当大量节点被更新或调用者不了解图拓扑（例如，图由库调用的流捕获产生）时，使用更新的cudaGraph_t更方便。当更改数量少且用户拥有需要更新的节点句柄时，首选单个节点更新，因为它跳过了未更改节点的拓扑检查和比较，在许多情况下效率更高。CUDA还提供了启用和禁用单个节点而不影响其当前参数的机制。
    图更新存在一些限制，主要针对特定类型的节点（如kernel节点、cudaMemset和cudaMemcpy节点），并且外部信号量等待节点和记录节点、条件节点也有特定限制，而主机节点、事件记录节点或事件等待节点的更新则不受限制。
    整体图更新通过cudaGraphExecUpdate()函数实现，要求更新图在拓扑上与原始图完全相同，包括依赖关系指定的顺序。为了确保sink节点（无依赖关系的节点）的一致排序，CUDA依赖特定API调用的顺序。
    单个节点更新允许直接更新已实例化图中的节点参数，消除了实例化和创建新cudaGraph_t的开销。如果需要更新的节点数量相对较少，单独更新节点更优。CUDA提供了专门的API来直接启用或禁用节点（如cudaGraphNodeSetEnabled()），以及查询节点状态。
    设备图启动允许在设备端发起图执行，适用于需要根据运行时数据依赖做出决策的工作流程，支持统一寻址的系统可以使用此功能。设备图可从主机和设备启动，而主机图只能从主机启动；设备图在设备上不能同时启动两次，同时从主机和设备启动行为未定义。

17. 设备图创建：要从设备上启动图，需通过在cudaGraphInstantiate()调用中传递cudaGraphInstantiateFlagDeviceLaunch标志来显式实例化。设备图的结构在实例化时固定，更新需要重新实例化，且只能在主机上执行。

18. 图上传：设备图执行前需上传至设备以准备必要资源。这可通过cudaGraphUpload()或在cudaGraphInstantiateWithParams()中请求上传实现，也可通过首先从主机启动图隐式完成上传。

19. 设备图更新与重上传：设备图仅能在主机上更新，并在更新可执行图后需重新上传到设备以应用变化。

20. 设备端启动：设备图支持从主机和设备启动，使用相同的cudaGraphLaunch()签名。设备上启动时必须来自另一个图，且为线程级操作，允许多线程并行启动。

21. 流管理：设备图不能在常规CUDA流中启动，只能在表示特定启动模式的命名流中启动，包括“fire and forget”（立即提交执行）模式。

22. 执行环境与同步：设备图启动会产生独立的执行环境，封装所有工作及生成的子工作。理解设备端同步模型需了解执行环境概念。主机启动图时，存在一个流环境作为父级。

23. 尾部启动（Tail Launch）：作为替代传统同步方法（如cudaDeviceSynchronize()）的机制，用于实现序列工作依赖。当图及其所有子图完成时，尾部启动的图将按顺序执行。

24. 条件节点：允许图中的条件执行和循环，支持动态迭代工作流程的图形化表示，提升CPU并行性。条件节点包含一个条件处理程序，用于评估是否执行其包含的图。
    这些要点概述了利用GPU设备图进行优化的关键方面，包括图的创建、资源管理、执行控制以及高级功能如条件执行的支持。

25. 条件IF节点（Conditional IF Nodes）：在图执行过程中，如果条件非零，则IF节点的主体图将执行一次。条件默认值通过上游内核设定，条件主体则利用图API进行填充。

26. 条件WHILE节点（Conditional WHILE Nodes）：WHILE节点的主体图将一直执行直至条件变为零。条件在节点执行时及主体图完成后被评估。同样，条件主体的构建也依赖于图API。

27. 事件（Events）：运行时允许应用异步记录程序中的事件以密切监控设备进度和进行精确计时。事件完成标志着其前所有任务或指定流中的所有命令已完成。流零中的事件在所有流的所有前置任务和命令完成后完成。

28. 事件的创建与销毁：示例代码展示了如何创建和销毁两个事件。

29. 持续时间测量：创建的事件可用来测量代码段的执行时间。

30. 同步调用（Synchronous Calls）：当调用同步函数时，主机线程会等待设备完成请求任务后才恢复控制。主机线程的行为（如让出、阻塞或自旋）可通过cudaSetDeviceFlags()预先设定。

31. 多设备系统：

    • 设备枚举：主机系统可能有多台设备，示例代码演示了如何枚举这些设备及其属性查询。
    • 设备选择：通过cudaSetDevice()，主机线程可随时更改当前操作的设备。
    • 流和事件行为：说明了不同设备间的流、事件交互规则，包括失败和成功的情况。
    • 对等内存访问（Peer-to-Peer Memory Access）：特定系统配置下，设备间可以直接访问对方内存，需要通过cudaDeviceCanAccessPeer()和cudaDeviceEnablePeerAccess()启用。
    • 统一虚拟地址空间（Unified Virtual Address Space）：在64位进程中，主机和计算能力2.0及以上设备共享单一虚拟地址空间，简化了内存管理和访问。
      以上摘要覆盖了GPU编程和优化中关于条件控制流、事件管理、同步操作、多设备资源管理和内存访问的关键概念。

32. 进程间通信（Interprocess Communication, IPC）：

    • GPU内存指针和事件句柄在同一流程内的线程间可直接引用，但不支持跨进程直接引用。
    • 要实现跨进程共享，需使用IPC API，该API仅支持Linux系统的64位进程及计算能力2.0以上的设备。注意，cudaMallocManaged分配的内存不支持IPC API。
    • 应用程序可通过cudaIpcGetMemHandle()获取内存指针的IPC句柄，通过标准IPC机制传递，并用cudaIpcOpenMemHandle()在其他进程中恢复有效指针。
    • 为防止信息泄露，推荐只共享大小为2MiB对齐的内存块。

33. 错误检查：

    • 所有运行时函数会返回错误码，但异步函数的错误码无法立即反映设备上的异步错误，需通过cudaDeviceSynchronize()同步后检查。
    • 运行时为每个主机线程维护一个错误变量，可通过cudaPeekAtLastError()查看而不清空，或用cudaGetLastError()查看并重置为成功状态。
    • 核心启动（kernel launch）不返回错误码，需在启动后立即调用错误检查函数以捕获预启动错误，并确保在检查前调用cudaGetLastError()重置错误变量。

34. 调用栈：

    • 计算能力2.x及以上设备可查询和设置调用栈大小。
    • 调用栈溢出会导致内核调用失败，调试模式下显示堆栈溢出错误，否则为未指定的启动错误。
    • 编译器无法静态确定栈大小时会发出警告，此时需要手动设置栈大小。

35. 纹理与表面内存：

    • 纹理与表面内存访问利用GPU图形文本硬件子集，相比全局内存访问，可带来性能优势。
    • 纹理对象API允许创建和管理纹理对象，控制访问模式、格式等。
    • 支持16位浮点纹理处理，需要通过特定函数进行类型转换。
    • 分层纹理提供一种组织纹理数据的方式，适用于一维或二维纹理数组，且支持在单一层内进行纹理过滤。
      以上内容覆盖了GPU优化中的关键概念，包括跨进程资源共享、错误管理和调试、程序调用栈管理以及如何利用纹理内存来提升数据读取性能。

36. 立方体贴图纹理（Cubemap Textures）：这是一种特殊的二维分层纹理，包含六层，分别代表立方体的六个面。它只能通过带有cudaArrayCubemap标志的cudaMalloc3DArray()函数创建。立方体贴图纹理通过texCubemap()函数访问，并且要求设备计算能力为2.0或更高。

37. 立方体贴图分层纹理（Cubemap Layered Textures）：这类纹理的每一层都是相同维度的立方体贴图。访问时使用一个整数索引和三个浮点纹理坐标，索引指定位子立方体贴图中的序列，坐标则定位该立方体贴图内的像素。它们需要通过带有cudaArrayLayered和cudaArrayCubemap标志的cudaMalloc3