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

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

GPU的优势：与CPU相比，GPU在相同价格和功耗范围内提供更高的指令吞吐量和内存带宽。这使得许多应用在GPU上运行速度更快，特别是在高度并行计算场景下。
GPU与CPU设计差异：CPU设计用于快速执行单一线程操作，并行处理少数线程；而GPU设计用于同时高效执行数千个线程，通过大量并行计算来抵消单线程性能较慢的问题，从而实现更高的吞吐量。
CUDA介绍：自2006年NVIDIA推出CUDA以来，它作为一种利用GPU并行计算引擎的通用并行计算平台和编程模型，使开发者能够以C++等高级语言编写程序，解决复杂计算问题，效率高于CPU。
可扩展的编程模型：CUDA编程模型通过线程组层次结构、共享内存和屏障同步三个核心抽象，简化了并行编程，支持细粒度和粗粒度的数据及任务并行性，实现了自动可扩展性和跨不同数量多核处理器的透明执行。
编程模型概述：文档结构中包含对编程模型的深入介绍，特别是如何在C++中表达CUDA概念，如定义并行执行的内核函数（kernels），使用特定语法指定执行配置，以及如何通过内建变量识别每个独立线程。
综上所述，GPU优化的关键在于理解并利用其高度并行架构，通过CUDA这样的平台，采用适合并行处理的算法和数据结构，以及有效管理内存访问和线程协作，以达到性能最大化。
线程层次结构：线程通过一个三维向量threadIdx来标识，这允许使用一维、二维或三维线程索引来形成线程块（thread block），自然地跨向量、矩阵或体积等域执行计算。
线程块限制：每个线程块中的线程数量受限，因为它们期望驻留在同一个流式多处理器核心上，并且必须共享该核心有限的内存资源。当前GPU中，线程块最多可包含1024个线程。
网格（Grid）组织：线程块组成一维、二维或三维的网格。网格中的线程块数量通常由待处理数据的大小决定，这通常超过系统中的处理器数量。
独立执行的线程块：要求线程块独立执行，意味着可以以任意顺序并行或串行执行，便于跨任何数量的核心调度，从而编写可随核心数扩展的代码。
线程间合作：通过共享内存和同步执行协调内存访问，线程块内的线程可以协同工作。__syncthreads()函数作为屏障同步块内所有线程，而Cooperative Groups API提供了丰富的线程同步原语。
线程块集群（Thread Block Clusters）：从NVIDIA Compute Capability 9.0起，引入了线程块集群的概念，保证集群内的线程块在GPU处理集群（GPC）上协同调度。集群支持动态和编译时定义的大小，并且集群内的线程块可以使用分布式共享内存进行协作。
内存层次结构：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编译器将内核代码编译为设备上可执行的二进制代码。

即时编译（Just-in-Time Compilation）：应用程序运行时加载的PTX代码会被设备驱动进一步编译为二进制代码，这称为即时编译。即时编译虽增加了应用加载时间，但能让应用受益于设备驱动内置的新编译器改进。它也是应用在编译时尚不存在的设备上运行的唯一途径。
二进制兼容性：二进制代码针对特定架构生成，确保了从小版本到下一个相邻小版本的兼容性，但不保证跨大版本或向后兼容。例如，为计算能力X.y生成的cubin对象仅能在计算能力X.z（z≥y）的设备上执行。
PTX兼容性：某些PTX指令仅支持较高计算能力的设备。PTX代码针对特定计算能力生成后，可编译为不小于该计算能力的二进制代码，但基于较早PTX版本的二进制可能无法利用新硬件特性，如Tensor Core指令。
应用兼容性：应用需加载与目标设备计算能力兼容的二进制或PTX代码。为了能在未来更高计算能力的架构上执行代码，应用需加载将被即时编译的PTX代码。
C++兼容性：编译器前端根据C++语法规则处理CUDA源文件，主机代码支持完整C++，而设备代码仅支持C++的子集。
64位兼容性：nvcc的64位版本以64位模式编译设备代码，要求主机代码也以64位模式编译。
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缓存的过程，详细说明了如何配置访问策略以优化内存访问性能。

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

默认流（Default Stream）：未指定流参数的内核启动和主机-设备内存拷贝，默认使用默认流执行，并保持顺序执行。通过编译选项--default-stream可设定默认流为每个线程独立的常规流或所有线程共用的特殊NULL流，后者会隐式同步。
显式同步（Explicit Synchronization）：包括cudaDeviceSynchronize()等待所有流命令完成，cudaStreamSynchronize()等待特定流的命令完成，cudaStreamWaitEvent()使流命令等待事件完成，以及cudaStreamQuery()查询流中所有前序命令是否完成。
隐式同步（Implicit Synchronization）：某些操作（如依赖检查、事件记录等）会导致不同流间的命令不能并发执行，应用需遵循指南以提升并发内核执行潜力。
重叠行为（Overlapping Behavior）：两个流间的执行重叠程度取决于命令发出的顺序及设备对数据传输与内核执行重叠、并发内核执行和并发数据传输的支持情况。
主机函数（Host Functions）：通过cudaLaunchHostFunc()可在流中插入CPU函数调用，该函数在流中所有之前命令完成后执行，且后续命令需等待该函数完成才开始。
流优先级（Stream Priorities）：通过cudaStreamCreateWithPriority()创建时可指定流的优先级，高优先级流的工作将优先于低优先级流执行。
程序化依赖启动与同步（Programmatic Dependent Launch and Synchronization）：针对计算能力9.0及以上设备，允许依赖于主内核结果的次内核在主内核完成前启动，利用内核执行中的非依赖部分实现并发，以提高性能，并引入新的API支持这一机制，减少启动延迟。
程序化依赖启动（Programmatic Dependent Launch）：在CUDA编程中，允许一个主核函数（primary kernel）和一个次核函数（secondary kernel）在同一CUDA流中启动。主核函数需在所有线程块执行完毕后，通过调用cudaTriggerProgrammaticLaunchCompletion来表明准备就绪，以便次核函数启动。次核函数必须使用可扩展启动API，并可设置cudaLaunchAttributeProgrammaticStreamSerialization属性，使得CUDA驱动能够在不等待主核函数完成及其内存刷新的情况下提前启动次核函数。若主核函数未显式触发，该触发会在所有主核函数线程块退出后隐式发生。在这种模式下，为确保数据一致性，次核函数必须使用cudaGridDependencySynchronize或其他机制来同步来自主核函数的结果数据。
CUDA图中的应用：程序化依赖启动也可应用于CUDA图中，通过流捕获或直接利用边数据实现。在CUDA图的两个核函数节点之间使用cudaGraphDependencyTypeProgrammatic类型的边，可以使得上游核函数对下游核函数中的cudaGridDependencySynchronize()可见。此类型要求使用特定的输出端口，如cudaGraphKernelNodePortLaunchCompletion或cudaGraphKernelNodePortProgrammatic。
CUDA图概述：CUDA图作为一种新的工作提交模型，允许将操作（如核函数启动）及其间的依赖关系定义为独立于执行流程的序列，从而实现定义一次、多次执行的效率提升。图的定义与执行分离，能够减少CPU启动开销，并为CUDA提供整个工作流程视图以进行潜在的优化，尤其对于短时执行的GPU核函数，能够显著降低总体执行时间的开销部分。
图结构与创建：CUDA图由操作节点和它们之间的依赖边组成，这些依赖关系限制了操作的执行顺序。图的创建可以通过显式API或流捕获机制实现，后者能够将现有基于流的API代码段转换为图。流捕获通过cudaStreamBeginCapture和cudaStreamEndCapture函数实现，期间向流中添加的工作被记录为图的一部分而非立即执行。
边数据（Edge Data）：自CUDA 12.3起，图中的边可以携带数据，用于修改依赖关系的行为，包括指定输出端口（触发时机）、输入端口（依赖哪部分节点）以及依赖类型。这为图中的依赖关系提供了更多控制，例如通过cudaGraphDependencyTypeProgrammatic类型支持核函数间的程序化依赖启动。边数据可以在创建和查询图的API中使用，同时也适用于某些流捕获API。
跨流依赖处理：流捕获能够管理通过cudaEventRecord()和cudaStreamWaitEvent()表达的跨流依赖，前提是等待的事件被记录在同一捕获图中。当事件在捕获模式的流中被记录时，它会生成一个捕获事件，该事件代表捕获图中的多个节点集。如果流等待一个被捕获的事件，则会将该流置于捕获模式，并对后续操作添加对该捕获事件节点的额外依赖。所有相关的流最终都会被合并回最初调用cudaStreamBeginCapture()的源头流。
禁止和未处理的操作：在流捕获期间，同步或查询处于捕获状态的流或捕获事件的执行状态是无效的，因为它们不代表已调度执行的项目。此外，在任何关联流正在进行捕获的情况下，使用遗留流（默认流）也是非法的。尝试合并两个不同的捕获图，或者在没有指定cudaEventWaitExternal标志的情况下从一个被捕获的流等待非捕获事件也是不允许的。某些异步API在流捕获模式下不支持，如cudaStreamAttachMemAsync()。
失效处理：若在流捕获过程中尝试执行非法操作，相关的捕获图会被失效。一旦捕获图失效，继续使用相关联的正在捕获的流或捕获事件会返回错误，直到通过cudaStreamEndCapture()结束捕获并使流脱离捕获模式。
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()），以及查询节点状态。
设备图启动允许在设备端发起图执行，适用于需要根据运行时数据依赖做出决策的工作流程，支持统一寻址的系统可以使用此功能。设备图可从主机和设备启动，而主机图只能从主机启动；设备图在设备上不能同时启动两次，同时从主机和设备启动行为未定义。
设备图创建：要从设备上启动图，需通过在cudaGraphInstantiate()调用中传递cudaGraphInstantiateFlagDeviceLaunch标志来显式实例化。设备图的结构在实例化时固定，更新需要重新实例化，且只能在主机上执行。
图上传：设备图执行前需上传至设备以准备必要资源。这可通过cudaGraphUpload()或在cudaGraphInstantiateWithParams()中请求上传实现，也可通过首先从主机启动图隐式完成上传。
设备图更新与重上传：设备图仅能在主机上更新，并在更新可执行图后需重新上传到设备以应用变化。
设备端启动：设备图支持从主机和设备启动，使用相同的cudaGraphLaunch()签名。设备上启动时必须来自另一个图，且为线程级操作，允许多线程并行启动。
流管理：设备图不能在常规CUDA流中启动，只能在表示特定启动模式的命名流中启动，包括“fire and forget”（立即提交执行）模式。
执行环境与同步：设备图启动会产生独立的执行环境，封装所有工作及生成的子工作。理解设备端同步模型需了解执行环境概念。主机启动图时，存在一个流环境作为父级。
尾部启动（Tail Launch）：作为替代传统同步方法（如cudaDeviceSynchronize()）的机制，用于实现序列工作依赖。当图及其所有子图完成时，尾部启动的图将按顺序执行。
条件节点：允许图中的条件执行和循环，支持动态迭代工作流程的图形化表示，提升CPU并行性。条件节点包含一个条件处理程序，用于评估是否执行其包含的图。
这些要点概述了利用GPU设备图进行优化的关键方面，包括图的创建、资源管理、执行控制以及高级功能如条件执行的支持。
条件IF节点（Conditional IF Nodes）：在图执行过程中，如果条件非零，则IF节点的主体图将执行一次。条件默认值通过上游内核设定，条件主体则利用图API进行填充。
条件WHILE节点（Conditional WHILE Nodes）：WHILE节点的主体图将一直执行直至条件变为零。条件在节点执行时及主体图完成后被评估。同样，条件主体的构建也依赖于图API。
事件（Events）：运行时允许应用异步记录程序中的事件以密切监控设备进度和进行精确计时。事件完成标志着其前所有任务或指定流中的所有命令已完成。流零中的事件在所有流的所有前置任务和命令完成后完成。
事件的创建与销毁：示例代码展示了如何创建和销毁两个事件。
持续时间测量：创建的事件可用来测量代码段的执行时间。
同步调用（Synchronous Calls）：当调用同步函数时，主机线程会等待设备完成请求任务后才恢复控制。主机线程的行为（如让出、阻塞或自旋）可通过cudaSetDeviceFlags()预先设定。
多设备系统：
- 设备枚举：主机系统可能有多台设备，示例代码演示了如何枚举这些设备及其属性查询。
- 设备选择：通过cudaSetDevice()，主机线程可随时更改当前操作的设备。
- 流和事件行为：说明了不同设备间的流、事件交互规则，包括失败和成功的情况。
- 对等内存访问（Peer-to-Peer Memory Access）：特定系统配置下，设备间可以直接访问对方内存，需要通过cudaDeviceCanAccessPeer()和cudaDeviceEnablePeerAccess()启用。
- 统一虚拟地址空间（Unified Virtual Address Space）：在64位进程中，主机和计算能力2.0及以上设备共享单一虚拟地址空间，简化了内存管理和访问。
  以上摘要覆盖了GPU编程和优化中关于条件控制流、事件管理、同步操作、多设备资源管理和内存访问的关键概念。
进程间通信（Interprocess Communication, IPC）：
- GPU内存指针和事件句柄在同一流程内的线程间可直接引用，但不支持跨进程直接引用。
- 要实现跨进程共享，需使用IPC API，该API仅支持Linux系统的64位进程及计算能力2.0以上的设备。注意，cudaMallocManaged分配的内存不支持IPC API。
- 应用程序可通过cudaIpcGetMemHandle()获取内存指针的IPC句柄，通过标准IPC机制传递，并用cudaIpcOpenMemHandle()在其他进程中恢复有效指针。
- 为防止信息泄露，推荐只共享大小为2MiB对齐的内存块。
错误检查：
- 所有运行时函数会返回错误码，但异步函数的错误码无法立即反映设备上的异步错误，需通过cudaDeviceSynchronize()同步后检查。
- 运行时为每个主机线程维护一个错误变量，可通过cudaPeekAtLastError()查看而不清空，或用cudaGetLastError()查看并重置为成功状态。
- 核心启动（kernel launch）不返回错误码，需在启动后立即调用错误检查函数以捕获预启动错误，并确保在检查前调用cudaGetLastError()重置错误变量。
调用栈：
- 计算能力2.x及以上设备可查询和设置调用栈大小。
- 调用栈溢出会导致内核调用失败，调试模式下显示堆栈溢出错误，否则为未指定的启动错误。
- 编译器无法静态确定栈大小时会发出警告，此时需要手动设置栈大小。
纹理与表面内存：
- 纹理与表面内存访问利用GPU图形文本硬件子集，相比全局内存访问，可带来性能优势。
- 纹理对象API允许创建和管理纹理对象，控制访问模式、格式等。
- 支持16位浮点纹理处理，需要通过特定函数进行类型转换。
- 分层纹理提供一种组织纹理数据的方式，适用于一维或二维纹理数组，且支持在单一层内进行纹理过滤。
  以上内容覆盖了GPU优化中的关键概念，包括跨进程资源共享、错误管理和调试、程序调用栈管理以及如何利用纹理内存来提升数据读取性能。
立方体贴图纹理（Cubemap Textures）：这是一种特殊的二维分层纹理，包含六层，分别代表立方体的六个面。它只能通过带有cudaArrayCubemap标志的cudaMalloc3DArray()函数创建。立方体贴图纹理通过texCubemap()函数访问，并且要求设备计算能力为2.0或更高。
立方体贴图分层纹理（Cubemap Layered Textures）：这类纹理的每一层都是相同维度的立方体贴图。访问时使用一个整数索引和三个浮点纹理坐标，索引指定位子立方体贴图中的序列，坐标则定位该立方体贴图内的像素。它们需要通过带有cudaArrayLayered和cudaArrayCubemap标志的cudaMalloc3DArray()创建，使用texCubemapLayered()函数获取数据，并且只支持计算能力2.0及以上的设备。
纹理聚集（Texture Gather）：这是一种针对二维纹理的特殊纹理获取方式，通过tex2Dgather()函数实现，可以同时返回四像素的特定分量值。此功能要求纹理数组在创建时带有cudaArrayTextureGather标志，且尺寸小于纹理聚集的最大限制，仅支持计算能力2.0以上的设备。
表面内存（Surface Memory）：对于计算能力2.0及以上的设备，带有cudaArraySurfaceLoadStore标志的CUDA数组可以通过表面对象进行读写操作。表面内存使用字节寻址，与纹理内存不同。
CUDA数组（CUDA Arrays）：专为纹理获取优化的不透明内存布局，支持一维、二维、三维，元素由1到4个组件组成，适用于各种数据类型。CUDA数组通过纹理内存或表面内存接口访问。
读/写一致性（Read/Write Coherency）：纹理和表面内存是缓存的，但在同一个内核调用中，缓存不会与全局内存或表面内存的写操作保持一致，因此在同一内核调用中对刚被写入的位置进行纹理获取或表面读取可能会得到未定义数据。
图形互操作性（Graphics Interoperability）：允许CUDA与OpenGL、Direct3D资源之间进行映射，以实现数据共享。资源需先注册到CUDA，通过特定函数映射和取消映射，映射时可指定读写标志以优化资源管理。映射后的资源可在内核中通过返回的设备内存地址访问。资源在映射状态下通过其他API或CUDA上下文访问会导致未定义行为。
以上内容涉及GPU优化的关键技术点，包括高级纹理特性、内存访问模式以及GPU与其他图形API的交互方式。

