CUDA计算能力与GPU兼容性指南

基本概念：CUDA计算能力 (Compute Capability) 概述

**CUDA计算能力（Compute Capability, CC）**是NVIDIA用于标识GPU硬件特性的版本编号。每款GPU都有一个CC版本号（如6.0、7.5、8.6等），由主版本和次版本组成，用于表示该GPU架构所支持的硬件功能、指令集、内存模型等。较新的计算能力通常引入更多硬件功能、性能改进和新的指令, 使开发者能够利用最新GPU架构的优势。反之，如果使用较低的计算能力生成代码，那么兼容的GPU范围更广，但无法利用新架构的最新特性。因此，理解计算能力对于确保应用兼容不同GPU并充分发挥性能非常重要。

**计算能力、GPU架构代号与产品名称的关系：**NVIDIA常以架构代号（如Pascal、Turing、Ampere等）来指代一代GPU微架构，每代架构对应一系列计算能力版本。商业GPU产品（如RTX 2080、RTX 4090、Tesla/Quadro系列等）会基于某一架构并具有相应的CC。下表列出计算能力6.x至9.x的主要架构代号及示例GPU，以便清晰对比三者关系：

计算能力 (CC)	架构代号	示例GPU型号
6.0	Pascal (旗舰)	Tesla P100 (服务器级, 6.0)
6.1	Pascal (消费)	GeForce GTX 1080 (6.1)
6.2	Pascal (嵌入)	Jetson TX2 (Tegra, 6.2)
7.0	Volta	Tesla V100 (7.0)
7.5	Turing	GeForce RTX 2080 Ti (7.5)
8.0	Ampere (服务器)	NVIDIA A100 (8.0)
8.6	Ampere (消费)	GeForce RTX 3090 (8.6)
8.7	Ampere (嵌入)	Jetson AGX Orin (8.7) (支持推理)
8.9	Ada Lovelace	GeForce RTX 4090 (8.9)
9.0	Hopper	NVIDIA H100 (9.0)

表：计算能力版本与GPU架构及示例产品的对应关系。 可以看到，同一架构可能支持多个CC次版本号。例如Ampere架构包括数据中心GPU A100（CC 8.0）和消费级RTX30系列（CC 8.6、8.7等）。一般来说，主版本号不同意味着全新的架构系列，例如Turing=7.x，Ampere=8.x，Hopper=9.x等；次版本号区别则用于架构内部的变种（如服务器芯片与消费芯片的小差异）。在分发软件时，我们需要关注客户GPU的计算能力范围，以确保生成兼容的二进制代码。

计算能力与向下/向上兼容性：较高版本的计算能力GPU通常能够运行较低版本CC编译的代码（硬件向后兼容旧指令），但无法运行高于自身CC的代码。例如，Turing架构GPU（CC7.5）可以执行针对Pascal (CC6.x)编译的程序，但Pascal GPU无法执行仅为CC7.5编译的代码。因此，在编译应用时应指定适当的最低计算能力，以覆盖目标用户的最老GPU型号。同时，可利用虚拟架构（PTX）提供一定的向前兼容，使应用在未来的新GPU上仍能运行（后文详细讨论）。

计算能力演进与主要特性 (CC 6.x 至 9.x 新特性)

从Pascal（CC 6.x）到Hopper（CC 9.0），每一代计算能力都引入了新的硬件功能和编程模型改进。下面按架构世代总结各主要CC版本的新特性，重点关注自定义CUDA内核开发相关的执行模型、Warp级原语、内存和指令集等方面。

Pascal架构 (计算能力6.x)

Pascal（2016年发布，包含GP100、GP104芯片）相对于前代Maxwell带来了多项改进：

• 硬件调度与抢占：Pascal引入了指令级和线程级抢占（Preemption）能力。这意味着GPU可以在执行内核时中断和切换任务（尤其在Tesla P100的6.0中实现），提高了多任务环境下的弹性。此前的架构在运行长时间内核时无法及时中断，Pascal的抢占特性对调试和多进程共享GPU非常有用。

