Related Resources
数据中心/云端

NVIDIA Blackwell 和 NVIDIA CUDA 12.9 引入基于系列的架构特性

2025年 5月 1日
By and

最早进入 NVIDIA GPU CUDA 平台的架构设计决策之一是支持 GPU 代码的向后兼容性。这种设计意味着，新 GPU 应该能够运行为之前的 GPU 编写的程序，而无需进行修改。它由 CUDA 的两个基本特性完成：

PTX 是面向 NVIDIA GPU 的虚拟 ISA。您可以将其想象成组装代码，但它不限于特定的物理芯片硬件架构，其设计足够通用，可以与未来的 GPU 架构兼容。

自 NVIDIA 创建 CUDA 平台使开发者能够为 GPU 编写通用程序以来，PTX 一直是 CUDA 不可或缺的一部分。为之前的 GPU 构建的 PTX 代码可以由当前的驱动进行 JIT 编译，并在当前的 GPU 上运行，无需修改。

举个例子。这是一段简单的代码，可打印 GPU 名称和计算能力，还可从 GPU 内核内部打印 hello。

#include <stdio.h>
#include <iostream>

__global__ void printfKernel()
{
    printf(">>>>>>>>>>>>>>>>>>>>\n" );
    printf("HELLO FROM THREAD %d\n", threadIdx.x );
    printf(">>>>>>>>>>>>>>>>>>>>\n" );
}

int main(int argc, char** argv)
{

    // Query and display device properties
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, 0);

    std::cout << deviceProp.name << std::endl;
    std::cout << "Compute Capability: " << deviceProp.major
        << "." << deviceProp.minor << std::endl;
    printfKernel<<<1,1>>>();
    cudaDeviceSynchronize();

    std::cout << "End Program" << std::endl;
    return 0;
}

当我们使用 CUDA 12.8 编译此代码并在配备 NVIDIA RTX 4000 Ada 的系统上运行时，我们会得到以下结果：

$ nvcc -o x.device_info device_info.cu

$ ./x.device_info
NVIDIA RTX 4000 Ada Generation
Compute Capability: 8.9
>>>>>>>>>>>>>>>>>>>>
HELLO FROM THREAD 0
>>>>>>>>>>>>>>>>>>>>
End Program

由于我们没有为 NVCC 指定任何编译器标志，因此它使用此版本编译器支持的最低 PTX 目标。您可以使用 cuobjdump 检查可执行文件，以查看代码中的 PTX 架构和 CUDA 二进制 (cubin) 架构 (为简洁起见，会对输出进行截图)：

$ cuobjdump x.device_info

Fatbin elf code:
================
arch = sm_52

>>> snipped <<<

Fatbin ptx code:
================
arch = sm_52

>>> snipped <<<

您可以看到 ELF (即 binary) 和 PTX。当您看到这样的输出时，表示 cubin 和 PTX 都嵌入到目标文件中。架构为 sm_52，即 Compute Capability (CC) 5.2。CC 由数字 X.Y 表示，其中 X 是主要修订版本号，Y 是次要修订版本号。

返回示例。GPU 为 CC 8.9，如运行代码时的打印输出所示，那么此代码如何在此 GPU 上运行？

这就是 JIT 编译发挥作用的地方。CUDA 驱动程序 JIT 编译 PTX 以在 CC 8.9 GPU 上运行。只要您的代码包含由等同于或更早于 GPU 架构的架构生成的 PTX，您的代码就能正常运行。

您可以通过稍微更改编译器标志来验证这一点。添加参数 -gencode arch=compute_75,code=compute_75。这将告知 NVCC 您希望它为您的应用程序构建版本为 compute_75 (计算能力 7.5) 的 PTX，然后将该 PTX 放入可执行文件并使用 cuobjdump 进行验证。有关 NVCC 如何构建 PTX 和二进制代码的更多信息，请参阅 理解 PTX (CUDA GPU 计算的汇编语言) 中的图 1。

您可以看到它正常运行。

$ nvcc -gencode arch=compute_75,code=compute_75 -o x.device_info device_info.cu