OpenGL与CUDA的互操作性：
- 可映射到CUDA地址空间的OpenGL资源包括缓冲区、纹理和渲染缓冲对象。
- 缓冲区对象通过cudaGraphicsGLRegisterBuffer()注册，然后在CUDA中作为设备指针使用，支持内核读写及cudaMemcpy()调用。
- 纹理或渲染缓冲对象通过cudaGraphicsGLRegisterImage()注册，在CUDA中表现为CUDA数组，可通过绑定为纹理或表面引用来读取，使用cudaGraphicsRegisterFlagsSurfaceLoadStore标志可写入。同时支持cudaMemcpy2D()访问。
- 进行OpenGL互操作API调用时，共享资源的OpenGL上下文需对主机线程当前有效。
- 创建无绑定OpenGL纹理后（如使用glGetTextureHandle*），无法再与CUDA注册。应用需先注册纹理进行互操作，再请求图像或纹理句柄。
- 对于Quadro GPU，在多GPU配置下，通过cudaWGLGetDevice()可提高OpenGL与CUDA间的互操作性能。
Direct3D与CUDA的互操作性：
- 支持Direct3D 9Ex、Direct3D 10、Direct3D 11。
- Direct3D设备需满足特定创建条件以实现互操作，如Direct3D 9Ex需设置D3DDEVTYPE_HAL和D3DCREATE_HARDWARE_VERTEXPROCESSING。
- Direct3D的缓冲区、纹理、表面可通过相应的cudaGraphicsD3D*RegisterResource()函数注册到CUDA地址空间。
SLI互操作性：
- 在多GPU系统中的SLI模式下，每个CUDA启用的GPU被视为单独设备，但存在内存消耗和数据传输的特殊考量。
- 应用应为每块GPU创建单独的CUDA上下文，以避免不必要的数据转移，并利用cudaD3D*GetDevices()和cudaGLGetDevices()来识别渲染使用的设备对应的CUDA句柄。
外部资源互操作性：
- 允许CUDA导入由其他API显式导出的资源，如通过操作系统句柄或NVIDIA软件通信接口。
- 可导入内存对象和同步对象，使用cudaImportExternalMemory()和cudaImportExternalSemaphore()，分别通过cudaDestroyExternalMemory()和cudaDestroyExternalSemaphore()释放。
- 对于Vulkan互操作，导入的内存和同步对象必须在相同的设备上创建和映射，通过比较设备UUID确保匹配。
  以上摘要涵盖了与GPU优化相关的API使用、互操作性配置以及在多GPU环境下的注意事项，对开发者理解和实施跨API资源管理和加速计算有重要指导意义。
  ####在Linux和Windows 10上，CUDA支持导入Vulkan导出的专用和非专用内存对象；而Windows 7仅支持导入专用内存对象。导入Vulkan专用内存对象时，需设置cudaExternalMemoryDedicated标志。
  Vulkan内存对象通过不同类型的句柄导出后，可被CUDA以相应方式导入：
使用VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT导出的，可通过文件描述符导入CUDA。
使用VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT或VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT导出的，分别通过NT句柄或全局共享的D3DKMT句柄导入CUDA。注意，NT句柄的生命周期管理由应用程序负责。
导入的内存对象可映射为设备指针或CUDA mipmapped数组，映射时需匹配创建时的偏移量、尺寸等参数，并且映射资源最终需使用cudaFree或cudaFreeMipmappedArray释放。
Vulkan信号量对象也可通过类似方式导入CUDA，并可以在CUDA中进行信号发送与等待操作，但对应的等待或信号发送操作必须在Vulkan中对应执行。
此外，文中提到OpenGL与Direct3D 12也可与Vulkan/CUDA进行互操作，特别是在内存和同步对象的导入与共享方面，通过特定扩展实现跨API的资源协调访问。例如，在Direct3D 12与CUDA的互操作中，需要确保在相同设备上导入和映射内存对象，并通过LUID匹配来确定相应的CUDA设备。
综上所述，该内容重点介绍了如何在CUDA中导入Vulkan、OpenGL及Direct3D 12的内存对象和同步对象，以及这些对象在CUDA中的使用方法，包括映射、信号控制等关键操作，强调了不同平台下的句柄类型处理和资源生命周期管理的重要性。
####在GPU优化的背景下，文章内容涉及了CUDA与Direct3D 12及Direct3D 11之间的互操作性，特别是在内存对象映射、多级纹理数组（mipmapped arrays）处理以及同步对象导入方面的高级知识点。以下为关键信息摘要：

映射多级纹理数组至导入内存对象：CUDA支持将多级纹理数组映射到通过Direct3D API创建的导入内存对象上，要求偏移量、尺寸、格式和多级数匹配。若该数组能在Direct3D 12中作为渲染目标使用，需设置cudaArrayColorAttachment标志。所有映射的多级纹理数组需通过cudaFreeMipmappedArray()释放。
导入同步对象：Direct3D 12的共享栅栏对象可通过其NT句柄导入CUDA中。应用需负责不再需要时关闭句柄，并在释放底层信号量前显式释放句柄。栅栏对象可在CUDA中进行信号设定与等待操作，但相关等待操作必须在Direct3D中发起，并且在信号设定之后。
Direct3D 11互操作性：
- 匹配设备LUID：导入由Direct3D 11导出的内存和同步对象时，需确保在相同设备上进行。通过比较CUDA设备与Direct3D 11设备的LUID来确定对应关系。
- 导入内存对象：通过特定标志创建的可共享Direct3D 11资源（如纹理和缓冲区），可使用NT句柄或全局共享D3DKMT句柄导入CUDA。同样，导入时需设置cudaExternalMemoryDedicated标志，并负责管理句柄生命周期。
- 映射缓冲区和多级纹理数组：类似地，设备指针可映射到导入的内存对象上，映射条件需符合Direct3D 11 API的创建参数。映射的多级纹理数组处理方式与Direct3D 12中相似。
- 导入Direct3D 11同步对象：包括共享栅栏对象和键控互斥锁（keyed mutex），均可通过NT句柄或命名句柄乃至全局共享D3DKMT句柄导入CUDA，同样强调句柄管理的重要性。
  这些知识点对于进行高性能计算和图形应用开发，特别是在跨API协同工作以优化GPU资源利用的场景下至关重要。

Direct3D 11同步对象导入与操作：文章讨论了如何在Direct3D 11环境中导入并操作同步对象（如fence和keyed mutex），包括信号发送（signaling）和等待（waiting）。对于fence对象，可以设置其值，并要求相应的等待操作必须在Direct3D 11中发出且在信号发送之后。Keyed mutex对象通过指定键值释放或等待，等待操作需匹配相同的键值，并且必须在信号发送后执行。
NVIDIA软件通信接口互操作性（NVSCI）：介绍了NvSciBuf和NvSciSync接口，设计用于特定目的，详细信息可参考NVIDIA官方文档。这些接口支持CUDA设备与NVIDIA驱动的其他组件间内存对象和同步对象的导入与映射，以实现跨系统资源共享和一致性维护。
NvSciBuf内存对象的分配与导入：说明了如何根据CUDA设备ID配置NvSciBuf属性以分配兼容内存，以及如何将分配的NvSciBuf对象导入CUDA环境。强调了在多驱动程序共享内存时使用NvSciSync对象作为屏障以保持缓存一致性的重要性。
映射缓冲区与映射mipmapped数组：阐述了如何将设备指针或CUDA mipmapped数组映射到导入的NvSciBuf内存对象上，并指出所有映射必须通过cudaFree或cudaFreeMipmappedArray释放。
NvSciSync同步对象的导入与操作：展示了如何生成与CUDA设备兼容的NvSciSync属性，创建NvSciSync对象，并将其导入CUDA。详细说明了如何对这些导入的同步对象进行信号发送和等待操作，同时介绍了跳过默认内存同步操作的标志（如cudaExternalSemaphoreSignalSkipNvSciBufMemSync），这在某些情况下可以提高性能。
版本控制与兼容性：解释了开发者在开发CUDA应用时需要考虑的两个版本号：计算能力和CUDA驱动API版本。强调了驱动API的向后兼容性，以及不同版本间的混合使用限制，特别是Tesla GPU产品引入的向前兼容升级路径。
综上所述，摘要内容涉及了GPU编程中的高级同步技术、资源管理和版本兼容性问题，这些都是进行高效GPU优化的关键方面。

计算模式（Compute Modes）：在使用Windows Server 2008及以后或Linux的Tesla解决方案上，通过NVIDIA System Management Interface (nvidia-smi) 可以将系统中的任何设备设置为三种模式之一。这影响了未明确调用cudaSetDevice()的主机线程可能关联的设备。对于Pascal架构及之后的设备，支持指令级粒度的计算抢占（Compute Preemption），以替代之前Maxwell和Kepler架构的线程块抢占，有助于防止长运行内核的应用程序独占系统或超时。但这也伴随着上下文切换开销，可以通过查询设备属性cudaDevAttrComputePreemptionSupported来确定设备是否支持该特性。
模式切换（Mode Switches）：具有显示输出的GPU会分配一部分DRAM给主表面（primary surface），用于刷新显示器。当用户改变显示分辨率或深度时，需要重新分配主表面内存，可能导致CUDA应用的内存分配被回收，进而使得CUDA运行时调用失败并返回无效上下文错误。
Tesla计算集群模式（Tesla Compute Cluster Mode）：针对Tesla和Quadro系列设备，Windows设备驱动可通过nvidia-smi设置为TCC模式，此模式移除了所有图形功能支持，专为计算任务设计。
硬件实现：NVIDIA GPU架构基于多线程流式多处理器（SMs）的可扩展阵列构建。采用单指令多线程（SIMT）架构并发执行大量线程，利用指令级并行和硬件多线程实现广泛的线程级并行。与CPU核心不同，SIMT架构中的指令按顺序发出，没有分支预测或推测执行。同时，GPU架构采用小端字节序表示法。
####SIMT架构概述了GPU执行线程的方式，其中线程以32个并行线程的组，即“warp”进行管理、调度和执行。一个warp中的所有线程从相同的程序地址开始，但拥有各自的指令地址计数器和寄存器状态，允许它们独立分支和执行。在Volta架构之前，warp内的线程因为数据依赖条件分支而发生分歧时，会降低效率，因为warp作为一个整体只能同时执行一条指令。而从Volta架构开始，引入了独立线程调度，使得线程间的并发性不再受warp限制，提高了灵活性，线程可以在子warp粒度上分散和重聚。
硬件多线程方面，每个被多处理器处理的warp的执行上下文（如程序计数器、寄存器等）在其整个生命周期内都保持在片上，因此上下文切换无成本。多处理器具有分配给各个warp使用的32位寄存器集以及分配给线程块的并行数据缓存或共享内存。内核能够同时在多处理器上驻留和处理的块和warp数量取决于内核使用的寄存器和共享内存量以及多处理器的可用资源。
性能优化策略主要围绕四个方面：最大化利用率、减少内存访问、管理内存使用和调整线程格大小。为了实现最佳性能提升，需要根据应用的具体瓶颈来选择合适的优化策略，并通过持续测量和监控（如使用CUDA Profiler工具）来指导优化工作，同时对比实际的浮点运算吞吐量或内存吞吐量与设备理论峰值，以识别改进空间。最大化利用率特别强调应用结构应能展现出尽可能多的并行性，并高效映射到系统各组件上，以保持其大部分时间忙碌。
####在应用层面，为了最大化主机、设备以及它们之间总线的并行执行，应使用异步函数调用和流，确保每个处理器执行其最擅长的工作类型：将串行工作负载分配给主机，而并行工作负载则分配给设备。算法中因线程间同步和数据共享而中断并行性的地方，应尽可能地在单个线程块内完成计算以减少额外的核调用和全局内存通信开销。
在设备层面，通过多流技术使多个核函数并发执行，以实现设备中多处理器的最大并行利用。
在更细的多处理器层面，应用程序应最大化单个多处理器内部各功能单元间的并行执行。关键在于通过线程级并行来充分利用功能单元，维持足够的活跃线程束（warps）以隐藏指令延迟，同时考虑寄存器依赖、内存访问延迟及同步点对并行度的影响。减少寄存器使用、控制内存访问模式和合理配置执行配置（如确保每个块的线程数为 warp 大小的倍数）对性能至关重要。
总之，有效的GPU优化策略包括跨层次的并行化措施，从宏观的异步执行和任务分配，到微观的线程调度与资源管理，旨在减少瓶颈、隐藏延迟并最大化所有计算资源的利用率。

5.2.3.1. 占用率计算器(Occupancy Calculator)

存在多个API函数帮助程序员根据寄存器和共享内存需求选择线程块大小和集群大小。
示例代码展示了如何计算MyKernel的占用率，通过并发线程束与多处理器最大线程束的比例报告占用水平。
另一示例展示了如何基于用户输入使用占用率来配置MyKernel的启动。
集群占用API的使用示例展示了如何找到给定大小的最大活动集群数，例如计算大小为2和每个块128个线程的集群。
计算能力9.0起支持8大小的集群，但硬件或MIG配置过小无法支持8个多处理器的GPU将减少最大集群大小。建议用户在启动集群内核前查询最大集群大小，可通过cudaOccupancyMaxPotentialClusterSize API查询。
CUDA Nsight Compute提供了一个独立的占用率计算器和启动配置器实现，位于<CUDA_Toolkit_Path>/include/cuda_occupancy.h，适用于不能依赖CUDA软件栈的场景。

5.3. 最大化内存吞吐量(Maximize Memory Throughput)

提升应用整体内存吞吐量的第一步是减少带宽较低的数据传输。
应最小化主机与设备间的数据传输，因这些传输的带宽远低于全局内存与设备间的传输。
也应通过最大化使用片上内存（共享内存和缓存，如L1、L2缓存以及所有设备上的纹理缓存和常量缓存）来减少全局内存与设备间的数据传输。
共享内存相当于用户管理的缓存，应用程序需显式分配和访问。
对于某些数据访问模式依赖于数据的应用，传统硬件管理的缓存更适合利用数据局部性。
计算能力7.x、8.x和9.0的设备允许在L1和共享内存之间配置内存比例。
内核的内存访问吞吐量受访问模式影响可能相差一个数量级，因此需要根据设备内存访问模式进行优化，尤其是全局内存访问。

5.3.1. 主机与设备间的数据传输

应用应尽量减少主机与设备间的数据传输，可以通过将更多代码从主机移至设备实现，即使这意味着运行的内核在设备上无法完全高效执行并行任务。
批量小传输为单一大传输总是表现更佳，因为每次传输都有开销。
使用页面锁定主机内存可提升具有前端总线系统的主机与设备间数据传输性能。
映射的页面锁定内存(Mapped Memory)无需分配设备内存或显式复制数据，访问映射内存时隐式执行数据传输，要求访问合并以达到最佳性能，适合集成系统中物理相同的主机和设备内存交互。

全局内存访问：全局内存访问通过32、64或128字节的内存事务进行，要求自然对齐。访问模式越分散，吞吐量越低。为了最大化吞吐量，应通过合并访问来减少交易数量，确保数据类型大小和对齐满足要求。
局部内存：局部内存访问具有与全局内存相似的延迟和带宽特性，但连续线程ID访问连续32位字会完全合并。计算能力5.x及以后的设备在L2缓存中缓存局部内存访问。
共享内存：共享内存因位于芯片上而具有高带宽和低延迟。它被划分为多个可同时访问的bank。避免bank冲突以实现最大性能是关键，因为冲突会导致访问序列化。
常量内存：常量内存存储于设备内存并被缓存在常量缓存中。访问时，请求会被分割以处理不同地址，影响吞吐量。
纹理和表面内存：这些内存类型同样位于设备内存中，并利用纹理缓存优化了2D空间局部性访问，适用于流式读取且具有固定延迟。通过纹理或表面读取可以减少DRAM带宽需求，适合特定的数据访问模式。
理解并优化这些内存访问模式对于提升GPU程序性能至关重要，特别是在处理大规模并行计算任务时。
####为了最大化指令吞吐量，应用程序应确保：

通过计算每周期每多处理器的操作数来衡量吞吐量，对于32线程束（warp），一个指令相当于32次操作。
多处理器的吞吐量需乘以设备中的多处理器数量以获得整个设备的吞吐量。

5.4.1. 算术指令优化

使用硬件原生支持的算术指令，并注意不同计算能力的设备可能有不同的实现。
编译时使用-ftz=true（将非规范化数字截断为零）可提高性能。
减少精度的除法和平方根运算（使用-prec-div=false和-prec-sqrt=false）可以提高性能。
使用__fdividef()函数进行单精度浮点除法比使用除号更快。
直接调用rsqrtf()而非通过1.0/sqrtf()计算单精度浮点倒数平方根以保持性能。
单精度平方根实现考虑了0和无穷大的特殊情况。
三角函数如sinf(x), cosf(x)等在大输入值时性能降低，因涉及复杂的近似计算及可能的局部内存使用，影响了吞吐量。
整数除法和取模操作成本较高，可被位操作替代（当除数为2的幂时）。
半精度浮点数运算推荐使用向量数据类型half2或__nv_bfloat162及相应的向量内建函数以提升效率。
类型转换应避免不必要的开销，如使用单精度浮点常量。

