Sunky Blog

Back
10 min
zh-CN
ptx /
cuda /
nvidia /
isa

PTX ISA 9.4：1. 简介

PTX ISA 9.4 第一章中文试译：GPU 数据并行计算、PTX 的设计目标、9.4 新特性与文档结构。

本文介绍 PTX：一种用于并行线程执行（Parallel Thread Execution）的低层次虚拟机及指令集架构（Instruction Set Architecture，ISA）。PTX 将 GPU 呈现为一个可进行数据并行计算的设备。

1.1. 利用 GPU 实现可扩展的数据并行计算#

在市场对实时、高清 3D 图形的持续需求推动下，可编程 GPU 已发展为一种高度并行、支持多线程的众核处理器，拥有强大的计算能力和极高的内存带宽。GPU 尤其适合处理既能表达为数据并行计算、又具有较高算术强度（arithmetic intensity）的问题。数据并行计算是指在多个数据元素上并行执行同一个程序；这里的算术强度是指算术运算次数与内存操作次数之比。由于每个数据元素都执行相同的程序，对复杂流程控制的需求较低；同时，大量数据元素上的计算与较高的算术强度，使得内存访问延迟可以通过计算来隐藏，而不必依赖大型数据缓存。

数据并行处理将数据元素映射到并行执行的线程。许多处理大规模数据集的应用都可以采用数据并行编程模型来加速计算。在 3D 渲染中，大量像素和顶点被映射到并行线程。类似地，图像与媒体处理应用，例如渲染图像的后处理、视频编码与解码、图像缩放、立体视觉和模式识别，也可以将图像块与像素映射到并行处理线程。事实上，数据并行处理所加速的算法远不限于图像渲染和处理领域，还包括通用信号处理、物理仿真、计算金融和计算生物学等。

PTX 为通用并行线程执行定义了虚拟机和指令集架构。PTX 程序在安装时被转换为目标硬件的指令集。PTX 到 GPU 的转换器与驱动程序，使 NVIDIA GPU 能够作为可编程的并行计算机使用。

1.2. PTX 的设计目标#

PTX 为通用并行编程提供稳定的编程模型和指令集。它旨在高效运行于支持 NVIDIA Tesla 架构所定义计算特性的 NVIDIA GPU。CUDA 和 C/C++ 等高级语言的编译器生成 PTX 指令，随后针对目标架构对这些指令进行优化，并将其转换为该架构的原生指令。

PTX 的设计目标包括：

  • 提供跨越多代 GPU 的稳定指令集架构。
  • 使编译后的应用程序获得与 GPU 原生代码相当的性能。
  • 提供与具体机器无关的指令集架构，作为 C/C++ 等编译器的目标。
  • 为应用程序和中间件开发者提供用于代码分发的指令集架构。
  • 为优化型代码生成器和转换器提供通用的源代码级指令集架构，由这些工具将 PTX 映射到具体的目标机器。
  • 便于手工编写库、高性能计算内核（kernel）和架构测试。
  • 提供可扩展的编程模型，覆盖从单个计算单元到多个并行计算单元的不同 GPU 规模。

1.3. PTX ISA 9.4 版本#