• 统一内存增强：Pascal扩展了统一内存（Unified Memory）的支持，实现CPU与GPU之间的自动内存分页和数据迁移硬件支持。例如，Pascal P100配备硬件页表和“页面迁移引擎”，使GPU能直接访问主机内存，大幅提升统一内存的性能和易用性。还加入了对系统范围的原子操作的支持（通过atomicXXX_system等）以配合统一内存的跨GPU一致性。

• **新数据类型与指令：**Pascal增加了原生半精度FP16运算指令，并在GP100上实现了FP16运算性能达到FP32的2倍（GP104略低）。同时Pascal GP104引入了加速8位整数点积运算的指令（如__dp4a用于4xINT8点积）用于推理应用。尽管这些特性主要面向深度学习，它们也为定制内核提供了更多低精度计算手段。

• 内存系统改进：Pascal延续了Maxwell的统一L1/纹理缓存设计，将L1缓存和共享内存结合以优化内存访问。Pascal的每SM共享内存总量为64KB（GP100）或96KB（GP104），并支持缓存本地内存（寄存器溢出）到L1，从而缓解寄存器不足时的性能损失。此外，Pascal支持原生共享内存原子操作（32位整数）和新增全局双精度原子加，实现原子性能的大幅提升。这些内存和原子特性的改进对自定义内核中大量线程间通信、同步的场景很有帮助。

(Pascal架构还有诸如HBM2高带宽显存、NVLink高速互连等硬件革新，此处不展开，侧重CUDA编程相关部分。)

Volta架构 (计算能力7.0)

Volta（2017年，Tesla V100等）是数据中心导向的架构，带来了CUDA编程模型的重要变革：

• 独立线程调度模型：Volta引入了Independent Thread Scheduling（独立线程调度），打破了以往GPU同一Warp内线程严格同步锁步执行的模型。在CC7.0及以后，每个线程有自己独立的程序计数器和调用栈，可以在warp内独立暂停/继续。这使得warp内线程可执行更加灵活的多分支代码和细粒度同步。不过，这一改变意味着以前依赖warp隐式同步的代码需做修改：开发者必须使用新的warp级原语（如__syncwarp()）来确保需要的同步，并改用带_sync后缀的warp内函数（如__shfl_sync等）来显式指定参与线程掩码。否则，不同线程可能不同时执行到某些代码段，导致数据不一致。总之，Volta改变了CUDA的执行模型，使warp内非锁步执行成为可能，这是开发者在编写内核时需要注意的重大变化。

• 新增加Tensor Core矩阵单元：Volta在每个SM中新增了Tensor Core张量运算单元（每SM 8个），可执行4x4矩阵乘加运算（FP16输入、FP32累加）以大幅加速深度学习矩阵计算。虽然Tensor Core主要用于AI领域，但在自定义CUDA内核中如果涉及矩阵运算，也可以通过Warp级矩阵乘API (wmma) 利用这些单元提升计算吞吐。

• 并行整数与浮点执行：Volta SM配备了独立的INT32运算单元（每SM 64个），能够与FP32单元并行执行整数运算。这意味着内核中可以同时执行整数地址计算和浮点计算而互不阻塞。例如可以在同一循环中并行更新指针和执行浮点计算。对于需要频繁地址索引或指针运算的内核，这提高了指令级并行度。

• 共享内存与缓存：Volta每个SM的共享内存提高到96KB（相比Pascal GP100的64KB提升50%），并且允许单个Thread Block使用全部96KB共享内存（在以前架构中，单Block最多用48KB）。这对需要大量共享内存的内核（如利用共享内存做块级缓存）非常有利。但需要注意，使用超过48KB共享内存需通过CUDA API显式opt-in（设置cudaFuncSetAttribute等）才能生效。Volta还保持了共享内存/L1缓存可配置划分的机制，不过其数据缓存和共享内存总容量达到128KB。