5.4.2. 控制流指令优化

避免条件分支导致线程束分歧，尽量使控制条件与线程束对齐，减少分歧的线程束数量。
编译器可能会通过条件预测（branch predication）优化if或switch结构，避免分歧。
使用#pragma unroll控制循环展开，进一步优化性能。

5.4.3. 同步指令优化

__syncthreads()指令的吞吐量依据设备计算能力而异，例如在计算能力6.0的设备上为每周期32次操作。
注意，__syncthreads()可能导致多处理器空闲等待，影响整体性能。

最小化内存颠簸：频繁地分配和释放内存的应用可能会遇到分配调用随时间逐渐变慢的情况。为优化性能，建议减少内存操作的频率，避免内存颠簸。
CUDA启用的GPU：NVIDIA官网列出了所有CUDA支持的设备及其计算能力。开发时可查询GPU的计算能力、多处理器数量、时钟频率、总设备内存等属性，以优化程序针对特定硬件的执行效率。
C++语言扩展：
- 函数执行空间指定符：如__global__表示定义的函数为内核函数，运行在设备上；__device__表示函数在设备上执行；__host__表示函数在主机上执行。
- 变量内存空间指定符：如__device__表示变量存储在设备内存中；__constant__用于声明常量内存中的变量，访问速度快但大小有限；__shared__用于声明共享内存中的变量，适合线程块内的数据共享。
- 内存管理与优化：使用__managed__指定符可以声明统一内存中的变量，自动处理主机与设备内存之间的数据迁移，但需注意其可能带来的额外开销。
- 限制指针(__restrict__)：通过限制指针的别名使用，帮助编译器进行更多的优化，如重排序和公共子表达式消除，但可能增加寄存器压力，影响CUDA代码的线程占用率（即占用）。
  这些知识点涵盖了内存管理、硬件选择、以及利用C++扩展来优化GPU代码的关键方面，对于提高CUDA程序的执行效率至关重要。
向量数据类型：包括char、short、int、long、longlong、float和double的向量类型，这些类型是基于基本整数和浮点数类型的结构体，可以通过x、y、z、w字段访问其1至4个分量。向量类型具有特定的对齐要求，且提供了构造函数来生成特定类型的向量。
dim3类型：用于指定维度的基于uint3的整数向量类型，在定义dim3类型的变量时，未明确指定的任何分量将初始化为1。
内置变量：用于指定网格和块的尺寸以及块和线程的索引，仅在设备上执行的函数中有效。包括gridDim（网格尺寸）、blockIdx（块索引）、blockDim（块尺寸）、threadIdx（线程索引）和warpSize（线程束大小）。
内存栅栏函数：用于在CUDA弱排序内存模型下强制执行内存访问的顺序一致性。不同函数作用于不同范围内的顺序保证，如__threadfence_block()、__threadfence()和__threadfence_system()分别在块、设备和系统级别上确保写操作的可见性。
同步函数：如__syncthreads()用于确保块内所有线程到达同一执行点，并使之前的全局和共享内存访问对块内所有线程可见。在计算能力2.x及更高版本的设备上，还支持带有条件判断的变体，如__syncthreads_count()、__syncthreads_and()和__syncthreads_or()。
内存栅栏与同步函数区别：内存栅栏仅影响线程内部的内存操作顺序，而同步函数除了排序外，还确保了内存操作结果在线程间的可见性，例如通过使用volatile变量。
数学函数与纹理函数：支持一系列C/C++标准库数学函数和特定于设备的内建数学函数，提供部分函数的精度信息。纹理对象和纹理获取API用于高级内存访问模式，特别是针对图像和数据密集型应用。
综上所述，GPU优化涉及理解并有效利用向量数据类型以提高内存访问效率、正确应用内置变量进行并行任务划分、合理安排内存栅栏和同步点以确保数据一致性与同步，以及选择合适的数学与纹理函数来优化计算性能。
####这些条目概述了CUDA编程中用于访问纹理内存的多种函数，重点在于不同维度、层状纹理、立方体贴图以及稀疏CUDA数组的处理。关键知识点包括：
tex1Dfetch(): 从一维纹理对象指定的线性内存区域获取数据，仅支持非归一化坐标和边界/钳制寻址模式。
tex1D(), tex2D(), tex3D(): 分别从一维、二维、三维CUDA数组或线性内存区域根据纹理坐标获取数据。
lod版本函数: 如tex2DLod()，在指定的细节级别(level-of-detail)上执行获取操作。
Grad版本函数: 如tex2DGrad()，根据纹理坐标及其梯度(dx, dy)计算细节级别进行获取。
Layered纹理功能: tex2DLayered()等函数允许从带有层次的二维纹理中获取数据，引入了“层”(layer)索引。
Sparse CUDA数组支持: 针对稀疏CUDA数组的函数变体，如tex2D() for sparse CUDA arrays，会检查texel是否驻留在内存中，并在未驻留时返回零值。
Cubemap和Cubemap Layered纹理: 特定于立方体贴图纹理的函数，如texCubemap()和texCubemapLayered()，支持基于(x,y,z)纹理坐标的立方体贴图纹理数据获取。
对于GPU优化，理解这些函数的适用场景和性能影响至关重要，例如选择正确的寻址模式、利用lod和梯度来高效地管理纹理采样，以及针对稀疏数据结构优化内存访问，从而减少带宽消耗并提升计算效率。
**表面函数（Surface Functions）**仅支持计算能力2.0及以上的设备，用于读写CUDA数组，并提供了多种维度和层面上的操作，如surf1Dread()、surf2Dwrite()等。这些操作支持不同的边界模式处理超出范围的坐标，包括钳制（clamping）、返回零、或触发错误。
**只读数据缓存加载函数（Read-Only Data Cache Load Function）**从地址加载数据到缓存中，支持计算能力5.0及以上设备，适用于提高全局内存访问效率，支持多种数据类型。
**使用缓存提示的加载函数（Load Functions Using Cache Hints）和存储函数（Store Functions Using Cache Hints）**同样要求计算能力5.0及以上，允许开发者利用缓存操作符来优化内存访问模式，提升数据加载和存储性能。
**时间函数（Time Function）**提供了一种测量内核在设备上完全执行所需时钟周期的方法，尽管它不直接反映指令执行时间，但可用于性能分析。
**原子函数（Atomic Functions）**实现了对全局或共享内存中的单个32位、64位或128位字的读改写原子操作，保证了多线程环境下的数据一致性。这些操作对于实现线程安全的并发更新至关重要，且只能在设备函数中使用。同时，展示了如何基于atomicCAS()实现不直接提供的原子操作，如在低版本计算能力设备上的双精度浮点数加法。
以上知识点涉及了GPU编程中的内存访问优化、性能监控以及并发控制的关键技术。
原子操作函数：
- atomicAdd() 支持16位、32位、64位整数及浮点数的原子加法，其中32位浮点数支持从计算能力2.x开始，64位浮点数从6.x开始，__half2和__nv_bfloat162从6.x开始，且保证每个元素的原子性而非整体。float2和float4仅在9.x及以上版本支持，且只针对全局内存地址。
- atomicSub() 实现原子减法。
- atomicExch() 实现原子交换，包括对128位数据的支持，但需满足特定条件且仅在9.x及以上版本支持。
- atomicMin() 和 atomicMax() 分别实现原子最小值和最大值更新，64位版本从5.0计算能力开始支持。
- atomicInc() 和 atomicDec() 实现原子递增和递减。
- atomicCAS() 实现比较并交换操作，也支持128位数据从9.x版本起。
按位操作函数：
- atomicAnd()、atomicOr()、atomicXor() 分别实现原子的按位与、或、异或操作，64位版本从5.0计算能力开始支持。
地址空间判定函数：
- 如 __isGlobal()、__isShared() 等用于判断指针所指向的内存区域类型。
地址空间转换函数：
- 如 __cvta_generic_to_global() 等用于在不同地址空间之间转换指针。
  这些知识点对于进行GPU编程时，尤其是在处理并发访问共享数据、控制内存访问模式以及优化内存使用效率方面至关重要。

内存访问转换函数: __cvta_constant_to_generic() 和 __cvta_local_to_generic() 分别通过执行PTX指令将常量和局部内存地址转换为通用指针，这对于优化内存访问模式和提高数据传输效率至关重要。
Alloca函数: 在计算能力5.2及以上版本支持，alloca() 函数在调用者的堆栈帧中动态分配内存，自动在设备代码调用时对齐16字节，并在调用者返回时自动释放。这有助于栈上快速分配临时或动态大小的内存，减少内存管理开销。
编译器优化提示函数:
- __builtin_assume_aligned(): 允许编译器假设指针参数至少按align字节对齐，有助于生成更高效的加载/存储指令。
- __builtin_assume(), __assume(): 让编译器根据提供的布尔表达式进行假设，以指导优化决策，但运行时不满足假设会导致未定义行为。
- __builtin_expect(): 提供给编译器分支预测信息，指导生成更优的代码路径。
- __builtin_unreachable(): 标记某段代码永远不会被执行，帮助编译器进行更激进的优化。
Warp投票函数: 如 __all_sync, __any_sync, __ballot_sync 等，用于线程束内执行布尔值的并行比较和广播操作，适用于计算能力7.x及以上的设备，需要同步掩码来指定参与线程，对于并行决策和分支控制有重要作用。
Warp匹配函数: 例如 __match_any_sync, __match_all_sync，在计算能力7.x及以上设备上提供线程束内变量值的广播和比较，用于数据一致性检查和同步。
Warp归约函数: 如 __reduce_sync, 支持计算能力8.x及以上设备，用于线程束内数据的并行归约操作（如加法、最小值、最大值等），要求同步掩码确保正确同步。
Warp洗牌函数: 包括 __shfl_sync, __shfl_up_sync, __shfl_down_sync, __shfl_xor_sync，在计算能力5.0及以上设备上实现线程束内数据交换，对于并行算法中的数据共享和分布处理非常重要，已废弃非同步版本，推荐使用同步版本。
这些知识点集中于内存管理、编译优化、线程间通信与同步等方面，是GPU编程和优化的关键技术点。

__shfl_sync() 内置函数：允许在同一warp内的线程间交换变量，无需使用共享内存。交换对于warp内所有活动线程同时进行，根据类型移动每线程4或8字节的数据。线程在warp中有索引，范围为0到warpSize-1。支持四种源lane寻址模式。
宽度参数：所有__shfl_sync()函数接受一个可选的宽度参数，用于改变函数行为。宽度值必须是[1, warpSize]范围内2的幂（例如，1, 2, 4, 8, 16或32）。其他值的结果未定义。
线程交互控制：通过掩码指定参与__shfl_sync()调用的线程，每个线程需在其掩码位设为1以确保硬件执行前正确收敛。目标线程非活动时，检索值未定义。
广播、扫描与归约示例：提供跨warp单值广播、子分区的包容性加法扫描及warp内归约操作的示例。
__nanosleep()函数：用于暂停线程大约ns纳秒，最大睡眠时间约为1毫秒，支持计算能力7.0及以上设备。
Warp矩阵函数：利用Tensor Cores加速特定形式的矩阵运算（D=A*B+C），支持混合精度浮点数，要求计算能力7.0及以上且warp内所有线程协作。在条件语句中使用这些操作需保证条件在整个warp中相同，否则可能导致挂起。
Tensor Cores扩展功能：
- 支持计算能力8.0及以上的设备上的替代浮点运算类型。
- 支持双精度浮点运算，需使用double类型的fragment，并采用.rn舍入模式。
- 预览特性：子字节操作，通过nvcuda::wmma::experimental命名空间访问，允许访问Tensor Cores的低精度能力，但API和数据结构可能随未来版本变化。
  以上内容覆盖了GPU编程中的线程间通信优化、时间管理、以及利用Tensor Cores进行矩阵运算优化的关键知识点。
Tensor Cores限制：不同主要和次要设备架构可能需要特殊的矩阵格式来利用Tensor Cores。线程仅持有矩阵片段（特定于架构的不透明ABI数据结构），开发者不能假设单个参数如何映射到参与矩阵乘积累加运算的寄存器上。片段是架构特定的，如果为不同但链接兼容的架构编译的函数A和B被链接到同一个设备可执行文件中，传递它们可能会导致大小和布局不一致，从而产生错误或潜在的数据损坏。例如，sm_70和sm_75之间碎片布局就不同。
避免链接问题：为避免架构不匹配问题，矩阵应始终存储到内存中通过外部接口传输，使用如wmma::store_matrix_sync()，然后以指针类型安全地传递给其他函数，例如float *dst。
元素类型与矩阵尺寸：Tensor Cores支持多种元素类型和矩阵尺寸组合，包括但不限于半精度、单精度、双精度以及实验性的亚字节操作。
DPX指令：DPX是一组函数，用于快速找到16位和32位整数参数的最大值、最小值以及融合加法和最大/最小值，适用于计算能力9.0及以上的硬件加速，或在旧设备上的软件仿真。DPX对于实现动态规划算法非常有用，如生物信息学中的Smith-Waterman或Needleman-Wunsch算法，以及路线优化中的Floyd-Warshall算法。
异步屏障：NVIDIA C++标准库引入了GPU版本的std::barrier，提供硬件加速的屏障操作，适用于计算能力8.0及以上设备，并与memcpy_async功能集成。该屏障支持细粒度的同步模式，允许程序在不同阶段进行线程同步，提高了异步编程的灵活性和效率。
CUDA Barrier的阶段管理：CUDA屏障（cuda::barrier）会根据参与线程调用bar.arrive()的预期到达次数递减计数，当计数达到零时，屏障完成当前阶段，并自动原子地重置计数，进入下一个阶段。通过bar.arrive()返回的cuda::barrier::arrival_token令牌与屏障当前阶段关联，用于在特定阶段阻塞或释放线程。
空间分区（Spatial Partitioning）：这是一种将线程块按空间划分的技术，使得不同warps可以执行独立计算，常用于生产者-消费者模式，其中一部分线程生成数据，另一部分并发消费。该模式要求双缓冲和两次单边同步来管理生产者与消费者间的数据缓冲区。
提前退出（Early Exit）：参与一系列同步操作的线程需要提前退出时，必须显式放弃参与，通过调用相应操作以确保不干扰其他线程的正常同步流程。
完成函数（Completion Function）：cuda::barrier可携带一个完成函数，在每个阶段最后一个线程到达后、任何线程从等待中解除阻塞前执行。此函数执行时，该阶段所有线程在屏障处执行的内存操作对执行函数的线程可见，反之亦然。
异步数据复制（Asynchronous Data Copies）：CUDA 11引入了memcpy_async API，允许设备代码显式管理数据移动的异步复制，实现计算与数据传输的重叠执行。memcpy_async操作与cuda::barrier、cuda::pipeline或cooperative_groups::wait同步原语协同工作，支持不同类型的异步数据复制场景，并在某些硬件上享受加速优势。
复制与计算模式：典型应用中，数据先从全局内存复制到共享内存，然后在共享内存上进行计算。通过使用或不使用memcpy_async，展示了如何表达和优化这一模式，特别是利用硬件加速特性减少内存访问延迟，提升性能。
memcpy_async的使用：memcpy_async是一个异步复制操作，用于从全局内存到共享内存的数据传输。在计算能力8.0或更高的设备上，这种传输可以利用硬件加速，避免通过中间寄存器传输数据。
cuda::barrier与异步数据复制：cuda::barrier的memcpy_async重载允许使用屏障同步异步数据传输。这使得所有参与屏障的线程完成复制操作并到达屏障时，屏障阶段才会推进，从而实现了异步数据复制与线程同步。
memcpy_async性能指导：
- 对齐要求：在计算能力8.0的设备上，为了获得最佳性能，需要确保共享和全局内存的指针都对齐到128字节，并且复制大小是4、8或16字节的倍数。
- 简单可复制类型：只有当复制的类型为TriviallyCopyable时，cp.async指令才能被用来加速memcpy_async。
- Warp纠缠：在同一线程束内的多个memcpy_async调用可能会导致“Warp纠缠”，影响性能。特别是在提交、等待和到达操作上，应尽量保证线程收敛以减少这种影响。
- 建议：推荐在执行提交和到达操作前，通过__syncwarp确保线程收敛，以避免因线程分歧带来的性能损失。
cuda::pipeline进行异步数据复制：CUDA提供了cuda::pipeline同步对象来管理并行的异步数据移动和计算。它作为一个双端多阶段队列，支持工作项的先进先出(FIFO)处理，适用于调度复杂的异步复制与计算重叠场景。
以上内容涵盖了使用memcpy_async进行高效内存管理的关键技术和注意事项，以及如何利用cuda::barrier和cuda::pipeline进一步优化异步数据处理的策略。
####本文档介绍了如何使用CUDA的cuda::pipeline特性来实现多阶段异步数据传输，以隐藏内存拷贝的延迟并重叠计算。关键点包括：
多阶段异步数据拷贝：通过cuda::pipeline，可以管理memcpy_async操作的批次序列，使得CUDA内核能够在执行计算的同时进行内存传输，从而隐藏数据传输的延迟。
管道对象（Pipeline Object）：作为双端队列，它按照先进先出（FIFO）顺序处理工作，生产者线程在管道头部提交工作，消费者线程则从管道尾部拉取工作。示例中，所有线程既扮演生产者也扮演消费者角色，一边提交新的memcpy_async操作，一边等待前一批操作完成。
资源管理与同步：cuda::pipeline_shared_state封装了有限资源，允许管道同时处理多个并发阶段的数据传输。当所有资源被占用时，生产者线程会阻塞直到消费者释放下一阶段的资源。管道内部使用共享内存中的屏障来进行线程间的同步。
性能优化：当所有线程均为生产者和消费者时，可以通过更紧密地集成循环的前序、后序部分与循环本身来简化代码，并利用pipeline<thread_scope_thread>结合__syncthreads()进一步优化性能。对于仅涉及同一warp内线程共享内存读取的操作，可以使用__syncwarp()。
Pipeline接口：提供了详细的API文档，支持灵活的线程参与模式，允许任意子集的线程成为生产者、消费者或两者兼备。同时，文档还提到了为C兼容性设计的Pipeline Primitives Interface。
Tensor Memory Access (TMA)：针对大规模数据移动的需求，Compute Capability 9.0引入了TMA，专为多维数组设计的高效数据传输机制，支持从全局内存到共享内存的快速拷贝。TMA特别适用于非连续访问模式的多维数组，且提供了对一维和多维数组的支持，通过硬件加速的地址计算减轻编程负担。TMA操作是异步的，支持多种完成机制，并针对特定架构进行了优化，如sm_90a上的多播功能。
综上所述，通过cuda::pipeline和TMA技术，开发者能够实现高效的内存管理和数据传输策略，以提升GPU计算的性能和效率，尤其是在处理大量数据移动和多维数组操作的应用场景中。