$ ./x.device_info
NVIDIA RTX 4000 Ada Generation
Compute Capability: 8.9
>>>>>>>>>>>>>>>>>>>>
HELLO FROM THREAD 0
>>>>>>>>>>>>>>>>>>>>
End Program

现在，如果您将 code=compute_75 更改为 code=sm_75，这将告知 NVCC 构建与之前 (arch=compute_75) 相同的 PTX。但是，NVCC 应将 PTX 编译为 SM_75 的 cubin，并将该 cubin 放入可执行文件中，而不是将其保留在可执行文件中用于 JIT 编译。同样，您可以使用 cuobjdump 进行验证。结果如下：

$ nvcc -gencode arch=compute_75,code=sm_75 -o x.device_info device_info.cu

$ ./x.device_info
NVIDIA RTX 4000 Ada Generation
Compute Capability: 8.9
End Program

如果仔细观察，您会发现 “HELLO FROM THREAD 0” 未打印。我们省略了所有错误检查代码，以使代码示例更清晰。

如果我们像在真实代码中一样包含错误检查，您将看到 GPU 内核未启动，并且返回的错误消息是 “No kernel image is available for execution on the device”。这意味着应用程序中没有与此 CC 8.9 设备兼容的内核代码，因此内核从未启动。

直至 CC 8.9 (含 CC 8.9) 并支持 CUDA 的所有 GPU (Tegra 除外， 因为它们遵循不同的规则 ) 应遵循的经验法则如下：

NVIDIA Hopper 中引入的架构特定功能集

从 NVIDIA Hopper 架构 (CC 9.0) 开始，NVIDIA 推出了一套高度专业化的小型功能集，这些功能被称为 特定架构 ，只能保证在特定的目标架构上存在。其中大多数功能与 Tensor Cores 的使用有关。

要使用这些功能，您必须在应用中嵌入 PTX 或 cubin 代码，使用 compute_90a 标志表示 PTX，或 sm_90a 标志表示编译中的 cubin。使用 a 后缀构建特定于架构的目标时，PTX 或 cubin 代码无法向前兼容任何未来的 GPU 架构。

例如，您使用以下 NVCC 行编译 CUDA 核函数：

$ nvcc -gencode arch=compute_90a,code=sm_90a -c kernel.cu

在本例中，您的代码仅在 CC 9.0 的设备上加载和运行。在使用特定于架构的 a 后缀时，PTX 或 cubin 都无法向前兼容。

NVIDIA Blackwell 中引入的家族特定功能集

从 NVIDIA Blackwell 架构和 CUDA 12.9 开始，我们引入了一类新功能： 特定于系列的功能 。

“该系列的特定功能类似于架构特定的功能，不同之处在于它们由具有多个次要计算能力的设备提供支持。一个系列的所有设备共享相同的主要计算能力版本。该系列的特定功能可确保在同一 GPU 系列中提供，其中包括具有相同主要计算能力和较高次要计算能力的后续 GPU。”

特定于系列的编译器目标类似于特定于架构的目标，但您可以使用 f 后缀，而不是使用 a 后缀的编译器目标。

如需详细了解同一系列中的 GPU，请参阅 编程指南 和 CUDA 计算能力 页面。如需详细了解该系列特定目标中包含的功能，请参阅 PTX ISA 中的表格。

例如，您使用以下 NVCC 行编译 CUDA 核函数，该行可调用系列特定的代码生成目标：

$ nvcc -gencode arch=compute_100f,code=sm_100 -c kernel.cu

在本例中，您为 sm_100f 系列生成特定于架构的 cubin 代码，并且您的代码将仅在具有 compute capability 10.x 的设备上运行。

目前，这是具有 10.0 和 10.3 计算能力的 GPU。如果引入具有 10.x 计算能力的新 GPU，这些 GPU 上的代码也将兼容，因为它们属于 sm_100f 系列。在本例中，code=sm_100 和 code=sm_100f 是彼此的别名，并将生成将在 sm_100f 系列设备上运行的相同 cubin。