• **其它改进：**Volta支持更大的寄存器文件（每SM 64K 32位寄存器，与Pascal相当）和最高64个warp并发，最大线程数与Pascal相同但因为单warp调度优化，实测可更有效地隐藏指令延迟。总体而言，大多数按照Pascal优化的代码无需改动即可在Volta上获得提速，但需要检查并调整任何依赖warp隐式同步的逻辑，以利用Volta的新调度模型。

Turing架构 (计算能力7.5)

Turing（2018年，GeForce RTX 20系列等）是Volta面向消费市场的延伸，计算能力7.5。它继承并发展了Volta的一些特性，同时针对图形和并行计算做了改进：

• **独立线程调度延续：**Turing沿用了Volta引入的独立线程调度模型。因此，Volta时代需要的warp同步措施（如__syncwarp）在Turing上依然适用。任何假设warp锁步执行的老代码在Turing上也可能遇到问题，需要同样的修改。

• 第二代Tensor Cores：Turing的Tensor Core为第二代，增加了对INT8和INT4低精度矩阵运算的支持（Volta仅支持FP16）。这使Turing在AI推理中可以利用更低精度获得更高吞吐。不过对于一般HPC自定义核而言，Tensor Core仍主要用于矩阵/向量计算场景。

• 统一内存路径 & 缓存架构调整：Turing对SM的内存子系统进行了重构与统一。SM的L1数据缓存、纹理缓存和共享内存之间采用统一的64KB+32KB配置（一种划分是64KB共享+32KB缓存，或反之）。相较Pascal/Titan时代，这种统一设计提高了共享内存和缓存的灵活利用率。Turing还将每个SM的L2缓存增大了一倍（旗舰芯片L2从Pascal的3MB增至6MB）来加速内存访问。此外，Turing每个线程块可使用完整64KB共享内存（Volta为96KB），这比早期架构的48KB限制有提升但低于Volta的上限。

• 新指令和性能改进：Turing引入了可变速率着色、光线追踪 RT Core等图形硬件，不直接影响CUDA核编程。不过Turing在CUDA方面还有一些细节改进，例如**硬件加速的多进程服务（MPS）**及隔离，方便多应用共享GPU；以及细粒度的时钟调度优化等。这些特性提升了Turing在高并发异构负载下的调度效率。对于自定义内核而言，Turing的FP32/INT32核心配置与Volta类似，每SM仍是64个FP32 CUDA Core，但FP64性能在消费卡中依然有限（1/32速率）。总体上，Turing提供了与Volta几乎相当的计算能力，并针对游戏和实时渲染做了专门增强。

Ampere架构 & Ada Lovelace架构 (计算能力8.x)

Ampere架构（2020年推出）涵盖数据中心和消费级产品，计算能力8.x系列。Ampere显著提升了硬件并发和带来了新的编程特性：

• 第三代Tensor Core & TF32: Ampere（例如A100 GPU）配备第三代Tensor Cores，增加支持Tensor Float 32 (TF32)和BF16等数据格式。TF32是一种10位指数、10位尾数的格式，在保持与FP32动态范围相近的同时减少精度，以加速深度学习训练。对HPC内核来说，如果涉及矩阵计算，Ampere的Tensor Core还能支持双精度矩阵运算（HMMA指令支持FP64输入），使某些科学计算也可利用Tensor Core加速。

• 异步数据拷贝和并发计算：Ampere引入了异步内存拷贝指令（如cp.async），允许线程直接将全局内存数据异步加载到共享内存，无需经过寄存器。这意味着内核可以在数据传输的同时重叠执行计算，从而更好地利用内存带宽和计算单元。配合异步拷贝，Ampere还提供了异步屏障（Asynchronous Barrier）机制，将同步操作拆分为“到达”和“等待”两阶段。这使得实现生产-消费模式变得更高效，并支持超越warp甚至线程块粒度的同步手段。对于需要在加载数据和计算之间pipeline重叠的自定义kernel（例如块内预取数据模式），这些新指令可以显著提高性能。