一维数组使用TMA传输

屏障初始化：通过参与线程数初始化屏障，确保所有线程到达后屏障才翻转。
异步批量复制：使用TMA进行大量数据的批量异步复制到共享内存，并在完成后更新共享内存屏障的事务计数。少量大尺寸的复制操作可提升性能。
事务会计：通过mbarrier.expect_tx指令告知屏障预期到达的字节数，确保所有数据到达后屏障翻转，数据可安全读取。
屏障等待与同步：使用mbarrier.try_wait等待屏障翻转，确保所有操作完成。
共享内存写入与同步：对缓冲区值的增益操作需确保写入共享内存对后续批量异步复制可见，使用特定指令排序写入操作。
写入全局内存与同步：由单一线程发起，使用线程局部机制等待所有操作完成，保证写入全局内存的可见性。

多维数组使用TMA传输

张量映射创建：通过CUDA驱动API（如cuTensorMapEncodeTiled）创建张量映射描述多维数组，包括基础指针、数组大小、行间步长等参数。
张量映射传递与使用：张量映射需位于不可变内存中，可通过常量内存或__grid_constant__参数传递给内核。示例展示了如何加载、修改多维数组的一块区域，并写回全局内存。
边界处理：超出界限的部分在读取时会导致共享内存对应区域被清零填充；写回时，部分区域可能越界，但起始索引不能为负。
尺寸与步长：强调了张量尺寸与步长的定义及其对齐要求，对理解多维数据布局至关重要。

PTX封装器

强调了用于多维TMA操作的PTX指令，如cp.async.bulk.tensor，提供了读取全局至共享内存及写入共享至全局内存的封装。
综上，文章详细阐述了利用TMA技术在GPU上进行高效的一维与多维数组数据管理与同步策略，涉及到了高级的内存操作与线程同步技术。

Profiler Counter Function: 每个多处理器有16个硬件计数器，应用程序可以通过调用__prof_trigger()函数单指令增量。计数器0到7的值可以通过nvprof工具获取，使用命令如nvprof --events prof_trigger_0x（x为0到7）。所有计数器在每个内核启动前重置。
Assertion: 仅支持计算能力2.x及以上的设备，当表达式等于零时停止内核执行。支持调试时触发断点或同步后向stderr打印错误信息。若表达式非零，则内核执行不受影响。推荐在生产代码中禁用断言以避免性能影响。
Trap Function: __trap()函数可由任何设备线程调用发起陷阱操作，导致内核执行中止并在主机程序中产生中断。
Breakpoint Function: __brkpt()函数允许从任何设备线程暂停内核函数的执行。
Formatted Output: 计算能力2.x及以上设备支持格式化输出。内核中的printf()函数类似于C库的printf()，将格式化的字符串输出到主机端流。支持特定的格式说明符，但最终格式化在主机上完成，可能依赖于主机操作系统。printf()限制最多接受32个除格式字符串外的参数，并且在64位Windows平台上的“%ld”格式可能因类型大小不同导致输出损坏。
Associated Host-Side API: 提供API函数来获取和设置用于传输printf()参数和内部元数据到主机的缓冲区大小，默认为1MB。需要注意，printf()可能导致线程执行顺序改变，因为其内部使用共享数据结构。
Examples: 示例代码演示了如何使用printf()函数在多线程环境中产生输出，以及如何通过条件语句控制输出，确保只有一条线程的输出被显示，展示了全局变量与局部变量在输出中的表现差异，以及如何管理输出缓冲区以避免数据丢失。
####动态全局内存分配与操作仅受计算能力2.x及以上版本的设备支持。这些功能允许从固定大小的全局内存堆中动态地分配和释放内存。CUDA内核中的malloc()函数可从设备堆中至少分配指定大小的字节，并返回指向已分配内存的指针，或在内存不足时返回NULL，且保证该指针对齐到16字节边界。而__nv_aligned_device_malloc()函数则允许指定内存对齐边界。free()函数用于释放由malloc()或__nv_aligned_device_malloc()先前分配的内存。
设备内存堆的大小是固定的，必须在使用malloc()、__nv_aligned_device_malloc()或free()的任何程序加载到上下文之前指定。默认情况下，如果没有明确指定堆大小且有程序使用了这些函数，则会分配8MB的堆。堆大小一旦在模块被加载到上下文中后就不能更改，且不会根据需求动态调整。通过设备端malloc()分配的内存不能直接用主机端的CUDA API（如cudaFree）释放，反之亦然。
执行配置定义了在设备上执行__global__函数时所使用的网格和线程块维度以及关联的流。调用__global__函数时必须指定执行配置，例如<<<Dg, Db, Ns, S>>>，其中Dg、Db、Ns、S分别代表网格维度、块维度、共享内存大小和流句柄。
计算能力9.0及以上允许编译时指定线程块簇维度，以便在CUDA中使用簇层次结构，这可以通过__cluster_dims__宏实现。同时，也可以在运行时指定簇维度并通过cudaLaunchKernelEx API启动带有簇的内核。
####GPU优化技术中的关键知识点包括：
启动界限（Launch Bounds）：通过在__global__函数定义中使用__launch_bounds__()限定符，开发者可以为编译器提供关于内核线程和块配置的额外信息，帮助编译器优化寄存器使用，减少寄存器溢出，并最小化指令计数。这有助于确保内核能够高效地利用多处理器资源，提升性能。启动界限还能防止内核因请求过多资源而失败。
寄存器使用控制：为了控制内核的寄存器使用量，除了__launch_bounds__外，还可以使用__maxnreg__函数限定符来直接指定单个线程的最大寄存器数量。同时，编译器选项maxrregcount可以控制文件中所有__global__函数的默认最大寄存器数量。
循环展开（Loop Unrolling）：使用#pragma unroll指令可以显式控制循环的展开程度，提高循环执行效率。编译器会根据指定的常量表达式或循环特性自动决定是否以及如何展开循环。
SIMD视频指令：针对计算能力3.0及以上的设备，PTX指令集提供了SIMD（单指令多数据）视频指令，用于加速处理16位值对和8位值四元组的操作。这些指令可以通过内联PTX汇编语法asm()在CUDA程序中使用，以实现更高效的并行处理。
诊断指令：CUDA编程中还支持诊断指令（pragmas），允许开发者控制编译器发出警告或错误的条件，例如抑制未使用变量的警告或改变特定诊断消息的严重性。但从CUDA 12.0起，没有nv_前缀的诊断指令在设备代码中将不再被支持，应使用带有nv_前缀的相应指令。
综上所述，通过合理设置启动界限、精确控制寄存器使用、有效利用循环展开以及应用特定硬件指令和编译器指令，开发者可以显著优化GPU代码的性能。

Cooperative Groups扩展：是CUDA编程模型的一个扩展，自CUDA 9起引入，旨在组织通信线程群体。它允许开发者明确线程通信的粒度，帮助实现更丰富、高效的并行分解。
历史同步限制与改进：传统上，CUDA编程模型仅提供了__syncthreads()内建函数来跨线程块的所有线程设置屏障以同步。Cooperative Groups扩展了这一功能，允许在更广泛的粒度上定义和同步线程群体，如单个warp或同一GPU上的线程块集，从而提高性能、设计灵活性及软件复用性。
安全高效的并行交互：为了解决并行交互模式的表达问题，Cooperative Groups提供了一个安全且面向未来的方法来实现高性能代码，避免了以往程序员需要自行编写不安全且易碎的同步原语的问题。
编程模型概念：
- 强调了将线程组作为一级程序对象，通过显式对象表示参与线程集合，改善了软件组成性，使程序员意图明确，有助于消除不稳定的架构假设，减少对编译器优化的限制，并提高了与新GPU代际的兼容性。
- 引入了新的启动API，施加特定限制以保证同步操作成功，支持新的合作并行模式，如生产者-消费者并行、机会并行以及整个Grid范围内的全局同步。
组类型与隐式组：介绍了线程块组（thread_block）作为一种已知的线程组，以及如何使用thread_block类进行同步操作。强调了隐式组的处理应谨慎，推荐在分支发生前创建组句柄以避免死锁或数据损坏。
实践建议：推荐使用专门化的组而非泛型，以利用编译时优化，并通过引用传递这些组对象给合作函数。要求CUDA 9.0或更高版本，并通过包含特定头文件和命名空间来使用Cooperative Groups。
代码示例：展示了如何使用Cooperative Groups改进块级求和归约操作的代码写法，通过显式定义和同步线程组来提升代码的清晰度和效率。
综上所述，Cooperative Groups通过提供一种机制以更细粒度和更安全地组织和同步GPU上的线程，促进了并行算法设计的灵活性和效率。

Cluster Group：
- 表示在单个集群中启动的所有线程，适用于计算能力9.0及以上的硬件。
- 提供同步（sync）、障碍到达（barrier_arrive）和障碍等待（barrier_wait）等API，以及查询线程、块排名、数量和维度的方法。
- 支持获取共享内存地址的块排名和映射其他块中的共享内存变量地址。
Grid Group：
- 代表单个网格中启动的所有线程，使用合作启动API可以跨网格同步。
- 包含网格、块、集群级别的排名和尺寸查询，以及块和集群索引信息。
- 同样提供同步和障碍操作相关的API。
Multi Grid Group（已废弃）：
- 代表多设备合作启动中所有设备上的线程，需要使用特定的启动API。
- 提供线程数量、排名、网格排名和网格数量的查询，但在CUDA 11.3中已被废弃。
显式组（Explicit Groups）- Thread Block Tile：
- 一种基于模板的分块组，允许编译时指定块大小，可能实现更优执行。
- 支持同步、线程数量和排名查询，以及元组大小和排名查询。
- 提供类似Warp Shuffle、Vote和Match功能，但部分功能受限于块大小（通常小于等于32）。
Warp-Synchronous Code Pattern：
- 开发者需明确指定线程块大小，代替以往对warp大小的隐式假设。
单线程组：
- 使用this_thread函数获取表示当前线程的组，可用于执行异步数据复制等操作。

####在CUDA的SIMT架构中，硬件层面的多处理器以32个线程为一组执行，称为warps。当应用程序代码存在数据依赖条件分支导致同一warp内的线程发生分歧时，该warp会串行执行每个分支，并禁用不在此路径上的线程。保持在路径上活跃的线程被称为聚集(coalesced)。Cooperative Groups具有发现并创建包含所有聚集线程的组的功能。
coalesced_threads()构建组句柄是机会主义的，它返回该时刻的活跃线程集合，但不保证返回哪些线程（只要它们是活跃的）或它们将在整个执行过程中保持聚集（它们将在执行集体操作时被重新聚集，但之后可能再次分歧）。
Coalesced Group类提供了一系列功能，如同步线程、获取组内线程总数、获取调用线程在组内的排名、执行warp shuffle函数、warp投票函数和warp匹配函数等。

CUDA SIMT架构：多处理器以32个线程组成的warps执行，遇到条件分支时可导致线程分歧，聚集线程(coalesced threads)指在特定路径上保持一致执行的线程。
Coalesced_Groups功能：允许开发者发现并创建包含当前活跃且聚集的线程组，通过coalesced_threads()方法实现。此功能是机会性的，不保证线程持续聚集。
线程同步与操作：提供了线程同步(sync)、获取组信息(num_threads, thread_rank等)以及执行warp级别的数据交换(shuffle)、投票(vote)和匹配(match)等集体操作。
组划分功能：
- tiled_partition：将父组划分为一维、行主序的子组，适用于需要对线程块进行细粒度控制的场景。
- labeled_partition：根据标签值将父组划分为线程聚集的子组，适用于按条件分组线程的需求。
- binary_partition：根据布尔型谓词将线程分为两组，是labeled_partition的特例化，标签只能是0或1。
组集体操作：包括同步（如barrier_arrive和barrier_wait）、数据传输等，要求组内所有线程参与完成操作，增强了并行任务的协同能力。
这些知识点围绕GPU优化中的线程管理、同步与数据操作展开，是提升CUDA编程效率和性能的关键技术。
memcpy_async：这是一个组内集体memcpy操作，利用硬件加速支持，实现从全局内存到共享内存的非阻塞内存事务。为了获得最佳性能，要求共享内存和全局内存都具有16字节对齐。此操作在一般情况下是memcpy，仅当源为全局内存、目标为共享内存且两者均可使用16、8或4字节对齐时才是异步的。异步复制的数据应在调用wait或wait_prior之后读取，以确保对应阶段已完成数据移动到共享内存。
等待机制：wait和wait_prior集合操作允许等待memcpy_async复制完成。wait会阻塞调用线程直到所有之前的复制完成；而wait_prior允许最新的NumStages个操作未完成，它会等待所有之前（不包括最后NumStages个）的请求完成。这两个函数都会同步命名的线程组。
数据处理：
- reduce：在提供的线程组中执行数据的归约操作，利用硬件加速（在计算能力8.0及以上设备上）进行算术加、最小值或最大值以及逻辑与、或、异或操作，并为旧硬件提供软件回退。仅4字节类型能得到硬件加速。
- inclusive_scan和exclusive_scan：在每个线程提供的数据上执行扫描操作，结果为每个线程根据其线程排名较低的线程数据的累积（exclusive_scan）或包含调用线程数据的累积（inclusive_scan）。支持加法、比较和位运算等操作对象。
执行控制：虽然具体细节未展开，但提到了与执行控制相关的章节标题，暗示了还包括有关如何控制和协调线程执行的内容。
这些知识点围绕CUDA编程中的数据传输、同步、数据处理和执行控制，特别强调了在现代GPU架构上实现高效数据移动和操作的技术。
invoke_one与invoke_one_broadcast：这两个函数允许在调用线程组中选择一个线程执行给定的可调用对象fn及参数args。其中，invoke_one_broadcast还会将调用结果广播给调用线程组内的所有线程。调用期间，组内通信被限制以保证前向进度，但与组外线程的通信是允许的。在计算能力9.0或更高的设备上，硬件加速可用于确定选择的线程。
网格同步（Grid Synchronization）：在引入Cooperative Groups之前，CUDA编程模型仅允许在核函数结束时在块间进行同步。现在，通过grid.sync()函数可在核函数内部实现网格级别的同步，需要使用cudaLaunchCooperativeKernel启动核函数。为了确保线程块在GPU上的共居，需精心设计启动的块数量，考虑如SM数量或使用占用率计算器来最大化并行度。
多设备同步：使用cudaLaunchCooperativeKernelMultiDevice API可以在多个设备上启用同步，要求为参与设备启用对等访问，并且代码需要分开编译。支持此特性的设备需具备6.0或更高计算能力，并且运行在Linux（无MPS）或Windows（TCC模式）上。
CUDA动态并行（Dynamic Parallelism）：这一模型扩展允许CUDA核函数在GPU上直接创建和同步新的任务，适用于计算能力3.5及以上的设备。它减少了主机与设备间控制和数据传输的需求，使得数据依赖的并行工作可以在运行时动态生成，适应动态变化的工作负载，特别适合递归、不规则循环或其他非单一层次并行结构的算法表达。
这些知识点涉及了GPU编程中的高级同步机制、多设备协作以及动态生成并行任务的能力，对于追求高性能计算和复杂并行算法实现的开发者来说至关重要。