PTX ISA 9.4 版本引入了以下新特性：

  • 新增对 sm_107 目标架构的支持。
  • 新增对 sm_107f 目标的支持，该目标支持架构家族特定特性（family-specific features）。
  • 新增对 sm_107a 目标的支持，该目标支持架构特定特性（architecture-specific features）。
  • 为 set 指令新增 .u8x4、.s8x4、.u16x2、.s16x2 指令类型支持。
  • 为 cvt 指令新增 .pzo 限定符支持。
  • 当 cvt 指令的目标类型为 .e4m3x2、.e5m2x2、.e2m3x2、.e3m2x2、.e2m1x2 时，新增 .rz 舍入模式支持。
  • 为 mbarrier.expect_tx、mbarrier.complete_tx、mbarrier.arrive、mbarrier.arrive_drop、cp.async.bulk、cp.async.bulk.tensor 和 tcgen05.commit 指令新增 .multicast::cluster::32b 限定符支持。
  • 为 tcgen05.commit 指令新增 .sync_restrict::shared::read::mma::a 限定符支持。
  • 为 cp.async.bulk.prefetch 和 cp.async.bulk.prefetch.tensor 指令新增 .level::eviction_priority 限定符支持。
  • 新增 applypriority.async.bulk 和 applypriority.async.bulk.tensor 指令支持。
  • 为 cp.async.bulk.tensor、cp.reduce.async.bulk.tensor、cp.async.bulk.prefetch.tensor 指令新增 .override::global_address 和 .override_attribute 限定符支持。
  • 为 cp.async.bulk.tensor 和 cp.reduce.async.bulk.tensor 指令新增 .im2col_no_offs::w 限定符支持。
  • 为 tcgen05.ld 和 tcgen05.ld.red 指令新增 .spcompress 限定符支持。
  • 新增 spcompress 和 spdecompress 指令支持。
  • 为 cvt 指令新增 .ue5m3x2 指令类型支持。
  • 为 cvt 指令新增 .scaled::n1::ue8m0 限定符支持。
  • 为 tcgen05.mma 和 tcgen05.mma.sp 指令新增 UE5M3 缩放因子类型。
  • 为 tcgen05.mma、tcgen05.mma.sp、tcgen05.mma.ws、tcgen05.mma.ws.sp 指令新增 .kind::ti16 限定符支持。
  • 新增 add、sub、mul 和 fma 指令对 .f16x2 / .bf16x2 / .f32x2 类型的支持。
  • 新增 add、sub、mul 和 fma 指令对 FP8/FP6/FP4 x4 打包类型的支持。
  • 为 tcgen05.alloc 和 tcgen05.dealloc 指令新增 .exclusive 限定符支持。
  • 为 mbarrier.test_wait 和 mbarrier.try_wait 指令新增 .phase_type::* 限定符支持。
  • 为 mbarrier.test_wait 和 mbarrier.try_wait 指令新增 reportPredicate 和 reportValue 操作数支持。
  • 为 mbarrier 指令新增 .layout 限定符支持，并新增 mbarrier.check_layout 指令支持。
  • 为 cp.async.bulk 和 cp.async.bulk.tensor 指令新增 .report_mechanism 限定符支持。
  • 为 tcgen05.mma 指令新增 .decompress::lut::b 限定符支持。
  • 为 tcgen05.mma 指令新增 .collector::b::* 限定符支持。
  • 新增 .loc_intermediate 指示符和 .nv_intermediate_source_section 结构支持，用于中间表示（IR）的源代码位置映射。
  • 为 prefetch 指令新增 .valid_addr 和 .L1::32B 限定符支持。
  • 为 atom、red、cp.reduce.async.bulk 和 multimem.cp.reduce.async.bulk 指令新增 .f32 类型下的 .noftz 限定符支持。
  • 为 ld 指令新增 .proxy::readonly 限定符支持。
  • 新增 .minperctamemory 指示符，以及 %perctamemoryoffset、%perctamemorysize 特殊寄存器支持，用于每个 CTA（协作线程数组）的全局内存。
  • 为 ldmatrix 指令新增 .m8n16 形状下的 .s8.s4 类型支持。
  • 为 fence.proxy.async::generic 指令新增 .sync_restrict::shared::cluster::read 限定符支持。
  • 为 fence.proxy 指令新增 .acquire、.release 限定符与 .alias 限定符配合使用的支持。
  • 新增 fabric.try_get.async.tensor、fabric.try_put.async.tensor 和 fabric.try_red.async.tensor 指令支持。
  • 新增 fabric.try_atom 指令支持。
  • 新增 createpolicy.range.fabric 指令支持。
  • 为 fabric.try_get、fabric.try_put 和 fabric.try_red 指令新增 .level::cache_hint 限定符支持。

1.4. 文档结构#

本文档的信息按以下章节组织：

参考文献

PTX ISA 9.4：9. 指令集

PTX ISA 9.4
EN