• 共享内存容量提升：Ampere大幅增加了共享内存容量。数据中心芯片A100每SM共有164KB的可配置L1/共享内存（较V100的96KB提高近70%），单个线程块最多可利用160KB共享内存。即使是消费级Ampere（如RTX 30系列，GA102芯片），每SM也支持128KB（其中100KB可用于单个线程块）。这为更大tile大小的并行计算（如更大的矩阵分块乘法）提供了空间。不过需要通过cudaFuncSetAttribute(cudaFuncAttributeMaxDynamicSharedMemorySize, ...)等显式设置，才能启用超过48KB的共享内存分配。

• **MIG和多实例GPU：在A100等数据中心Ampere GPU上，NVIDIA引入了多实例GPU (Multi-Instance GPU, MIG)**功能，允许将一块物理GPU分割为多个逻辑GPU实例，各实例拥有独立的SM、内存和缓存资源。虽然MIG属于GPU管理范畴，不直接影响CUDA内核的编写，但在部署应用时可能需要考虑实例化后的GPU资源上限（如可用SM数、内存大小）。对于面向HPC服务的应用，MIG有助于在一台服务器上并行运行多个作业而互不干扰。

• Ada Lovelace架构 (8.9): Ada Lovelace是Ampere的后继消费级架构（RTX 40系列，CC 8.9）。Ada继承了Ampere的大部分CUDA特性，并进一步提升性能和能效。Ada的SM设计与Ampere类似，但在第四代Tensor Core中加入了对FP8矩阵运算的支持。据NVIDIA发布的数据，Ada的Tensor Core通过FP8“Transformer Engine”使AI推理吞吐提升最高可达5倍。虽然Ada主要改进集中于光线追踪、着色执行重排（Shader Execution Reordering, SER）等图形特性，但对于CUDA开发者来说，其计算部分与Ampere基本相同。需要注意的是，Ada (SM 8.9)仍被视作8.x系列，可以运行为8.6编译的代码。然而为了充分利用Ada的新Tensor Core性能，深度学习代码可使用适配Ada的新库（如TransformerEngine）以开启FP8加速。总之，在HPC自定义内核场景下，Ada的主要益处是更高的时钟、更大的L2缓存和微架构优化，直接的编程模型变化不大。

Hopper架构 (计算能力9.0)

Hopper（2022年，Tesla/Datacenter H100等）作为最新的计算架构，引入了面向超大规模计算的全新特性：

• 线程块簇 (Thread Block Cluster)：Hopper扩展了CUDA的调度层级，引入线程块簇这一新概念。线程块簇是由多个线程块组成的调度单元，一组簇内的多个Block可以部署在多个SM上但具有紧密的邻近性。集群内各线程块不仅可以像以往那样在各自SM的共享内存中合作，还能直接访问彼此SM上的共享内存！这被称为分布式共享内存 (Distributed Shared Memory, DSMEM)。通过DSM，簇内不同线程块之间可以直接读写共享数据、执行原子操作等，并能使用特殊同步指令在跨块范围内屏障同步。线程块簇相当于CUDA层次中介于Block和Grid之间的第4级。在H100上，一个簇最多可包含16个线程块（默认可移植配置下为8个)。这种特性非常适合超大规模并行算法，例如需要跨Block共享中间结果的分布式矩阵分解、大规模排序或归约等，因为以前这些需要借助全局内存+内核间同步才能完成，现在在单kernel簇内即可完成跨SM协作。

• Tensor Memory Accelerator (TMA)：Hopper引入了新的Tensor Memory Accelerator单元，用于在共享内存和全局内存之间进行异步的批量传输。TMA可以在硬件级别处理多维（最多5维）张量的数据搬运，并支持在从共享内存写回全局内存时直接进行元素级的规约计算和按位操作。TMA由CUDA 12提供接口（通过cuda::memcpy_async等）供开发者使用。相较Ampere的cp.async需要软件管理每次复制的字节数，Hopper的TMA更智能高效，能够显著降低大规模数据搬运对SM计算单元的占用。这对带宽受限的内核（典型如大矩阵转置、张量重排）将带来性能改进。