9.2.1.2 CUDA原语的作用域

CUDA运行时在主机和设备上提供了API，用于启动内核并通过流和事件跟踪启动之间的依赖关系。
主机系统中，所有线程共享启动状态及引用流和事件的CUDA原语；进程独立执行，不可共享CUDA对象。
设备上，已启动的内核和CUDA对象对网格中的所有线程可见。

9.2.1.3 同步

CUDA运行时操作（包括内核启动）对网格中的所有线程可见，允许父网格中的调用线程同步控制由网格中任何线程在任一流上启动的网格顺序。
网格的执行直到网格中所有线程的所有启动完成才被视为结束。

9.2.1.4 流和事件

流和事件控制网格启动间的依赖：同一流中的网格按序执行，事件可用于在流之间创建依赖。
在网格内创建的流和事件具有网格作用域，超出该网格使用则行为未定义。
所有工作在网格退出时隐式同步，包括流中的工作及其适当解决的依赖。
主机上创建的流和事件在内核中使用时行为未定义。

9.2.1.5 顺序和并发

设备运行时的内核启动遵循CUDA流顺序语义。
网格内的内核到同一流的启动（忽略特殊流）按序执行，依赖于网格内的线程调度。
隐式NULL流只在同一个线程块内的所有线程间共享，跨线程块并发需使用显式命名流。

9.2.1.6 设备管理

设备运行时不支持多GPU，仅能在当前执行的设备上操作，但允许查询系统中任何CUDA兼容设备的属性。

9.2.2 内存模型

父子网格共享全局和常量内存存储，但拥有独立的局部和共享内存。

9.2.2.1 一致性与连贯性

全局内存：父子网格间有连贯访问，弱一致性保证。只有在子网格被父线程调用的时刻，其内存视图与父线程完全一致。
零拷贝内存：与全局内存有相同的一致性和连贯性保证。
常量内存：不可在设备上修改，主机修改同时被网格访问的行为未定义。
局部和共享内存：私有于线程块或线程，不在父子间可见或连贯，越界引用行为未定义。
注意，向cudaMemcpy*Async()或cudaMemset*Async()传递共享或局部内存指针是非法的，会返回错误，因为这些API可能为保持流语义而在设备上调用新的子内核。

本地内存（Local Memory）：本地内存是执行线程的私有存储空间，不对外部可见。尝试将指向本地内存的指针作为参数传递给子核函数启动是非法的，且子核函数访问此类指针的解引用结果未定义。编译器何时将变量置于本地内存中有时难以让程序员察觉，因此建议所有分配给子核函数使用的存储都应显式地从全局内存堆中分配，如使用cudaMalloc()、new()或在全局作用域声明__device__存储。
纹理内存（Texture Memory）：对映射了纹理的全局内存区域的写入与纹理访问不一致。纹理内存的一致性在子网格启动时和子网格完成时强制执行。这意味着，在子核函数启动前对内存的写入会反映在子核函数的纹理内存访问中。然而，子核函数对内存的写入并不保证能被父核函数的纹理内存访问看到。要访问子网格线程修改的内容，必须通过cudaStreamTailLaunch流中启动的核函数。父母和孩子并发访问可能导致数据不一致。
编程接口（Programming Interface）：描述了支持动态并行性的CUDA C++语言扩展的变化和新增内容。设备运行时（Device Runtime）为CUDA内核提供了类似于主机上CUDA运行时API的接口和API。这使得可以在主机或设备环境中运行的例程代码重用变得容易。设备运行时API是每个线程的代码，允许每个线程独立决定接下来执行哪个核函数或操作，且无需在线块内的线程间进行同步即可调用任何提供的设备运行时API。
设备端核函数启动（Device-Side Kernel Launch）：使用标准CUDA <<< >>> 语法从设备上启动核函数。这些启动是异步的，与启动线程立即返回，并继续执行直到遇到隐式的启动同步点（例如在cudaStreamTailLaunch流中启动的核函数）。子网格配置（如共享内存和L1缓存大小）将从父级继承。
流（Streams）：设备运行时支持命名和未命名（NULL）流。命名流可在网格中的任何线程中使用，但流句柄不能传递给其他子/父核函数。设备端不支持主机端NULL流的跨流屏障语义，且所有设备流必须使用cudaStreamCreateWithFlags()并传入cudaStreamNonBlocking标志创建。使用cudaStreamTailLaunch流的核函数可作为等待子核函数完成的替代方案。
遗忘流（Fire-and-Forget Stream）：提供了一种快速、低开销的方式来启动不需要显式跟踪的“遗忘”工作。它比为每次启动创建新流更高效，但不能用于记录或等待事件。使用遗忘流需要64位模式编译，并且不支持在定义了CUDA_FORCE_CDP1_IF_SUPPORTED时使用。
尾部启动流（Tail Launch Stream）：cudaStreamTailLaunch允许一个网格在完成执行后调度新的网格启动。通常情况下，它能够替代cudaDeviceSynchronize()的功能。每个网格拥有独立的尾部启动流，且非尾部启动的工作在尾部流启动前会隐式同步。网格通过尾部流启动时，必须等待父网格及其所有子工作（包括普通流、每线程流或fire-and-forget流中的工作）全部完成。
事件（Events）：仅支持CUDA事件的跨流同步功能，即支持cudaStreamWaitEvent()，但不支持cudaEventSynchronize()、cudaEventElapsedTime()和cudaEventQuery()。创建事件时需使用cudaEventCreateWithFlags()并设置cudaEventDisableTiming标志。事件对象在同一网格的所有线程间可共享，但对其他网格不可见，且事件句柄在不同网格间不保证唯一性。
同步（Synchronization）：程序需负责执行足够的线程间同步，如使用CUDA事件，以确保调用线程与从其他线程启动的子网格同步。父线程无法显式同步子工作，因此无法确保子网格中的更改对父网格线程可见。
设备管理（Device Management）：仅运行内核的设备可由该设备控制，意味着cudaSetDevice()等设备API不受设备运行时支持。从GPU视角（通过cudaGetDevice()）看到的活跃设备号与主机系统相同。虽然cudaDeviceGetAttribute()可请求其他设备信息，但设备运行时不提供cudaGetDeviceProperties()这样的全面API，属性需单独查询。
内存声明（Memory Declarations）：
- 设备和常量内存：使用__device__或__constant__指定的内存，在设备运行时表现相同。所有内核可读写设备变量，无论最初由主机还是设备运行时启动。
- 纹理和表面（Textures and Surfaces）：动态创建的纹理和表面对象可在主机上创建、传递给内核使用、然后从主机销毁。设备代码内不允许创建或销毁这些对象，但由主机创建的对象可在设备上自由使用。
- 共享内存变量声明：支持静态大小的文件作用域或函数作用域的共享内存变量声明，以及由运行时配置确定大小的外部变量声明。
- 符号地址：设备端符号可通过&直接引用，全局作用域的设备变量位于内核的可见地址空间中。这也适用于__constant__符号，但访问是只读的。
API错误与启动失败：任何函数都可能返回错误码，通过cudaGetLastError()获取最近的错误。设备端启动也可能失败，用户需调用cudaGetLastError()检查错误，但无错误不代表子内核成功执行。设备端异常会被返回到主机。
API参考：详细说明了设备运行时支持的CUDA运行时API部分，强调了主机和设备API的语法相似性及特定语义差异。
从PTX的设备端启动：为面向PTX的编程语言和编译器实现者提供了低级细节，涉及在PTX级别支持内核启动。
Kernel Launch APIs：
- cudaLaunchDevice()：用于在设备端启动指定的内核，需要一个通过cudaGetParameterBuffer()获取并填充了启动内核所需参数的参数缓冲区。若内核无需参数，可不调用cudaGetParameterBuffer()。
- cudaGetParameterBuffer()：用于获取参数缓冲区，需在PTX级别声明。参数指定了缓冲区的对齐和大小需求，当前实现保证返回的缓冲区至少64字节对齐，但建议传入正确的对齐值以确保未来的可移植性。
Parameter Buffer Layout：参数缓冲区禁止参数重排序，每个参数必须按其大小对齐放置。缓冲区最大4KB。
Toolkit Support for Dynamic Parallelism：
- Including Device Runtime API：CUDA设备运行时API的原型在编译时会自动包含，无需显式包含cuda_device_runtime_api.h。
- Compiling and Linking：使用动态并行性时，nvcc会自动链接静态设备运行时库libcudadevrt。程序需要链接此静态库，并提供了命令行编译和分步编译链接的示例。
Programming Guidelines：
- Basics：设备运行时是主机运行时的功能子集，支持API级别的设备管理、内核启动等。动态并行编程与标准CUDA编程相似。
- Performance：动态并行启用的内核可能因设备运行时的执行跟踪和管理软件而产生额外开销，影响性能。
- Implementation Restrictions：包括内存足迹、待处理的内核启动限制、配置选项、内存分配与生命周期、SM ID和Warp ID的不稳定性以及ECC错误处理等硬件和软件资源的限制。
  这些内容覆盖了动态并行计算在GPU编程中的关键API使用、参数处理、性能考量及实施限制等方面，对进行GPU优化和深入理解CUDA编程模型有重要指导意义。

CUDA动态并行性(CDP)差异与兼容性

CDP2与CDP1的主要区别：
- CDP2不再支持显式设备端同步，必须使用隐式同步（如尾部启动）。
- 在CDP2或计算能力9.0及以上的设备上，查询或设置cudaLimitDevRuntimeSyncDepth会导致错误。
- CDP2没有为不适应固定大小池的待处理启动提供虚拟化池，需要通过设置cudaLimitDevRuntimePendingLaunchCount避免启动槽溢出。
- CDP2限制了同时存在的事件总数，等于待处理启动计数的两倍，需合理设置避免事件槽不足。
- CDP2按网格跟踪流，而非按线程块，允许工作被启动到其他线程块创建的流中。
- CDP2引入了尾部启动和fire-and-forget命名流，并仅支持64位编译模式。
兼容性和互操作性：
- CDP2是默认模式，但可以通过编译选项选择在低版本设备上使用CDP1。
- CDP1和CDP2函数可以在同一上下文中同时加载和运行，各自使用专属特性，但不能交叉调用，否则会报错。

CUDA动态并行性(CDP1)概述

执行环境与内存模型：
- CUDA基于线程、线程块和网格的执行模型，动态并行性扩展了在设备上配置、启动和同步新网格的能力。
- 父网格与子网格的定义与嵌套执行，确保父网格完成前所有子网格完成，即使没有显式同步。
- CUDA原语的作用域在设备上遵循特定规则，如流和事件在同一线程块内共享，但跨线程块则行为未定义。
- 同步机制确保线程块内的所有启动完成才认为该块执行结束，且提供了控制依赖关系的机制。
- 流和事件的使用需注意作用域限制，以及在不同上下文中的未定义行为。
- 订单和并发性遵循CUDA流的顺序语义，尽管动态并行性易于表达并发，但设备运行时并未引入新的并发保证，尤其在不同线程块之间或父线程块与其子网格之间。
  综上，CDP2带来了对并行管理的改进，包括同步机制、资源管理以及对流和事件使用的更严格规范，同时也强调了与旧版CDP1接口的兼容性策略及限制。
设备管理：在当前文档版本中，设备运行时并不支持多GPU。尽管如此，允许查询系统中任何CUDA兼容设备的属性。
内存模型：
- 父网格和子网格共享相同的全局和常量内存存储，但具有独立的局部和共享内存。
- 全局内存一致性与连贯性：父网格与子网格对全局内存的访问是一致的，但在子网格与父网格间有较弱的一致性保证。子网格执行前后存在两个点（启动时和父线程同步于子网格完成时）其内存视图与父线程完全一致。
- 零拷贝内存：零拷贝系统内存与全局内存具有相同的连贯性和一致性保证，遵循相同语义。内核不能分配或释放零拷贝内存，但可以使用从主机程序传递进来的零拷贝内存指针。
- 常量内存：常量不可变，即使在父网格和子网格之间也不能从设备端修改。所有__constant__变量的值必须在启动前由主机设置。常量内存自动被所有子内核从其相应的父内核继承。
- 共享与局部内存：共享内存私属于线程块，局部内存私属于线程，两者在父网格和子网格间不可见且不连贯。超出所属作用域引用这些位置的对象行为未定义，可能导致错误。
- 局部内存：局部内存是执行线程的私有存储，对线程外不可见。将指向局部内存的指针作为参数传递给子内核启动是非法的，如果子内核中访问这样的局部内存地址，其结果是未定义的。
编程接口：介绍了CUDA C++扩展的更改和新增内容以支持动态并行性。设备运行时API在语法和语义上尽可能保留了CUDA运行时API的风格，以便于代码在主机和设备环境中的重用。
设备端内核启动：内核可以通过标准CUDA <<< >>> 语法从设备上启动，且所有设备端内核启动都是异步的，与发起线程立即返回，直到遇到如cudaDeviceSynchronize()之类的显式同步点才会等待子网格完成。

启动环境配置(CDP1): 全局设备配置设置（如共享内存和L1缓存大小）会从父级继承。对于主机启动的内核，主机设置的每内核配置优先于全局设置，并在设备上启动时同样适用。无法从设备重新配置内核环境。
流(Streams, CDP1): 设备运行时支持命名和未命名流。命名流可在线程块内的任何线程使用，但句柄不能传递给其他块或子/父内核。设备上不支持主机端NULL流的跨流屏障语义，所有设备流必须使用cudaStreamCreateWithFlags() API创建并传入cudaStreamNonBlocking标志。设备运行时不支持cudaStreamSynchronize()和cudaStreamQuery()，应使用cudaDeviceSynchronize()来确认流启动的子内核已完成。
隐式(NULL)流(CDP1): 设备运行时为每个线程块提供一个单一的隐式未命名流，但由于命名流需使用非阻塞标志，NULL流中的工作不会对其他流中的待处理工作产生隐式依赖。
事件(Events, CDP1): 仅支持CUDA事件的跨流同步能力。cudaEventSynchronize(), cudaEventElapsedTime(), 和 cudaEventQuery()不被支持，事件必须通过cudaEventCreateWithFlags()并带cudaEventDisableTiming标志创建。事件对象是线程块本地的，且句柄在不同块间不保证唯一。
同步(Synchronization, CDP1): cudaDeviceSynchronize()会同步线程块中所有线程直至调用点的所有工作。需要额外的线程间同步（如通过__syncthreads()）以确保与其他线程启动的子网格同步。
块宽同步(Block Wide Synchronization, CDP1): cudaDeviceSynchronize()不隐含块内同步，需要显式同步如__syncthreads()来确保所有线程已提交工作。
设备管理(Device Management, CDP1): 只有运行内核的设备可由该内核控制，意味着cudaSetDevice()等设备API不受设备运行时支持。活动设备号在GPU和主机系统中相同。可以查询其他设备属性，但不提供cudaGetDeviceProperties()。
内存声明(Memory Declarations, CDP1): 文件作用域中使用__device__或__constant__声明的内存空间在使用设备运行时时行为一致。所有内核可读写设备变量。动态创建的纹理对象始终有效，可以从父内核传递给子内核。
共享内存变量声明(Shared Memory Variable Declarations, CDP1): 共享内存可以在文件作用域或函数作用域静态声明，或作为外部变量由运行时确定大小，两种声明方式在设备运行时下都有效。
设备端符号引用（Device-side Symbols）: 在CUDA内核中，通过&操作符可直接引用全局作用域的__device__符号，包括__constant__符号，但对__constant__空间的引用是只读的。这表明在运行中的内核中，即使是子内核启动前，常量数据也无法修改。
API错误与启动失败（API Errors and Launch Failures）: 设备端内核启动可能因多种原因失败，如无效参数等。需使用cudaGetLastError()检查是否产生错误，但无错误不代表子内核成功执行。设备端异常会返回到主机，而非由父内核的cudaDeviceSynchronize()捕获。
启动设置APIs（Launch Setup APIs）: 设备端运行时库通过cudaGetParameterBuffer()和cudaLaunchDevice()API直接暴露内核启动机制。应用程序可直接调用这些API，需遵循PTX相同要求，并负责正确填充所有数据结构。
API参考（API Reference）: 描述了在设备运行时支持的CUDA运行时API部分，主机和设备运行时API语法相同，除非特别说明，语义也一致。
从PTX的设备端启动（Device-side Launch from PTX）: 面向实现动态并行性的PTX目标编程语言和编译器开发者，提供在PTX级别支持内核启动的底层细节。
内核启动APIs（Kernel Launch APIs）: 设备端内核启动可通过PTX访问的cudaLaunchDevice()和cudaGetParameterBuffer()实现。前者用于启动内核并传递参数缓冲区，后者获取参数缓冲区，参数缓冲区布局有特定要求以保证对齐和顺序。
工具包对动态并行的支持（Toolkit Support for Dynamic Parallelism）: 包括自动包含设备运行时API原型、编译和链接时自动链接静态设备运行时库libcudadevrt等，简化开发过程。

编程指南(CDP1)