在 NVCC 中思考这些不同的特征集的方式如下：

  • 无后缀：您的 PTX 或 cubin 兼容性一如既往。
  • 后缀 f： 无论您是停留在 PTX 还是从该代码生成 cubin，该代码都兼容在具有相同主要计算能力版本以及具有相同或更高次要计算能力版本的 GPU 设备上运行。
  • a 后缀： 代码仅在该特定 CC 的 GPU 上运行，不得在其他 GPU 上运行。

开发者指南

现在，我们已经解释了如何使用 NVCC 构建架构和特定系列的代码目标，我们希望为您在构建应用程序时应该执行的操作提供建议。

通常，您应构建能够在尽可能多的架构上运行的代码。只要您没有使用架构或系列特定功能，就不必在应用中包含架构或系列特定目标，而且您可以继续像往常一样构建代码。即使您使用的是使用架构或系列特定功能的库，只要这些库以二进制形式分发，它们也能正常运行。

那么，您何时需要使用 family 或特定于 architecture 的编译器目标？

如前所述，在使用主要与 Tensor Cores 相关的功能时，特别是通过 PTX 对 Tensor Cores 进行编程时，会使用这些目标。如果您直接编写 PTX 并使用系列或架构特定的功能，则必须分别使用 f 或 a 标志构建代码，具体取决于您使用的 PTX 指令是否在 f 特征集中，或者它们是否仅在 a 特征集中。

如果您希望在不同 CC 的 GPU 之间实现可移植性，则必须在代码中包含适当的防护措施，以确保在不具备这些功能的不同 GPU 架构上运行时，有可用的备用代码路径。根据您正在使用的系列和架构特定功能，使用以下宏控制代码路径：

  • __CUDA_ARCH_FAMILY_SPECIFIC__
  • __CUDA_ARCH_SPECIFIC__

这些宏的定义类似于 __CUDA_ARCH__。有关更多信息，请参阅《 CUDA Programming Guide 》。

例如，如果您正在构建应用程序并使用 CUTLASS 等头文件库或任何包含 CUTLASS 的库（例如 cuBLASDx ），并且您在 CC 9.0 （NVIDIA Hopper）或更高版本上运行应用程序，则应针对将运行代码的 GPU 设备构建特定于架构的目标。

CUTLASS 专为实现高性能而设计，具有特殊的代码路径，可使用特定于架构的功能来更大限度地提高性能。这些库在内部已经有后备路径，可以与其他架构完全兼容。

换言之，如果您使用的是库，则无需担心使用宏的后备路径。

付诸实践

现在，我们已经讨论了架构和系列特定目标，以及何时使用它们，我们将把所有内容放在一起。

一般情况

首先要确定的是，您的代码使用的是架构还是系列特定功能。您可能知道自己是否在使用这些功能，因为您要直接编写 PTX，或者包含像 CUTLASS 这样的头文件库。如果不是这样，对于大多数开发者来说都是如此，那么构建应用程序就像以前一样。

为了提供最佳性能和未来兼容性，通常的指导是为您知道代码将运行的每个架构构建 binary code。这可提供最佳性能。

您还应嵌入适用于最新架构的 PTX，以提供出色的未来兼容性。例如，您可能知道自己的代码将在 CC 8.0、9.0 和 10.0 的设备上运行。以下代码示例展示了如何为这些架构编译二进制文件，以及如何为将来的兼容性编译 CC 10.0 PTX。

$ nvcc -gencode arch=compute_80,code=sm_80 
       -gencode arch=compute_90,code=sm_90
       -gencode arch=compute_100,code=sm_100 
       -gencode arch=compute_100,code=compute_100 -c kernels.cu

系列特定功能

如果您选择使用在不同架构中不可移植的特定功能来优化代码，则应首先确定这些功能是否属于该系列的特定功能集。

如果是，那么您可以使用后缀 f 构建目标，并且您将与该系列兼容。如果您希望可移植到系列以外的 GPU，则必须为任何使用系列特定功能的代码添加备用代码路径。

通常，这是通过应用中的条件宏保护特定系列的代码来实现的。在扩展先前的示例并为 CC 10.0 添加系列特定功能后，您对 NVCC 的使用可能类似于以下代码示例：