• 显存层次与容量提升：Hopper使用HBM3显存，带宽提高到3 TB/s，比A100的HBM2e高出50%。L2缓存容量增至60MB（Ampere为40MB），SM的L1/共享内存合计容量增至256KB每SM。共享内存方面，每SM最多228KB可用于共享内存（相比A100的164KB进一步提升），单线程块可占用高达227KB共享内存。如此巨大的共享内存结合线程块簇特性，使得以前受限于48/96KB共享内存的算法现在可以在SM内部或跨SM协作中使用上百KB的高速内存，从而降低对全局显存的访问频率。

• 新指令和数学单元：Hopper配备第四代Tensor Core（与Ada同代），支持FP8精度训练和推理，FP8吞吐相当于FP16的2倍。此外，Hopper还增加了一些特定算法加速指令，例如用于生物信息学Smith-Waterman序列比对的DPX指令集合。这些指令对特定领域应用有明显加速效果。对于一般CUDA开发者来说，Hopper在FP32方面每SM能执行的运算也翻倍（128个FP32 ALU vs Ampere的64个），意味着常规浮点计算也有性能提升。

总结来说，Hopper (CC 9.0)结合了更强大的硬件（更多计算单元、更高带宽）和更灵活的编程模型（块簇+DSMEM、TMA），为高性能计算和大规模并行算法提供了前所未有的支持。利用这些新特性需要CUDA 12及以上版本的工具链支持，并通常需要显式使用新的API或编程模型来获益。

编译与部署实践：架构兼容性的策略

理解了不同计算能力和特性，下一步就是确保我们的CUDA应用程序在编译和分发时兼容目标GPU。这里将解释CUDA编译过程中真实架构代码与虚拟架构代码的区别，以及JIT编译、胖二进制、最佳实践编译策略等内容，并给出CMake配置示例。

真实架构 (SASS) vs 虚拟架构 (PTX)

当使用nvcc编译CUDA代码时，会涉及两种层次的目标代码：

• SASS代码（机器代码）：即针对具体GPU硬件架构（如“sm_75”）生成的低级二进制指令代码。这是真正运行在GPU流处理器上的指令序列，针对特定计算能力的GPU进行了优化。SASS代码通常不可跨架构使用——不同计算能力的GPU有不同的指令集和微架构，不能直接执行彼此的机器码。

• PTX代码（虚拟ISA代码）：PTX是CUDA的并行线程汇编语言，是一种虚拟GPU指令集架构。PTX代码类似中间表示（IR），不直接绑定具体硬件，而是在GPU驱动层再编译为实际硬件指令执行。换言之，PTX提供了一层虚拟架构抽象。例如“compute_75”表示兼容计算能力7.5架构特性的PTX代码。运行时，NVIDIA驱动可以将PTX再编译（JIT）为当前GPU的原生指令。

在实践中，nvcc可以生成SASS和PTX的组合。例如，我们可以让编译器输出针对sm_75的SASS二进制，同时embed一份compute_75的PTX代码。当程序在一块RTX 2080 Ti (sm_75)上运行时，驱动会优先使用匹配的SASS代码直接执行；如果程序在未来某个兼容PTX的GPU上运行（假设sm_80+），没有合适的SASS时，驱动将检测到PTX代码并即时编译成该GPU的SASS来执行。概括来说，SASS=已编译好的机器码（快速执行），PTX=待编译的虚拟码（灵活兼容）。

即时编译（JIT）与兼容性

JIT (Just-In-Time) 编译是CUDA驱动在运行时将PTX代码编译为目标GPU机器码的过程。如果我们的应用二进制中不包含某GPU架构的SASS，但包含适用的PTX，那么在第一次加载该CUDA内核时，驱动会自动触发JIT编译，将PTX转译为该GPU的本地代码并缓存起来。下次运行时，可直接使用缓存的代码，避免重复编译开销。JIT机制确保了向前兼容性：比如用CUDA 11编译的程序只含sm_75的SASS和7.5的PTX，也能在sm_86的Ampere上运行——首次运行时驱动JIT出适配sm_86的代码（前提是PTX版本不低于目标GPU所需）。