设备运行时是主机运行时的功能子集，提供API级别的设备管理、内核启动、设备内存复制、流管理和事件管理。
对于有CUDA经验的开发者来说，为设备运行时编程应较为熟悉，其语法和语义大多与主机API相同，本文档中已详细说明了任何例外情况。
建议仅在需要在线程块结束前与子内核同步时调用cudaDeviceSynchronize()，以避免性能损失。

性能(CDP1)

线程间的同步可能影响同一线程块中其他线程的性能，即使这些线程本身未调用cudaDeviceSynchronize()。隐式地在线程块结束时同步子内核通常比显式调用cudaDeviceSynchronize()更高效。
控制动态启动的系统软件可能会对任何同时运行的内核（无论是否自行调用内核启动）施加开销，这可能导致性能下降，特别是在设备端调用库函数时对比主机端调用。

实现限制与约束(CDP1)

动态并行保证了文档中描述的所有语义，但某些硬件和软件资源的实现依赖性限制了使用设备运行时的程序的规模、性能等属性。
设备运行时系统软件为各种管理目的预留内存，特别是用于同步子启动期间保存父网格状态的预留，以及跟踪待处理网格启动的预留。可通过配置选项减少这些预留的大小，但会受到一定的启动限制。
最大嵌套深度限制为24层，但实际上限制因素是系统为每一新层所需的内存量（见内存占用）。超过最大深度的启动将失败。
可通过cudaDeviceSetLimit()控制最大同步深度及预留存储，以避免过度分配备份存储。
设备运行时还管理着一个固定大小和虚拟化的启动池来追踪内核的配置和参数数据，这些池的大小可以通过API进行配置。

内存分配与生命周期(CDP1)

cudaMalloc()和cudaFree()在主机和设备环境中的语义不同。在设备环境中，它们映射到设备端的malloc()和free()，这意味着可分配内存受限于设备malloc()堆的大小，且不能跨环境混用分配和释放操作。

SM Id和Warp Id(CDP1)

注意PTX中的%smid和%warpid被定义为易失值。设备运行时可能会为了更有效地管理资源而重新调度线程块到不同的SM上，因此依赖%smid或%warpid在整个线程或线程块生命周期中保持不变是不安全的。

ECC错误处理：CUDA内核代码无法直接接收到ECC（Error Correction Code）错误的通知。ECC错误在完成整个启动树后，于主机端报告。执行嵌套程序期间出现的ECC错误会根据错误类型和配置产生异常或继续执行。
虚拟内存管理API：
- 引入于CUDA 10.2，这些API允许应用直接管理CUDA提供的统一虚拟地址空间，以映射物理内存到GPU可访问的虚拟地址。
- 提供了与OpenGL、Vulkan等其他进程和图形API交互的新方式，并提供了用户可调优的新型内存属性以适应应用需求。
- 解决了历史问题，即cudaMalloc等内存分配调用返回的GPU内存地址无法按需调整大小，引入了类似malloc但缺乏realloc功能的问题。虚拟内存管理API分离了地址和内存的概念，允许应用独立处理它们。
- 允许应用程序选择特定的内存分配让目标设备访问，而不是默认映射所有cudaMalloc分配到对等设备，从而避免了不必要的运行时成本。
- 提供细粒度控制，让用户可以管理应用中的GPU内存，包括分配、释放、映射、解映射物理内存以及导出和导入共享句柄等操作。
- 支持压缩内存类型分配，加速访问具有非结构化稀疏性和其他可压缩数据模式的数据，节省DRAM带宽和L2缓存资源，需要设备支持计算数据压缩特性。

使用前提：

确保系统支持统一虚拟寻址（UVA）。
在使用虚拟内存管理API前，应用需检查目标设备是否支持CUDA虚拟内存管理。

关键API功能：

cuMemCreate：创建没有设备或主机映射的物理内存块。
cuMemAddressReserve：预留虚拟地址范围。
cuMemMap：将物理内存映射到预留的VA范围。
cuMemSetAccess：设置映射内存的访问权限。
cuMemExportToShareableHandle 和 cuMemImportFromShareableHandle：用于跨进程通信的内存导出和导入。

注意事项：

分配的内存大小需符合特定的对齐粒度要求。
需查询并确保所请求的内存类型和句柄类型在目标设备上受支持。
应用需负责释放通过cuMemRelease分配的内存。

10.4 虚拟地址范围预留

虚拟内存管理中，地址和内存是分离的概念。应用程序需预留一个地址范围以容纳由cuMemCreate创建的内存分配。预留的地址范围至少要与用户计划放置其中的所有物理内存分配的总大小一样大。通过向cuMemAddressReserve传递适当参数来预留虚拟地址范围，此范围不会有任何设备或主机物理内存关联。预留的虚拟地址范围可映射到系统中任何设备的记忆块，为应用提供由不同设备内存支持和映射的连续VA范围。应用程序应使用cuMemAddressFree将虚拟地址范围返回给CUDA，并确保在调用前整个VA范围未映射。

10.5 虚拟别名支持

虚拟内存管理API允许通过多次调用cuMemMap并使用不同的虚拟地址，为同一分配创建多个虚拟内存映射或“代理”，即虚拟别名。除非PTX指令集另有说明，对分配的一个代理的写入在写入设备操作（网格启动、memcpy、memset等）完成之前被认为是与其他任何代理不一致和不协调的。在写入设备操作完成之前位于GPU上的网格，但在写入设备操作完成后进行读取的，也被视为具有不一致和不协调的代理。

10.6 内存映射

已分配的物理内存和前两节中划分的虚拟地址空间体现了虚拟内存管理API引入的内存与地址的区别。为了使分配的内存可用，用户必须首先将其放置在地址空间中。通过使用cuMemMap，将从cuMemAddressReserve获得的地址范围和从cuMemCreate或cuMemImportFromShareableHandle获得的物理分配相互关联。用户可以将来自多个设备的分配映射到连续的虚拟地址范围中，只要他们预留了足够的地址空间。为了分离物理分配和地址范围，用户必须使用cuMemUnmap取消映射映射的地址。

10.7 控制访问权限

虚拟内存管理API允许应用程序明确地使用访问控制机制保护其VA范围。使用cuMemMap将分配映射到地址范围的区域并不会使该地址可访问，如果被CUDA内核访问会导致程序崩溃。用户必须特别使用cuMemSetAccess函数选择访问控制，该函数允许或限制特定设备对映射地址范围的访问。

11. 流顺序内存分配器

流顺序内存分配器允许应用程序将内存分配和释放与CUDA流中启动的其他工作（如内核启动和异步拷贝）进行排序，从而减少GPU跨所有执行CUDA流的同步。这通过利用流排序语义来重用内存分配，改善了应用程序的内存使用。分配器还允许应用程序控制其内存缓存行为，支持内存分配在进程间的轻松安全共享，有助于提升应用程序性能和内存管理效率。

11.3. API基础（cudaMallocAsync和cudaFreeAsync）

cudaMallocAsync和cudaFreeAsync是分配器的核心API，分别用于分配和释放内存。
这两个API接受流参数来定义内存何时可用或不可用。
cudaMallocAsync同步确定返回的指针值，并可用于构建后续任务。它根据指定的内存池或提供的流决定内存驻留设备，而非当前设备/上下文。
在非分配流中使用分配的内存时，必须保证访问发生在分配操作之后，否则行为未定义。可通过同步分配流或使用CUDA事件来确保生产与消费流的同步。
cudaFreeAsync()在流中插入释放操作，用户需保证释放操作在分配及任何使用后进行，否则行为未定义。应使用事件或流同步操作确保其他流对分配的访问在开始释放前完成。

11.4. 内存池与cudaMemPool_t

内存池管理虚拟地址和物理内存资源，按池属性和特性分配。
所有cudaMallocAsync调用都使用内存池资源。未指定内存池时，使用流设备的当前内存池。
可通过cudaDeviceSetMempool设置设备的当前内存池，并通过cudaDeviceGetMempool查询。默认情况下，当前内存池为设备的默认内存池。
cudaMallocFromPoolAsync允许用户为分配指定内存池，而不将其设为当前池。
cudaDeviceGetDefaultMempool和cudaMemPoolCreate提供获取内存池句柄的API。

11.5. 默认/隐式池

设备的默认内存池可通过cudaDeviceGetDefaultMempool获取，其分配的内存不可迁移且位于该设备上。默认池的访问性可修改。
因无需显式创建，默认池有时称为隐式池。默认内存池不支持IPC（进程间通信）。

11.6. 显式池

cudaMemPoolCreate创建显式池，允许应用请求超出默认池特性的属性，如IPC能力、最大池大小、特定CPU NUMA节点上的分配等。

11.7. 物理页缓存行为

分配器默认尝试最小化池拥有的物理内存。应用需通过释放阈值属性(cudaMemPoolAttrReleaseThreshold)配置每个池的内存占用。
释放阈值是池尝试释放内存回操作系统之前的保持量。超过此阈值时，同步操作将尝试释放内存给OS。设置为UINT64_MAX可防止每次同步后尝试缩小池大小。
cudaMemPoolTrimTo允许显式减小内存池的内存占用。

11.8. 资源使用统计

CUDA 11.3新增了查询池内存使用的属性，如已保留和已使用内存的当前和最高值，并可重置这些属性的高水位标记。

11.9. 内存重用策略

驱动程序尝试通过cudaFreeAsync释放的内存来满足新的分配请求，以重用内存。
分配策略受内存池属性控制，如遵循事件依赖、允许机会主义重用和允许内部依赖性重用。这些策略可由更新驱动程序改变或增强。

禁用重用策略：尽管可控的重用策略可以改善内存重用，但用户可能需要禁用它们。允许机会性重用（如cudaMemPoolReuseAllowOpportunistic）会根据CPU和GPU执行的交错情况引入运行到运行的分配模式变化。内部依赖插入（如cudaMemPoolReuseAllowInternalDependencies）可能会在用户更愿意显式同步事件或流以应对分配失败时，以不可预测的方式序列化工作。
多GPU支持的设备可访问性：内存池分配的可访问性不由cudaDeviceEnablePeerAccess或cuCtxEnablePeerAccess控制，而是通过cudaMemPoolSetAccess API修改哪些设备可以访问来自某个池的分配。默认情况下，分配只能从进行分配的设备访问，且这种访问不能被撤销。要允许其他设备访问，访问设备必须与内存池设备具有对等能力；使用cudaDeviceCanAccessPeer检查。若未检查对等能力，设置访问权限可能会因cudaErrorInvalidDevice而失败。如果池中尚未进行任何分配，即使设备不是对等的，cudaMemPoolSetAccess调用也可能成功；在此情况下，从池中的下一次分配将失败。
IPC内存池：IPC功能的内存池允许进程间轻松、高效且安全地共享GPU内存。CUDA的IPC内存池提供与虚拟内存管理API相同的安全优势。共享分为两个阶段：首先是共享池访问，然后是共享特定的池内分配。创建和共享IPC内存池涉及使用cudaMemPoolExportToShareableHandle API获取池的OS原生句柄，通过常规OS IPC机制传输该句柄，并在导入进程中使用cudaMemPoolImportFromShareableHandle API创建导入的内存池。
导入进程中的访问设置：导入的内存池最初仅可从其驻留设备访问，不继承导出进程设置的任何可访问性。导入进程需要启用来自任何计划访问内存的GPU的访问权限（使用cudaMemPoolSetAccess）。
同步API操作：与CUDA驱动程序集成的优化之一是与同步API的集成。当用户请求CUDA驱动程序同步时，驱动程序会等待异步工作的完成，并确定哪些释放操作保证已完成，这些释放的内存将被重新分配，而不考虑指定的流或禁用的分配策略。
附录内容：包括了关于cudaMemcpyAsync的当前上下文/设备敏感性、对已通过cudaFreeAsync释放的分配执行cuPointerGetAttribute查询导致的未定义行为，以及cuGraphAddMemsetNode不适用于通过流排序分配器分配的内存的情况，但这些分配的memset操作可以被捕获到流中。
cuPointerGetAttributes查询适用于流有序分配，这些分配与上下文无关。尽管可以查询CU_POINTER_ATTRIBUTE_CONTEXT属性，但返回值将是NULL，因为它们不与特定上下文关联。CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL属性可用于确定分配的位置，这对于使用cudaMemcpyPeerAsync进行p2h2p复制时选择上下文很有帮助。在CUDA 11.3中添加了CU_POINTER_ATTRIBUTE_MEMPOOL_HANDLE属性，有助于调试和确认分配来源池，特别是在执行IPC之前。
图内存节点允许图创建并拥有内存分配，具有GPU有序生命周期语义，定义了设备上何时允许访问内存，支持驱动管理的内存重用，并与创建图时可能捕获的流有序分配API（如cudaMallocAsync和cudaFreeAsync）的语义相匹配。图中的分配在整个图的生命周期（包括重复实例化和启动）中具有固定地址，便于直接引用而无需图更新。图内生命周期不重叠的分配可共享相同的物理内存。
支持与兼容性：图内存节点需要11.4版本或更高版本的CUDA驱动程序以及GPU上对流有序分配器的支持。通过代码片段检查特定设备是否支持此特性。
API基础：图内存节点作为表示内存分配或释放操作的图节点存在。通过cudaGraphAddMemAllocNode和cudaGraphAddMemFreeNode API显式创建。图分配在节点创建时分配虚拟地址，这些地址在分配节点的生命周期内固定，但分配内容在释放后不会持久且可能被覆盖。
流捕获：通过捕获对应的流有序分配和释放调用（如cudaMallocAsync和cudaFreeAsync），可以在图中创建图内存节点。捕获的分配API返回的虚拟地址可在图内的其他操作中使用，且由于捕获了流顺序依赖关系，图内存节点会与捕获的流操作正确排序。
图外访问与释放：图分配不必由创建图释放，未被图释放的分配在图执行结束后仍然存在，可通过后续CUDA操作访问或释放。可以通过事件、单一流或外部事件节点等机制确保正确的排序。
cudaGraphInstantiateFlagAutoFreeOnLaunch：该标志允许包含未释放内存分配的图被重新启动，自动插入异步释放操作。这简化了需在每次重新启动前手动释放所有图拥有的分配的算法，特别适用于单生产者多消费者场景，提高内存管理效率和算法的灵活性。
优化的内存重用：CUDA通过两种方式重用内存：一是基于图内存节点的自动管理机制，二是驱动管理的内存重用策略，进一步提升资源利用率和性能。
地址重用在图内：CUDA允许在图内通过为生命周期不重叠的不同分配赋予相同的虚拟地址范围来重用内存。这意味着指向不同生命周期的分配的指针可能不是唯一的。
物理内存管理和共享：CUDA负责在GPU顺序中的分配节点到达前将物理内存映射到虚拟地址。作为一种减少内存占用和映射开销的优化，如果多个图不会同时运行，它们可以使用同一块物理内存进行不同的分配。但是，如果物理页面绑定到一个或多个正在执行的图，或者绑定到尚未释放的图分配，则不能重用这些页面。CUDA可以在图实例化、启动或执行的任何时候更新物理内存映射，并可能在未来的图启动之间引入同步以防止活动图分配引用相同物理内存的问题。
性能考虑：当多个图在同一流中启动时，CUDA尝试为它们分配相同的物理内存，因为这些图的执行不能重叠。作为优化措施，图的物理映射会在启动之间保留，以避免重新映射的成本。如果之后其中一个图在可能与其他图重叠执行的流中启动（例如，如果它在不同的流中启动），则CUDA必须执行一些重映射操作，因为并发图需要不同的内存以避免数据损坏。
首次启动与cudaGraphUpload：物理内存不能在图实例化期间分配或映射，因为未知该图将在哪个流中执行。映射在图启动时完成。调用cudaGraphUpload可以将分配成本与启动分离，立即执行该图的所有映射，并将图与上传流关联。如果该图随后在同一流中启动，则无需额外的重映射。
物理内存占用：异步分配的池管理行为意味着销毁包含内存节点的图（即使其分配已释放）不会立即将物理内存返回给操作系统供其他进程使用。要明确地将内存返回给操作系统，应用程序应使用cudaDeviceGraphMemTrim API。
对等访问：图分配可以配置为允许多个GPU访问，CUDA会根据需要将分配映射到对等GPU上。CUDA允许需要不同映射的图分配重用相同的虚拟地址。这意呀着一个分配有时可能允许比创建时请求的更多的对等访问，但依赖于这些额外映射仍然是错误的。
数学函数：设备代码支持的C/C++标准库数学函数以及所有仅在设备代码中支持的内部函数的列表、描述及部分函数的精度信息被提供。设备代码中的数学函数不设置全局errno变量，也不报告任何浮点异常以指示错误，因此需要错误诊断机制的用户应实现额外的输入和输出筛查。用户负责指针参数的有效性，且不应向数学函数传递未初始化的参数，因为这可能导致未定义行为。