$ nvcc -gencode arch=compute_80,code=sm_80 
       -gencode arch=compute_90,code=sm_90
       -gencode arch=compute_100f,code=sm_100 
       -gencode arch=compute_100,code=compute_100 -c kernels.cu

这让您的代码能够在 CC 8.0、9.0 和 10.0 的设备上运行，并具有 10.0 的系列特定功能。通过嵌入式 PTX，您的代码也将在未来的设备上运行。

使用特定于系列的功能的另一种可能情况是，您知道您的应用必须利用这些功能，而应用仅设计为在该系列的设备上运行。例如，如果您将代码设计为仅使用 100f 系列的功能，并且只想在此系列的设备上运行，则应用程序构建类似于以下代码示例：

$ nvcc  -gencode arch=compute_100f,code=sm_100 -c kernels.cu

在这种情况下，您的代码只能在此系列的设备中进行移植。

架构特定功能

如果您已确定系列特定功能不足以满足您的应用需求，并且必须使用架构特定功能集中的功能，则必须使用 a 标志进行构建。

与使用 f 进行构建的情况类似，您必须使用应用内的条件宏保护代码，从而确定必须在应用中构建何种代码可移植性。为了获得与上一个示例相同的可移植性，请按照以下代码示例构建代码：

$ nvcc -gencode arch=compute_80,code=sm_80 
       -gencode arch=compute_90,code=sm_90
       -gencode arch=compute_100a,code=sm_100a 
       -gencode arch=compute_100,code=compute_100 -c kernels.cu

您的代码将具有相同的兼容性，可在 CC 8.0、CC 9.0、CC 10.0 及更高版本上运行，

与仅为特定系列设计应用类似，您也可以选择针对特定架构进行设计和优化。如果您设计和编写的应用程序使用特定于架构的功能，并且知道它不必在任何其他 GPU 上运行，则可以构建类似于以下代码示例的应用程序：

$ nvcc  -gencode arch=compute_100a,code=sm_100a -c kernels.cu

您的应用程序将仅在 CC 10.0 上运行，不兼容任何其他 GPU。

总结

总结一下，以下简单流程说明了您应如何考虑构建您的代码：

  • 您是直接编写 PTX，还是直接调用 CUTLASS 等库？否则，您不需要添加 f 或 a 标志。像往常一样构建代码。
  • 如果您要编写 PTX，或包含仅包含头文件的库，则需要确定该库是否使用架构或系列特定特征集中的特征，如果是，则确定要使用的 f 或 a 标志。您需要参阅该库的文档，以确定为您的架构构建的最佳方式。例如， CUTLASS 构建说明指定在构建 CC 9.0 和 10.0 的设备时使用 a 标志

本文介绍了很多内容，向您展示了如何构建可以使用架构和系列特定功能的代码。我们想明确的是，在 NVCC 编译行中使用 a 和 f 后缀并不是一种神奇的优化技术。要具体使用这些功能，您必须直接编写 PTX 或调用具有此功能的库。

立即下载 CUDA 12.9 ，开始在 Blackwell 上的代码中使用家族和架构特定的功能。

致谢

感谢以下 NVIDIA 贡献者：Cody Addison、Vyas Venkataraman、Rob Armstrong、Girish Bharambe 和 Mridula Prakash。

 

相关资源

标签

开发与优化 | 数据中心/云端 | 模拟/建模/设计 | Blackwell | CUDA | Deep dive | Development & Optimization | HPC/科学计算 | Intermediate Technical | nvcc | 仿真/建模/设计 | 数据中心/云

关于作者

Jonathan Bentz 领导 NVIDIA 的 CUDA 技术营销工程团队，其团队专注于创建和提供引人入胜的内容，并与 CUDA 开发者建立联系。Jonathan 拥有爱荷华州立大学化学博士学位和计算机科学硕士学位。
Tony Scudiero 是 CUDA 平台的技术营销工程师。他致力于将 CUDA 带给各种类型和能力的开发者。在 NVIDIA 任职期间，他曾使用过大型 HPC 系统和应用、实时声学模拟 (VRWorks Audio) 和 Omniverse RTX 渲染器。