然而，JIT编译会带来一定的启动延迟，“即时”编译复杂内核可能耗时从几百毫秒到几秒不等，特别在多次启动或多Kernel场景下累积的延迟不可忽视。因此，如果我们预先在编译期生成好目标GPU的SASS，就可以绕过运行时JIT，获得更快的启动和稳定的性能表现。

NVIDIA提供了JIT缓存机制，默认情况下驱动会将JIT生成的二进制缓存到本地磁盘（如~/.nv/ComputeCache），以便后续运行直接加载。但在分发软件时，我们无法假定客户一定启用了缓存，也不能控制首次JIT的时间。所以，对于面向终端用户的发布版本软件，一般希望尽量减少JIT发生。

胖二进制 (Fat Binary) 的创建与作用

为同时兼容多种GPU架构且降低JIT开销，CUDA引入了胖二进制（Fat Binary）的概念。胖二进制就是在一个可执行文件或库中，打包包含了针对多个计算能力的SASS代码，以及可能还有一份PTX代码的二进制。。当应用加载时，CUDA运行时会根据当前GPU的计算能力自动选择匹配的最佳代码路径执行；如果没有直接匹配的SASS，它会退而使用PTX进行JIT编译。因此，胖二进制可以被视为同时携带了多个GPU“版本”的函数实现。

如何生成胖二进制？ 使用nvcc时可以通过多次指定-gencode选项来编译多个目标。例如命令：

nvcc mykernel.cu -gencode arch=compute_75,code=sm_75 \
                 -gencode arch=compute_86,code=sm_86 \
                 -gencode arch=compute_75,code=compute_75

这将生成同时包含sm_75的SASS、sm_86的SASS，以及compute_75 PTX的二进制文件。前两项是两种架构的真实代码，最后一项是一个通用PTX（用于将来超过8.6的新架构JIT）。在运行时，RTX 2080 Ti会用sm_75的代码直接运行，RTX 3090会用sm_86代码，而假如跑在Ada Lovelace (SM 8.9)上，没有对应SASS则驱动会JIT compile那份PTX来运行。

胖二进制的优点是显而易见的：用户无需自行编译，本机即可针对自身GPU运行优化过的代码，无启动延迟。同时也保证了单个可执行程序可以支持广泛的GPU型号（只要涵盖的架构都打包进去了）。其缺点则主要是体积变大——每多支持一种架构，二进制文件就增添一份机器代码。例如一个大型库如果支持5种架构，可能会比只支持1种架构时大几倍。此外，加载二进制到GPU内存时也需时间，虽然一般相比计算时间可以忽略不计。

总的来说，胖二进制是CUDA软件发布中的最佳实践：“用空间换时间”，预植多架构代码以获得即开即用的性能。值得注意的是，如果非常注重安装包大小，开发者也可以只打包最近几年常用的架构，把更老或极新的通过PTX/JIT方式覆盖，这样在常见GPU上无延迟，而在少见GPU上仍可运行（但首次有延迟）。

编译兼容性的最佳实践策略

结合以上原理，这里给出在使用CMake构建CUDA项目时兼顾最大兼容性和性能优化的建议：

1. 确定支持的最低和最高计算能力：根据客户群的GPU情况，选择一个最低CC作为支持下限。例如如果客户至少有Pascal显卡，则最低选6.0或6.1；如果要覆盖老Kepler/Maxwell则可能更低。但一般Pascal (6.x)已足以涵盖近几年的GPU。最高则取决于目前CUDA工具链支持的上限，如CUDA 12支持到sm_90（Hopper）。选择适当范围可以避免包含完全不需要的过老架构代码（增加体积）。