标准函数误差界限：在设备上执行的标准函数具有特定的错误界限，这些界限基于广泛测试得出但非绝对保证。单精度浮点运算的加法和乘法遵循IEEE标准，最大误差为0.5 ulp。推荐使用rintf()而非roundf()进行单精度浮点数到整数的转换，因为前者映射为单个指令，效率更高。同理，在双精度操作中，推荐使用rint()而不是round()。
内建函数（Intrinsic Functions）：仅能在设备代码中使用的内建函数提供了某些标准函数的更快但精度较低的版本，通过前缀__标识（如__sinf(x)）。它们通过映射到较少的原生指令实现加速，并可通过编译器选项-use_fast_math全局启用以牺牲精度换取速度。单精度内建函数如__fadd_[]和__fmul_[]不参与FMAD（融合乘加）操作，而直接加法和乘法运算符可能被编译器合并为FMAD。
C++语言支持：CUDA源文件支持混合主机和设备代码，前端编译器旨在模拟ISO C++标准行为，并扩展了CUDA特定的构造。支持从C++11到C++20的特性，但存在特定限制，如不支持主机编译器特有扩展、__Complex类型仅限主机代码、__int128和__float128类型的使用条件等。同时，对变量声明的内存空间指定符（如__device__、__shared__）应用了特定规则。
针对GPU优化，了解和利用这些信息可以指导开发者在精度和性能之间做出权衡，特别是在处理数学运算、内存管理及编译配置方面。

Volatile Qualifier: 使用volatile关键字可以防止编译器对全局或共享内存的读写操作进行优化，确保任何时候访问这些变量都会实际执行内存读写指令。这对于需要精确控制内存访问时机和顺序的优化场景特别重要。
Pointers: 指针使用上的限制强调了主机代码与设备代码间访问内存的不同行为，错误使用会导致未定义行为，如段错误。正确管理指针访问范围是优化数据传输和内存访问效率的关键。
Operators and Assignment: __constant__变量只能从主机端通过运行时函数赋值；__shared__变量声明时不能初始化；内置变量不能被赋值或取地址，理解这些限制有助于避免不必要的内存访问和优化内存使用模式。
Run Time Type Information (RTTI), Exception Handling, 和 Standard Library 在设备代码中不受支持，提醒开发者在GPU代码中不能依赖这些特性，需寻找替代方案以保证性能和兼容性。
Namespace Reservations: 规定了一系列保留的命名空间，开发者应避免在这类命名空间中添加声明或定义，以免引发未定义行为。
Functions: 强调了外部链接、隐式声明及默认函数的执行空间确定规则，以及__global__函数参数传递的限制（如大小限制、不支持变参、不支持引用传递等），这些知识对于编写高效且兼容的内核函数至关重要。
Function Parameters and Argument Processing: 特别指出__global__函数参数处理的特殊性，包括不同启动环境下的复制行为，以及对大于4KB参数的支持要求特定版本的工具包和驱动，这对设计大型数据结构传递给内核函数时的优化策略有直接影响。
Static Variables within Function 和 Function Pointers 的规则指导了如何在函数内部有效使用静态变量和函数指针，特别是它们在不同执行空间的适用性和限制。
Function Recursion, Friend Functions, Operator Function, 以及 Allocation and Deallocation Functions 的限制，提醒开发者在设计复杂类和运算符重载时，需要遵循特定规则，以避免不可预知的行为或性能瓶颈。

####在GPU优化的上下文中，以下几点是关键知识点：

虚拟函数的匹配规则：当派生类中的函数覆盖基类的虚拟函数时，覆盖函数和被覆盖函数的执行空间说明符（如__host__、__device__）必须一致。
传递对象至全局函数限制：不允许将具有虚拟函数的类的对象作为参数传递给__global__函数。此外，如果对象在主机代码中创建，那么在设备代码中调用其虚拟函数会导致未定义行为；反之亦然。
虚拟基类的使用限制：同样地，不支持将从虚拟基类派生的类的对象作为参数传递给__global__函数。
匿名联合的限制：命名空间范围内的匿名联合的成员变量不能在__global__或__device__函数中引用。
Windows平台特定约束：CUDA编译器遵循的类布局与Microsoft主机编译器不同，可能导致类型C在CUDA和Microsoft编译器下计算的布局和大小有差异。跨主机和设备代码传递类型C的对象会导致未定义行为。
模板使用的限制：在__global__函数模板实例化或__device__/__constant__变量实例化中，类型或模板若不满足特定条件则不能被使用。
长双精度类型：设备代码中不支持使用long double类型。
废弃注解：支持使用[[deprecated]]等注解来标记废弃的实体，并在特定条件下生成废弃诊断信息。可以使用特定的编译器标志或Pragma来抑制这些警告或将其转化为错误。
noreturn注解：支持在主机和设备代码中使用noreturn属性来标记不会返回的函数。
【likely】/【unlikely】标准属性：这些属性可用于指导设备编译器优化器关于代码路径执行可能性的提示。
const和pure GNU属性：支持对设备函数使用这些属性以向优化器提供函数纯净度的额外信息。
Intel主机编译器特定：使用Intel编译器作为主机编译器时，需要通过宏定义来启用某些内建函数的支持。
C++11特性支持：nvcc支持默认由主机编译器启用的C++11特性，并且可以通过指定-std=c++11标志来进一步开启所有C++11特性。
这些知识点对于进行GPU编程和优化至关重要，特别是在涉及跨平台编译、代码纯净性优化、以及确保代码兼容性和可移植性方面。
Lambda表达式执行空间指定：编译器根据包含lambda表达式的最小作用域（块、类或命名空间）自动推导闭包类成员函数的执行空间。若无包围函数作用域，默认为__host__。闭包类型不能在__global__函数模板实例化中的类型或非类型参数中使用，除非lambda定义于__device__或__global__函数内。
std::initializer_list：默认情况下，CUDA编译器将std::initializer_list的成员函数视为具有__host__ __device__执行空间，允许直接从设备代码调用。使用--no-host-device-initializer-list标志会将其限制为仅__host__。
Rvalue引用：类似地，std::move和std::forward默认为__host__ __device__，允许设备代码直接调用。使用--no-host-device-move-forward标志则限制为__host__。
Constexpr函数与模板：默认情况下，不同执行空间的函数间不能调用constexpr函数。实验性标志--expt-relaxed-constexpr放松此限制，允许跨执行空间调用。编译器定义宏__CUDACC_RELAXED_CONSTEXPR__来标识该标志已启用。
Constexpr变量：未标注执行空间（如__device__）的constexpr变量被视为主机变量。标量类型的constexpr变量值可直接用于设备代码，非标量类型的常量元素可在constexpr函数内部使用，但设备代码不能直接引用或取地址。
Inline命名空间：如果CUDA翻译单元包含特定实体的定义，且这些实体在inline命名空间内，可能需要使用唯一名称以避免编译时的命名冲突。
Inline无名命名空间：限制某些实体声明于inline无名命名空间中。
thread_local存储类别：不支持在设备代码中使用。
__global__函数与模板：限制包括lambda表达式作为模板参数、不能声明为constexpr、不能有std::initializer_list或va_list类型的参数、不能有右值引用类型参数以及对变参模板的额外限制。
__managed__和__shared__变量：不能标记为constexpr。
默认构造函数的执行空间指定：首次声明时显式默认的函数忽略执行空间指定符，而后续的显式默认不会被忽略。
C++14至C++20特性支持：概述了对C++14至C++20各版本特性的支持与限制，包括函数返回类型推导、变量模板、内联变量、结构化绑定等，并指出模块和协程在CUDA设备代码中不受支持。
三向比较运算符：支持在主机和设备代码中使用，但某些用法依赖于主机实现提供的标准模板库功能。为了消除警告并使功能满足设备代码的要求，可能需要指定--expt-relaxed-constexpr标志。
consteval函数：允许跨执行空间调用，当被调用的函数被声明为consteval时，不触发编译器诊断。例如，__device__或__global__函数可以调用__host__consteval函数，反之亦然。
多态函数包装器(nvstd::function)：nvfunctional头文件提供了一个类模板，用于存储、复制和调用任何可调用目标，如lambda表达式。它可以在主机和设备代码中使用，但有特定初始化限制，如不能跨执行空间传递实例，且不能用于__global__函数的参数类型中。
扩展Lambda：通过--extended-lambda标志，允许在lambda表达式中明确指定执行空间注解。这包括__device__和__host__ __device__的扩展lambda，并且可以在__global__函数模板实例化时作为类型参数使用。这些扩展lambda有类型特征检测编译时宏以及一系列使用限制。
this指针捕获：对于在非静态类成员函数内部定义的lambda，如果lambda体引用了类成员变量，C++11/14要求隐式捕获this指针。CUDA编译器支持C++17引入的“*this”捕获模式，允许在__device__和__global__函数中的lambda以及特定条件下主机代码中的扩展__device__ lambda直接捕获对象副本，解决潜在的运行时错误问题。

纹理获取(Texture Fetching)
这一节介绍了根据纹理对象的各种属性（如纹理内存和表面内存中的描述）计算纹理函数返回值的公式。纹理通过非规范化或规范化纹理坐标进行获取，超出范围的坐标会根据寻址模式映射到有效范围内。
最近点采样(Nearest-Point Sampling)
在此过滤模式下，纹理获取返回的值基于最近的整数坐标计算。适用于整数纹理时，返回值可选地映射到[0.0, 1.0]区间。
线性过滤(Linear Filtering)
仅对浮点纹理可用，通过计算周围像素的加权平均值来实现。权重存储为9位定点格式，具有8位小数部分。
表查找(Table Lookup)
利用纹理过滤实现表查找，确保TL(0)对应纹理的第一个元素，TL®对应最后一个元素。
计算能力(Compute Capabilities)
计算设备的规格和特性由其计算能力决定。不同计算能力级别对应不同的架构特征和技术规格，包括对IEEE浮点标准的遵循情况及特定功能的引入与可用性。例如，新功能通常会在后续架构中保留，但某些高度专业化的功能可能不保证在所有后续计算能力中都可用。
全局内存访问(Global Memory Accesses)
对于计算能力5.x的设备，全局内存访问默认缓存在L2中，只读数据可通过__ldg()函数缓存在统一的L1/纹理缓存中。非只读数据在某些条件下可通过特定机制启用L1/纹理缓存。
以上内容涉及GPU编程中的纹理操作、数据获取方式、计算能力差异以及内存访问优化策略，对于进行GPU优化工作具有参考价值。

共享内存(Shared Memory)

共享内存有32个bank，每个连续的32位字映射到连续的bank中。每个bank每时钟周期有32位的带宽。
即使两个地址落在同一个bank中，一个warp的共享内存请求在访问同一32位字内的任何地址时不会产生bank冲突。读取访问时，该字会广播给请求的线程；写入访问时，由其中一个线程（哪个线程执行写入是未定义的）进行。
计算能力6.x及7.x的设备中，共享内存的行为与计算能力5.x设备相似，但7.x设备中的共享内存可从统一数据缓存中配置不同大小。

计算能力6.x架构

流式多处理器(SM)包含：多个调度器、统一的L1/纹理缓存、纹理单元等。
统一的L1/纹理缓存也服务于纹理单元，实现各种寻址模式和数据过滤。所有SM共享L2缓存，用于缓存对局部或全局内存的访问，包括临时寄存器溢出。
L1/L2缓存的行为可以通过加载指令的修饰符按访问基础部分配置。

计算能力7.x架构

引入了独立线程调度，允许warp内的线程进行以前不可用的同步模式，简化了从CPU代码移植时的代码更改，但也可能因开发者对先前硬件架构的warp同步性假设而导致执行代码的线程集与预期不同。
提供了新的内联函数来支持独立线程调度，并给出了迁移至Volta安全代码的建议修正措施。
共享内存可配置为不同的大小，由统一数据缓存划分，剩余部分作为L1缓存使用。
允许单个线程块寻址全部共享内存容量，具体数值依据具体架构而定，并且需要动态共享内存分配和显式设置属性以支持超过特定阈值的共享内存使用。

全局内存(Global Memory)

在计算能力5.x、6.x、7.x的设备上，全局内存的行为保持一致。

流式多处理器（SM）结构：SM由多个调度器组成，静态地将其线程束分配给这些调度器。每个调度周期，每个调度器为其分配好的、准备好执行的线程束发出一条指令。
统一数据缓存与共享内存：在NVIDIA Ampere架构中，统一数据缓存大小为192 KB（计算能力8.0和8.7）或128 KB（计算能力8.6和8.9），其中一部分可配置为共享内存，大小可在一定范围内动态调整。应用程序通过cudaFuncSetAttribute()设置共享内存预留量（carveout）。不同计算能力的设备允许的最大共享内存容量有所不同，且部分保留用于系统使用。
全局内存：全局内存的行为与计算能力5.x的设备一致，未提供详细变化信息。
计算能力9.0的特性：
- 在NVIDIA H100 Tensor Core GPU架构中，统一数据缓存大小为256 KB，共享内存容量更加灵活，最大可达228 KB。
- 强烈推荐通过CUDA-X库（如cuBLAS、cuDNN、cuFFT）或CUTLASS模板库来利用复杂的矩阵乘积累（MMA）加速功能，这些功能仅通过内联PTX在CUDA编译工具链中可用。
驱动API：驱动API是CUDA的底层接口，需要初始化并通过cuInit()开始使用。应用程序应加载PTX代码而非二进制代码以确保对未来架构的兼容性。驱动API支持对CUDA上下文、模块和内核执行等对象的细粒度控制。
综上，关键点涉及GPU架构中的SM调度、共享内存配置与优化、全局内存访问以及针对特定计算任务（如矩阵运算）的硬件加速特性，还包括了使用CUDA驱动API进行程序开发时的初始化、对象管理和代码加载的最佳实践。
####CUDA上下文类似于CPU进程，管理着所有资源及执行的操作，并在销毁时自动清理这些资源。每个上下文有独立的地址空间，不同的CUdeviceptr值引用不同的内存位置。主机线程同一时间只能有一个设备上下文当前有效，通过cuCtxCreate()创建并激活上下文。若没有有效的上下文绑定到线程，大多数CUDA函数会返回CUDA_ERROR_INVALID_CONTEXT错误。主机线程维护一个当前上下文的栈，cuCtxCreate()将新上下文压入栈顶，cuCtxPopCurrent()可使上下文脱离当前线程并恢复前一个上下文。上下文有使用计数，cuCtxCreate()创建时为1，cuCtxAttach()增加，cuCtxDetach()减少，计数归零时上下文被销毁。
模块是动态加载的设备代码和数据包，由nvcc输出，类似Windows的DLL，其中所有符号（函数、全局变量等）在模块范围内命名，便于第三方编写的模块在同一CUDA上下文中交互。通过cuLaunchKernel()函数可以配置并启动内核执行，参数传递可通过指针数组或额外选项进行。参数对齐要求需匹配设备代码中的类型对齐规则。
运行时API与驱动API之间存在互操作性，应用可以混合使用两者。驱动API创建的上下文会被后续的运行时调用所继承，反之亦然，运行时初始化隐式创建的上下文也可由驱动API管理。设备内存可在任一API下分配和释放，CUdeviceptr与常规指针可互相转换，这使得基于驱动API的应用能够调用基于运行时API的库（如cuFFT, cuBLAS等）。
此外，Driver Entry Point Access APIs允许从CUDA 11.3开始，通过函数指针的方式调用驱动API函数，为用户提供类似dlsym或GetProcAddress的功能，以动态获取CUDA驱动函数的地址。

17.5.2. 驱动函数类型定义

CUDA工具包通过提供访问包含所有CUDA驱动API函数指针定义的头文件，帮助检索CUDA驱动API入口点。这些头文件安装在CUDA工具包中，并在包含目录下可用。表中总结了包含每个CUDA API头文件typedefs的文件。这些头文件不定义实际的函数指针，而是为函数指针定义typedef。例如，cudaTypedefs.h为cuMemAlloc驱动API定义了typedef。CUDA驱动符号采用基于版本的命名方案，除第一个版本外，名称中带有_v*扩展名。当特定CUDA驱动API的签名或语义发生变化时，相应的驱动符号版本号会增加。

17.5.3. 驱动函数获取

利用驱动程序入口点访问API和适当的typedef，可以获取任何CUDA驱动API的函数指针。使用方法包括直接调用驱动API或运行时API来根据CUDA版本获取兼容的ABI版本的函数地址。

17.5.3.1. 使用驱动API

驱动API需要CUDA版本作为参数以获取请求的驱动符号的ABI兼容版本。每个函数有特定的ABI，通过_v*扩展名表示。为了可移植性，指定高于所需版本的CUDA版本来获取特定版本的驱动API地址可能不可行。

17.5.3.2. 使用运行时API

cudaGetDriverEntryPoint使用CUDA运行时版本获取兼容的ABI版本。而cudaGetDriverEntryPointByVersion允许用户更具体地控制所请求的ABI版本。

17.5.3.3. 获取每线程默认流版本

某些CUDA驱动API支持配置为默认流或每线程默认流语义，可以通过特定后缀（如_ptsz或_ptds）的名称获取。

17.5.3.4. 访问新的CUDA特性

即使没有最新CUDA工具包，也可以通过更新CUDA驱动来使用新功能，通过API动态获取函数指针实现。

17.5.4. cuGetProcAddress的潜在影响

讨论了使用cuGetProcAddress与隐式链接、编译时与运行时版本使用差异等场景下的潜在问题，强调了不同CUDA版本和API版本间的行为差异及可能的未定义行为风险，特别是在处理API/ABI兼容性时。

API版本更新与兼容性问题：在CUDA的不同版本中，API的修改（如从CUDA 11.4到11.6）可能导致应用编译时使用的typedef与实际运行时获取的函数指针不匹配，从而引发未定义行为。即使进行了显式的版本检查，也无法安全地涵盖同一CUDA主版本内的次版本升级所带来的兼容性问题。
运行时API使用的问题：当应用程序动态链接到不同版本的CUDA驱动时，运行时API（如cudaApiGetDriverEntryPoint）可能返回与应用程序typedef不匹配的函数指针版本，特别是在使用较新Runtime和较旧驱动的组合下，这可能导致类型不一致和潜在的ABI不兼容。
动态版本控制带来的复杂性：不同编译、运行时版本与驱动版本的组合可能导致函数指针与实际使用的API版本不匹配，增加了开发和维护的复杂度，尤其是在API签名变更时，可能会导致ABI不兼容和未定义行为。
API/ABI不匹配的含义：API不匹配不仅限于功能上的差异，还可能涉及到ABI（应用二进制接口）的变化，如cuCtxCreate从_v2到_v3的变化引入了新的参数，如果typedef与实际函数指针不匹配，将导致ABI不兼容和程序行为未定义。
cuGetProcAddress失败原因分析：cuGetProcAddress可能因API使用错误或无法找到请求的驱动API而失败。通过返回的错误码和查询结果，开发者可以区分是由于版本不足还是符号找不到等原因导致的失败。
统一内存编程：CUDA统一内存为CPU和GPU提供了透明的数据访问，简化了数据管理和迁移，提高了编程效率。虽然数据移动仍然存在，但可以通过提示进行性能优化，且统一内存的物理位置对程序透明，保证了访问的正确性和一致性。统一内存可通过显式分配或通过库（如CuMemPool）获得，并要求系统支持特定级别的地址转换服务。
系统需求与查询支持等级：使用统一内存需要系统支持特定级别的地址转换服务，应用程序可通过查询设备属性来确定系统对统一内存的支持等级，以确保代码的可移植性和正确运行。
综上所述，GPU优化需关注API兼容性、运行时与驱动版本的匹配、理解动态版本控制的复杂性、处理API/ABI不匹配问题、有效分析函数查询失败原因，并利用统一内存特性简化数据管理，同时考虑系统支持的需求。
####CUDA统一内存简化了GPU编程模型，不再需要主机和设备之间的独立分配及显式内存传输。程序可通过以下方式分配统一内存：

系统分配内存：包括malloc()、mmap()、C++ new操作符分配的内存，以及CPU线程栈、线程局部、全局变量等。首次访问时（即“首次触摸”），物理内存按访问线程所在的处理器分配，如GPU线程访问则分配GPU物理内存，CPU线程访问则分配NUMA节点的CPU物理内存。
CUDA托管内存：通过cudaMallocManaged()分配，该API允许CPU和GPU并发访问，且不需要手动复制数据。托管内存也可使用cudaFree()释放。托管变量（__managed__）简化了全局变量的处理，自动在主机和设备间共享数据。
性能提示与优化：
数据预取：cudaMemPrefetchAsync可异步迁移数据至指定处理器附近，提高访问速度。
数据使用提示：cudaMemAdvise提供数据访问模式的提示，如常读、首选位置等，帮助CUDA做出更好的性能决策。
查询内存使用属性：可以查询托管内存区段的属性，如是否设置为常读或首选位置，以进一步优化数据布局和访问模式。
关键点总结：
统一内存自动管理主机与设备内存的交互，减少程序员的手动操作。
托管内存通过cudaMallocManaged自动在CPU与GPU间迁移数据，提升编程效率。
性能提示API（如数据预取和使用建议）为开发者提供了优化内存访问性能的手段，虽不影响程序语义，但需根据实际情况谨慎使用以避免负面影响。
__managed__关键字的引入简化了全局变量的跨平台共享，但需注意与CUDA上下文的关联及特定约束。
####### 系统分配内存深入示例与GPU优化
本文节选了关于系统分配内存(System-Allocated Memory)的深入示例及其在CUDA统一内存环境下的应用，特别关注于GPU优化的相关知识点。

主要内容概要：

系统分配内存示例：通过一个打印字符数组前8个字符到标准输出的内核函数，展示了多种调用方式，包括如何访问文件作用域或全局作用域变量。强调了栈变量、文件作用域和全局作用域变量必须通过指针由GPU访问，并讨论了全局变量直接访问与间接访问（通过指针）的区别及原因。
文件支持的统一内存：说明了在具有完整CUDA统一内存支持的系统中，设备可以直接访问主机进程拥有的任何内存，包括由物理文件支持的内存。此部分展示了如何修改示例以直接从输入文件读取并打印字符串到GPU，同时指出在某些系统上对文件支持内存的原子访问不被支持。
统一内存的跨进程通信（IPC）：探讨了在管理单个GPU每进程的应用场景中，使用统一内存进行超订阅和多GPU访问的需求。虽然CUDA IPC不直接支持托管内存共享，但在支持统一内存的系统中，系统分配的内存可以实现跨进程通信。提到了在Linux下创建IPC能力系统分配内存的各种方法，并指出不同主机间无法使用此技术共享内存。
性能调优：
- 强调了为实现统一内存的良好性能，理解内存分页、页面大小及虚拟地址空间的重要性。
- 建议遵循统一内存性能调整提示，但注意不当使用可能导致性能下降。
- 讨论了内存分页和页面大小对性能的影响，指出小页面尺寸减少内存碎片但增加TLB（转换旁路缓冲区）未命中，大页面尺寸则相反，且GPU上的TLB未命中成本显著高于CPU。选择合适的页面大小是性能调优的关键因素之一，需平衡内存碎片、TLB未命中率和内存迁移成本。

统一内存访问：展示了如何在不同的上下文中访问统一内存，包括全局变量、文件支持内存及跨进程通信，强调了通过指针访问的重要性。
性能考虑：详细介绍了内存分页和页面大小对GPU性能的影响，指出适当调整页面大小以减少TLB未命中，特别是针对GPU的优化，因为GPU上的TLB未命中成本更高。
跨进程共享：说明了在支持的系统中，系统分配内存能够用于跨进程通信，拓展了统一内存的应用范围，尽管存在一些限制如不同主机间的共享限制。
综上所述，这些内容为GPU优化工程师提供了关于在CUDA环境下利用统一内存进行高效编程和性能调优的深入见解。
####### GPU与CPU页面表：硬件一致性与软件一致性
NVIDIA Grace Hopper等硬件一致性系统为CPU和GPU提供了逻辑上统一的页面表。这对GPU访问系统分配内存至关重要，因为GPU使用CPU为请求的内存创建的页面表项。如果该页面表项使用CPU默认的4KiB或64KiB页面大小，则对大型虚拟内存区域的访问会导致严重的TLB未命中，从而显著降低性能。通过配置大页面可以确保系统分配的内存使用足够大的页面大小以避免此类问题。
在CPU和GPU各自拥有独立逻辑页面表的系统中，则需考虑不同的性能调优方面：为保证一致性，这些系统通常在处理器访问映射到其他处理器物理内存的地址时使用页面错误。硬件一致性系统在CPU和GPU线程频繁并发访问同一内存页的情况下，相比软件一致性系统提供了显著的性能优势。

主机直接访问统一内存

某些设备具有硬件支持，可直接从主机对GPU驻留的统一内存进行一致性的读取、存储和原子访问。这些设备的cudaDevAttrDirectManagedMemAccessFromHost属性设置为1。硬件一致性系统中NVLink连接的设备都具备此属性。在这些系统上，主机可以直接访问GPU驻留内存，无需页面错误和数据迁移。

主机本地原子操作

包括硬件一致性系统中NVLink连接设备在内的某些设备支持对CPU驻留内存的硬件加速原子访问，这意味着对主机内存的原子访问不需要通过页面错误来模拟。对于这些设备，cudaDevAttrHostNativeAtomicSupported属性设置为1。

在不完全支持CUDA统一内存的设备上的统一内存

计算能力6.x设备上的统一内存

具有计算能力6.x但无分页内存访问能力的设备完全支持CUDA托管内存，并且是协调的。编程模型和性能调优与完全支持CUDA统一内存的设备类似，但系统分配器不能用于分配内存。

计算能力5.x设备或Windows平台上的统一内存

计算能力低于6.0的设备或Windows平台支持CUDA托管内存v1.0，具有有限的数据迁移、一致性和内存超分配支持。这些平台上的托管内存使用和优化有其特定细节，包括数据迁移与一致性限制、GPU内存超分配限制以及多GPU交互的带宽和兼容性问题。

结论

文章重点介绍了GPU优化中的关键概念，包括硬件一致性页面表的优势、直接从主机访问GPU统一内存的能力、主机本地原子操作的支持，以及不同计算能力设备上统一内存的特性和限制。这些知识点对于开发者理解和优化涉及CPU与GPU协同工作的应用程序性能至关重要。
####在使用预6.x架构的GPU时，为了确保统一内存编程模型的连贯性，对CPU和GPU同时执行时的数据访问进行了限制。实际上，只要任何内核操作正在执行，GPU就独占所有托管数据的访问权，不论该内核是否实际使用这些数据。当使用cudaMemcpy*()或cudaMemset*()访问托管数据时，系统可能会选择从主机或设备访问源或目标，这将限制在执行cudaMemcpy*()或cudaMemset*()期间CPU对该数据的并发访问。对于具有concurrentManagedAccess属性设置为0的设备，在GPU活跃期间，CPU不允许访问任何托管分配或变量，即使CPU和GPU访问不同的托管内存分配也会导致段错误。
为了避免冲突，程序必须在访问托管数据（如变量y）之前与GPU进行显式同步。在6.x之前的架构上，无论GPU内核是否实际触碰同一数据，CPU线程在执行内核启动和后续同步调用之间均不得访问任何托管数据。逻辑GPU活动用于判断GPU是否空闲，因此即使内核运行迅速并在CPU访问数据前完成，也需要进行显式同步。
有效的同步方式包括cudaDeviceSynchronize()、确保GPU完成工作的流同步函数（如cudaStreamSynchronize()在指定流是唯一仍在GPU上执行的流时），以及某些情况下的事件同步函数。依赖关系可以在流或事件之间创建，以推断通过同步流或事件的其他流的完成情况。在特定条件下，CPU可以在流回调中访问托管数据，但必须确保没有其他可能访问托管数据的流在GPU上活跃。
通过cudaStreamAttachMemAsync()函数，可以将托管内存与CUDA流关联起来，从而实现更细粒度的控制，允许基于程序特定数据访问模式的并发。这样，程序员可以指示基于内核是否在指定流中启动来使用数据。如果分配未与特定流关联，则默认对所有运行的内核可见。通过这种方式，可以在保持数据一致性的同时，增强CPU与GPU之间的并行性，并可能实现统一内存系统中的数据传输优化。
示例说明了如何通过流关联来控制数据可见性，以及在低版本计算能力设备上必须注意哪些数据对哪些流可见，以避免访问冲突和未定义行为。正确管理流与数据的关联对于实现安全且高效的CPU+GPU并发至关重要。

Stream Attach与多线程主机程序： cudaStreamAttachMemAsync()函数主要用于在CPU线程中实现独立的任务并行。每个CPU线程通常会为它产生的所有工作创建自己的流，以避免使用CUDA的NULL流时在不同线程间产生依赖性。此功能通过将线程的托管分配与其自身的流关联起来，减少多线程程序中托管数据全局可见性可能引起的交互问题。
模块化程序与数据访问限制的高级主题： 使用cudaMallocManaged()时指定cudaMemAttachHost标志可以创建初始时对设备端执行不可见的分配，确保数据在分配和被特定流获取之间不会意外与其他线程执行产生交互。这有助于线程间安全独立操作，避免新分配的内存因其他线程启动的内核正在运行而被视为在GPU上使用，从而影响CPU侧的即时访问。
与流关联的统一内存的Memcpy()/Memset()行为： 在并发管理访问未设置的设备上，统一内存的Memcpy*/Memset*操作遵循特定规则，确保数据在满足一定条件时从主机或设备访问，并强调在GPU操作期间任何与该流相关联或具有全局可见性的数据的CPU访问可能会导致段错误，除非适当地同步以确保操作已完成。
懒加载（Lazy Loading）： 懒加载推迟了CUDA模块和内核从程序初始化到接近内核执行时的加载时间，减少未使用的内核加载，降低初始化时间和内存开销。通过设置环境变量CUDA_MODULE_LOADING为LAZY启用。此特性针对CUDA运行时用户，且在CUDA 11.7及更高版本中提供，包括延迟模块加载和延迟内核加载直到首次使用。要求应用程序使用11.7+的CUDA运行时，并且对用户透明，前提是遵循CUDA编程模型。
懒加载版本支持： 需要R515及以上版本的用户模式库以及CUDA 11.7+工具包来利用懒加载特性。对于编译器，懒加载不需特别支持，但需要使用11.7+的CUDA运行时以享受其益处。触发内核加载自动进行，但也可以通过API更细粒度地控制加载时机，如CUDA驱动API中的cuModuleGetFunction()调用。

20.3.2. CUDA Runtime API管理

CUDA Runtime API自动处理模块管理，推荐使用cudaFuncGetAttributes()来引用内核，确保内核加载时状态不变。

20.4. 检查懒加载是否开启

通过CUresult cuModuleGetLoadingMode ( CUmoduleLoadingMode* mode )检查懒加载是否启用。注意，运行此函数前必须初始化CUDA。

20.5. 懒加载可能遇到的问题

懒加载设计上不需应用程序修改即可使用，但存在一些注意事项，特别是应用未完全遵循CUDA编程模型时。

并发执行：加载内核可能导致上下文同步，若程序错误地将内核并发执行视为保证，且一个内核依赖另一个内核执行以返回，可能导致死锁。
分配器：懒加载延后代码加载至执行阶段，此过程需要内存分配。若应用启动时尝试分配全部VRAM（如自定义分配器），可能导致无足够内存加载内核。
自动调优：部分应用通过多次启动实现相同功能的内核以确定最快版本。懒加载下，包括内核加载时间会影响结果准确性。

21. 扩展GPU内存(EGM)

EGM利用NVLink-C2C提高单一节点系统中GPU对系统内存的访问效率。适用于集成CPU-GPU的NVIDIA系统，允许任何GPU线程访问分配的物理内存。EGM确保所有GPU能以GPU-GPU NVLink或NVLink-C2C速度访问资源，通过NVSwitch结构，GPU线程可访问包括CPU附加内存和HBM3在内的所有内存资源。

21.1. EGM预备知识

支持的拓扑：单节点单GPU、单节点多GPU、多节点单GPU。
Socket标识符：EGM利用操作系统分配的NUMA节点标识符，与设备序号不同，与最近的主机节点关联，可通过cuDeviceGetAttribute获取。
分配器与EGM支持：映射系统内存为EGM不会导致性能问题，访问远程Socket系统内存更快，因为EGM确保流量通过NVLink路由。cuMemCreate和cudaMemPoolCreate支持EGM。
内存管理API扩展：EGM内存可通过虚拟内存或流有序内存分配器映射，新增属性类型以识别NUMA节点位置。

EGM接口使用

单节点单GPU：现有CUDA主机分配器及系统分配内存均可用于高带宽C2C访问。
单节点多GPU：用户需提供放置信息，通过NUMA节点ID表达，使用VMM API或CUDA内存池分配管理EGM内存。
以上内容涵盖了GPU优化中的CUDA Runtime API使用、懒加载机制的验证、潜在问题分析，以及EGM技术在扩展内存访问和提高数据传输效率方面的应用详情。
####### 使用CUDA内存池进行GPU优化
在实现显存管理（EGM）时，用户可以通过创建节点上的内存池并授权给对等节点来使用。关键点是明确指定cudaMemLocationTypeHostNuma作为位置类型及numaId作为位置标识符。示例代码展示了如何使用cudaMemPoolCreate创建内存池。
对于直接连接的对等访问，可利用cudaMemPoolSetAccessAPI来设置访问权限。提供了一个示例代码片段，说明了如何为访问设备配置权限。
内存池创建并分配访问权限后，用户可以将其设置到驻留设备上，并开始使用cudaMallocAsync异步分配内存。

多节点单GPU场景下的远程对等访问

除了内存分配外，远程对等访问在EGM方面没有特定改动，遵循CUDA跨进程（IPC）协议。用户应通过cuMemCreate分配内存，同样需要明确指定CU_MEM_LOCATION_TYPE_HOST_NUMA作为位置类型和numaID作为位置标识符。此外，还需定义CU_MEM_HANDLE_TYPE_FABRIC作为请求的句柄类型。
代码示例展示了如何在节点A上分配物理内存，并通过cuMemExportToShareableHandle导出该句柄至其他节点（节点B）。在节点B上，使用cuMemImportFromShareableHandle导入句柄，并像处理其他Fabric句柄一样对待它。接着，用户可以在节点B上预留地址空间并本地映射内存。最后，为节点B上的每个本地GPU赋予适当的读写权限。