2. 为主要架构生成SASS代码：针对范围内每一代主要架构，都生成对应的SASS。比如要覆盖Pascal、Turing、Ampere、Ada、Hopper，则生成sm_61, sm_75, sm_86, sm_89, sm_90等。这样每代GPU都会有量身定制的代码路径，发挥硬件特性（如使用了适当的新指令）。

3. 包含一份高版本PTX作为后备：为了未来兼容性，可加入一个较高compute_XY的PTX。通常选最高的架构做PTX即可。例如在CUDA 12编译时，可用arch=compute_90,code=compute_90，这样当将来出现Compute 9.x新GPU（如Blackwell架构）时，驱动能够从这份PTX JIT出代码运行。有了PTX，“未来-proof”就更好，避免应用在新GPU上报*“no kernel image available for this device”*错误。

4. 平衡二进制体积和性能：支持的架构越多，二进制越大。我们可以有策略地裁剪：如果确定某些老架构极少使用，可以只提供PTX而不提供它们的SASS。例如“只提供sm_75+的SASS，低于7.5的GPU统一用PTX JIT”。这样在新卡上无损性能，在老卡上第一次运行有延迟但仍能跑。如果倾向**“优先支持新一代”**（如用户提示的要求），可以考虑舍弃对非常老GPU的支持以减少负担，把精力和空间更多放在新架构上。当然，这需要权衡客户实际需求。

5. 利用CUDA Toolkit的新特性：CUDA工具链在不断进步，例如CUDA 12.4引入了nvFatbin库支持运行时组合fatbin，以及NVCC对新架构的默认支持等。紧跟官方最佳实践，比如在新版CUDA中使用CMAKE_CUDA_ARCHITECTURES "all-major"可以让CMake自动包含所有主流架构。密切关注CUDA发行说明，调整编译策略以获得最佳兼容性和性能。

概括而言，推荐的策略是：“尽可能在编译期涵盖客户GPU的架构SASS，并附带高版本PTX保兼容”。只有在确有需要时才依赖运行时JIT，以避免给最终用户带来不必要的性能损失或等待时间。

CMake编译配置示例：使用CUDA_ARCHITECTURES

现代CMake提供了对CUDA架构的简化设置。可以通过CMAKE_CUDA_ARCHITECTURES变量或目标属性来定义需要编译的架构列表。下面给出一个配置示例：

cmake_minimum_required(VERSION 3.18)
project(MyCudaApp LANGUAGES CXX CUDA)

# 指定支持的CUDA架构列表，例如支持: 6.1(Pascal), 7.5(Turing), 8.6(Ampere), 8.9(Ada), 9.0(Hopper)
set(CMAKE_CUDA_ARCHITECTURES 61 75 86 89 90)

add_executable(my_app main.cpp kernel.cu)
# （CMake会根据上面列表自动添加对应的-gencode选项生成fat binary）

在上述配置中，CMake将为每个架构号既生成SASS又生成PTX（相当于每项都arch=compute_X,code=[sm_X,compute_X]）。生成的可执行my_app即是包含多架构的胖二进制。部署该应用时，无需额外操作，不同计算能力的NVIDIA GPU都会从中挑选合适的代码路径运行。如果还需要更精细控制（比如只生成SASS不含PTX，或排除某架构），也可以使用TARGET_PROPERTY CUDA_ARCHITECTURES并加后缀-real/-virtual进行设置。一般情况下，上述简单列表已经满足需求。

通过以上配置，我们的CUDA加速C++应用就能同时兼容多代NVIDIA GPU，并在每代硬件上发挥出最佳性能。这既保障了软件分发给不同客户时的广泛适用性，又避免了用户端繁琐的重新编译或性能不佳的问题。

参考资料： 本文内容参考并引用了NVIDIA官方文档、开发者博客以及社区经验帖等，以确保信息准确和更新。如需深入了解，可以查阅CUDA编程指南附录关于各Compute Capability的详细差异，以及NVIDIA技术博客对新架构特性的解读。