GPU异构计算和CUDA程序简介¶

Author:: 吴超，中国科大超级计算中心

背景¶

近年来，随着人工智能、高性能数据分析和金融分析等计算密集型领域的兴起，传统通用计算已经无法满足对计算性能的需求，异构计算越来越引起学术界和产业界的重视。

异构计算是指采用不同类型的指令集和体系架构的计算单元组成系统的计算方式。相比传统CPU，异构计算可以实现更高的效率和更低的延迟。目前的异构计算引擎主要有图形处理器（GPU，Graphics Processing Unit）、现场可编程门阵列（FPGA，Field Programming Gate Array）、专用集成电路（ASIC）等。

当前的通用CPU设计得已经很复杂，配有几十个核心，运行频率高达几GHz，每个核心有自己的独立缓存。通常CPU已具备一级、二级、三级缓存。而GPU是目前科研领域比较常用的硬件计算工具。GPU的计算核心数通常是CPU的上百倍，运行频率尽管比CPU的低，但是核心数量多，整体性能好。所以，GPU比较适合计算密集型应用，比如视频处理、人工智能等，现在传统的科学计算、工程计算等也开始越来越适合在GPU上运行。相比来说，CPU的缺点就是太通用了，数据读写、计算、逻辑等各种功能都得照顾，反而影响了计算性能。

GPGPU¶

通用图形处理器（GPGPU，General Purpose Graphics Processing Unit）最早由NVIDIA公司的Mark J. Harris于2002年提出。基于图形渲染管线的流水线特征，GPU本质上是一个可同时处理多个计算任务的硬件加速器。由于GPU中包含了大量的计算资源，Mark J. Harris自2002年就开始尝试在GPU上做通用并行计算方面的研究。在此阶段，由于架构及编程平台的限制，研究人员采用将目标计算算法转换为图形运算算法的方式，使用GPU来实现通用并行计算需求。

NVIDIA公司提出Tesla统一渲染架构以及CUDA（Compute Unified Device Architecture，计算统一设备架构）编程模型后，NVIDIA公司的GPU开始了对通用并行计算的全面支持。在CUDA提出近两年之后，开放计算语言标准OpenCL 1.0发布，这标志着利用GPU进行通用并行计算已基本成熟。目前市场上应用甚广的GPU芯片除了完成高质量的图形渲染之外，通用并行计算也已经成为一个主流应用。GPGPU在各个方面得到了不同GPU厂家为GPU通用计算提供的编程模型与平台，如CUDA和OpenCL，这些编程模型在C/C++基础之上做了面向大规模通用并行计算的语法扩展，为程序员提供了更好的、面向GPU的编程接口。

GPGPU通常由成百上千个架构相对简易的基本运算单元组成。在这些基本运算单元中，一般不提供复杂的诸如分支预测、寄存器重命名、乱序执行等处理器设计技术来提高单个处理单元性能，而是采用极简的流水线进行设计。每个基本运算单元可同时执行一至多个线程，并由GPGPU中相应的调度器控制。GPGPU作为一个通用的众核处理器，凭借着丰富的高性能计算资源以及高带宽的数据传输能力在通用计算领域占据了重要的席位。虽然各个GPGPU厂商的芯片架构各不相同，但几乎都是采用众核处理器阵列架构，在一个GPU芯片中包含成百上千个处理核心，以获得更高的计算性能和更大的数据带宽。

GPU中执行的线程对应的程序通常成为内核（kernel），这与操作系统中的内核是完全不同的两个概念。除此之外，GPU中执行的线程与CPU或者操作系统中定义的线程也有所区别，GPU中的线程相对而言更为简单，所包含的内容也更为简洁。在GPU众核架构中，多个处理核心通常被组织成一个线程组调度执行单位，线程以组的方式被调度在执行单元中执行，如NVIDIA的流多处理器、AMD的SIMD执行单元。同一个线程组中的线程执行相同的程序指令，并以同步的方式执行，每个线程处理不同的数据，实现数据级并行处理。不同GPU架构对线程组的命名各不一样，如NVIDIA将线程组称为warp，AMD将线程组称为wavefront。线程组中包含的线程数量各不相同，从4个到128个不等。除此以为，线程组的组织执行模式也各不相同，常见的执行模式有SIMT（Single Instruction Multiple Threads，单指令多线程）执行模式和SIMD（Single Instruction Multiple Data，单指令流多数据流）执行模式两种。

在一个GPU程序中，避免不了对数据的加载和存储，同时也避免不了条件分支跳转指令。这两类指令通常会引起程序以不可预测的情况执行。对于前者，在第一级高速缓存命中缺失的情况下，指令的执行周期将不可预测。为了避免执行单元因为数据加载或者存储原因而造成运算资源的浪费，GPU的每个执行单元通常设置线程组缓冲区，以支持同时执行多个线程组。线程组之间的调度由线程组硬件调度器承担，与软件调度器不同的是，硬件调度过程一般为零负载调度。在执行单元中，即将执行的线程组首先被调度到缓冲区中，以队列的方式组织，当线程组被调度执行时，调度器从线程组队列中选择一个准备好的线程组启动执行。采用这种线程调度执行方式，可有效解决指令之间由于长延时操作所引起的停顿问题，更高效的应用执行单元中的计算资源。对于后者，在线程级并行执行过程中，条件分支指令的执行特点决定了程序执行的实际效率。无论是SIMD执行模式或是SIMT执行模式，当一组线程均执行相同的代码路径时可获得最佳性能。若一组线程中的每个线程各自执行不同的代码路径，为了确保所有线程执行的正确性，线程组中的多线程指令发送单元将串行地发送所有的指令代码，代码的执行效率将受到严重的影响。GPU架构采用各种控制方法来提高条件分支指令的执行效率。

背景知识大部分内容引自 [1]。

GPU异构计算¶

CPU-GPU协同是实现高性能计算的必要条件，称为CPU-GPU异构计算（HC，Heterogeneous Computing）。它通过将应用程序的计算密集型部分卸载到GPU来提供更高的性能，而其余代码仍然在CPU上运行，能智能地结合CPU和GPU的最佳特性以实现高计算增益，旨在将每个应用程序的需求与CPU/GPU架构的优势相匹配，并避免两个处理单元的空闲时间。需要新的优化技术来充分发挥HC的潜力并朝着百亿级性能的目标迈进。

了解CPU和GPU之间差异的一种简单方法是比较它们处理任务的方式。CPU由几个针对顺序串行处理优化的内核组成，而GPU具有大规模并行架构，由数千个更小、更高效的内核组成，旨在同时处理多个任务。

在GPU上解决计算问题原则上类似于使用多个CPU解决问题。手头的任务必须拆分为小任务，其中每个任务由单个GPU内核执行。GPU内核之间的通信由GPU芯片上的内部寄存器和内存处理。CUDA或OpenCL等特殊编程语言不是使用消息传递进行编程，而是提供主机CPU之间的数据交换和同步GPU内核的机制。

一个现代超级计算系统实际上可能由大量节点组成，每个节点包含2到32颗常规CPU以及1到16个GPU。通常还会有一个高速网络和一个数据存储系统。该系统的软件可以使用传统编程语言（如C/C++、Fortran等）的组合编写，结合用于CPU并行化的消息传递系统以及用于GPU的CUDA或OpenCL。所有这些组件都必须进行调整和优化，以实现整个系统的最佳性能。

CUDA编程框架¶

CUDA是由NVIDIA公司推行的一套并行编程框架，目前只有NVIDIA的GPU支持该框架，其开发语言主要为CUDA C。作为一种GPU的并行开发语言，CUDA的API涉及设备管理、存储管理、数据传输、线程管理、事件管理等功能。

CUDA的存储模型主要分为全局存储（global memory）、局部存储（local memory）、共享存储（shared memory）、常量存储（constant memory）和纹理存储（texture memory）等存储类型。不同的存储类型，其存储容量、可见程度、读写速度差异巨大，需要在程序设计中根据各自特点和应用问题的需求合理调配。

CUDA的编程模型和执行模型按照层次结构分层设计。CUDA的执行模型由3个层级组成，最基础的执行单位是线程（thread），多个线程组成一个线程块（block），多个线程块形成线程网格（grid）。

CUDA编程模型作为一个异构模型，其中使用了CPU和GPU。在CUDA中，主机（host）指的是CPU及其存储器，设备（device）是指GPU及其存储器。在主机上运行的代码可以管理主机和设备上的内存，还可以启动在设备上执行的内核函数（kernel）。这些内核由许多GPU线程并行执行。

鉴于CUDA编程模型的异构性， CUDA C程序的典型操作序列是：

声明并分配主机和设备内存。
初始化主机数据。
将数据从主机传输到设备。
执行一个或多个内核。
将结果从设备传输到主机。

对于更多的CUDA编程细节可以在NVIDIA官网CUDA开发者页面 [2] 找到详细的资料，读者可以自行查阅。

NVCC编译引擎¶

CUDA应用程序的源代码由传统的C++主机代码和GPU设备函数混合组成。CUDA编译过程将设备函数与主机代码分开，使用专有的NVIDIA编译器和汇编器编译设备函数，使用可用的C++主机编译器编译主机代码，之后将编译过的GPU函数嵌入到主机对象（object）文件。在链接阶段，向最终生成的可运行二进制文件添加特定的CUDA运行时库函数，如支持远程SPMD过程调用的运行时库函数、显式GPU操作的运行时库函数（如分配GPU内存缓冲区，主机与设备之间的数据传输等）。

编译过程涉及每个CUDA源文件的拆分、编译、预处理和合并步骤。为了将上述复杂的编译过程向开发人员隐藏，NVIDIA公司设计了CUDA编译器引擎程序nvcc。nvcc接受一系列常规编译器选项，例如宏定义和头文件、函数库路径设置，支持编译过程的组合。所有非CUDA编译步骤都被转发到nvcc支持的主机C++编译器，编译选项到主机C++编译选项的转换也由nvcc自动完成。

nvcc预定义宏¶

预定义宏	含义
__NVCC__	编译C/C++/CUDA源文件时预定义
__CUDACC__	编译CUDA源文件时预定义
__CUDACC_VER_MAJOR__	NVCC主版本号
__CUDACC_VER_MINOR__	NVCC次版本号
__CUDACC_VER_BUILD__	NVCC编译版本号

支持的输入文件后缀¶

输入文件后缀	描述
.cu	CUDA源文件，包含主机代码和设备函数
.c	C源文件
.cc, .cxx, .cpp	C++源文件
.o, .obj	目标文件（object file）
.a, .lib	库文件（library file）
.res	资源文件（resource file）
.so	共享目标文件（shared object file）

常用编译选项¶

文件和路径配置

选项	描述
-o file	配置输出文件名和路径
-l library,…	配置链接阶段链接的库文件
-D def,…	定义预处理阶段使用的宏
-U def,…	取消宏定义
-I path,…	配置头文件搜索路径
-L path,…	配置库文件搜索路径
-cudart {none\|shared\|static}	配置CUDA运行时库的类型，默认使用静态（static）
-cudadevrt {none\|static}	配置CUDA设备运行时库的类型，默认使用静态（static）

编译器/链接器选项

选项	描述
-pg	生成供gprof使用的可执行代码
-g	编译带调试信息的主机代码
-G	编译带调试信息的设备代码
-O level	指定主机代码的优化级别
-dopt kind	允许设备端代码优化。当不指定-G选项时，设备端代码优化是默认的行为
-shared	生成共享库
-x {c\|c++\|cu}	显式指定待编译的输入文件的编程语言，而不是由编译器根据文件后缀自动判断
-std {c++03\|c++11\|c++14\|c++17}	指定c++标准的版本

特定阶段编译选项

下表列举了可以直接传递给 nvcc 封装的内部编译工具的编译选项。通过这些选项的应用， nvcc 不需要具备对内部编译工具的过多细节的了解。

选项	描述
-Xcompiler options,…	指定直接传递给编译器/预处理器的编译选项
-Xlinker options,…	指定直接传递给链接器的编译选项
-Xarchive options,…	指定直接传递给库管理器的编译选项
-Xptxas options,…	指定直接传递给ptxas（PTX优化汇编器）的编译选项
-Xnvlink options,…	指定直接传递给nvlink（设备链接器）的编译选项

GPU代码生成选项

选项	描述
-arch {arch\|native\|all\|allmajor}	指定编译阶段使用的虚拟GPU类型
-code code,…	指定汇编优化使用的具体GPU类型
-use_fast_math	使用快速数学计算库

为了实现架构演进，NVIDIA的GPU以不同的世代（generation）发布。新一代产品在功能和/或芯片架构方面进行重大改进，同时同一代产品中的GPU型号仅在配置方面存在次要差别，对功能、性能的影响适中。不同代的GPU其应用程序的二进制兼容性是没有保证的。例如，为Fermi GPU编译的CUDA应用程序很可能无法在Kepler GPU上运行（反之亦然）。这是因为每一代的指令集和指令编码与其他世代的指令编码都不相同。同一代的GPU由于共享相同的指令集，在满足特定条件下其二进制兼容性可以得到保证。特定条件通常是指两个没有功能差异的GPU版本之间的情况（例如，当一个版本是另一个版本的缩减版），或者当一个版本在功能上完全包含在另一个版本中。后者的一个例子是基础Maxwell版本sm_52，其功能是所有其他Maxwell版本的一个子集：任何为sm_52编译的代码将可以在所有Maxwell GPU上运行。

nvcc 编译命令总是使用两个架构：一个虚拟的中间架构，加上一个真实的GPU架构（指定代码将运行的平台）。要使 nvcc 命令有效，真实架构必须是虚拟架构的实现。

虚拟GPU完全由提供给应用程序的能力和特征定义。虚拟架构提供了一个通用的指令集合，并且不涉及二进制编码格式。虚拟架构列表如下：

虚拟架构（-arch参数）	特征描述
compute_35 compute_37	Kepler架构支持统一内存编程支持动态并行
compute_50 compute_52 compute_53	+Maxwell架构支持
compute_60 compute_61 compute_62	+Pascal架构支持
compute_70 compute_72	+Volta架构支持
compute_75	+Turing架构支持
compute_80 compute_86 compute_87	+Ampere架构支持

在CUDA的命名方案中，GPU被命名为sm_xy，其中x表示GPU的世代编号，y表示该世代的版本。为了便于比较GPU的能力，执行特定的命名设计规则，如果x1y1 <= x2y2，那么sm_x1y1的所有非ISA相关能力都包括在sm_x2y2中。由此可见，sm_52确实是基础麦克斯韦模型，这也解释了为什么表格中的高条目总是对低条目的功能扩展（表格中用加号表示）。

真实架构（-code参数）	特征描述
sm_35 sm_37	Kepler架构支持统一内存编程支持动态并行
sm_50 sm_52 sm_53	+Maxwell架构支持
sm_60 sm_61 sm_62	+Pascal架构支持
sm_70 sm_72	+Volta架构支持
sm_75	+Turing架构支持
sm_80 sm_86 sm_87	+Ampere架构支持

本节介绍了CUDA的 nvcc 编译引擎，并列举了 nvcc 一些基础和常用的编译选项。关于nvcc完整的介绍请参考官方指南NVIDIA CUDA Compiler Driver NVCC [3]。

一个简单的例子¶

代码示例¶

下面介绍一个简单的CUDA C程序例子，演示如何在瀚海22上编译运行CUDA代码。

例子展示的是两个向量相加的CUDA代码 add.cu 。

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>

__global__ void add(int n, float *x, float *y, float *z)
{
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if (i < n) z[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20;
  float *x, *y, *z, *d_x, *d_y, *d_z;
  x = (float*)malloc(N*sizeof(float));
  y = (float*)malloc(N*sizeof(float));
  z = (float*)malloc(N*sizeof(float));

  cudaMalloc(&d_x, N*sizeof(float));
  cudaMalloc(&d_y, N*sizeof(float));
  cudaMalloc(&d_z, N*sizeof(float));

  for (int i = 0; i < N; i++) {
        x[i] = rand()*1.0/RAND_MAX;
        y[i] = rand()*1.0/RAND_MAX;
  }

  cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);

  add<<<(N+255)/256, 256>>>(N, d_x, d_y, d_z);

  cudaMemcpy(z, d_z, N*sizeof(float), cudaMemcpyDeviceToHost);

  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
        maxError = max(maxError, fabs(z[i]-x[i]-y[i]));
  printf("Max error: %f\n", maxError);

  cudaFree(d_x);
  cudaFree(d_y);
  cudaFree(d_z);
  free(x);
  free(y);
  free(z);

  return 0;
}

函数 add 是在GPU上并行运行的内核， main 函数是宿主代码。

程序解读¶

main函数声明3个数组。

float *x, *y, *z, *d_x, *d_y, *d_z;
x = (float*)malloc(N*sizeof(float));
y = (float*)malloc(N*sizeof(float));
z = (float*)malloc(N*sizeof(float));

cudaMalloc(&d_x, N*sizeof(float));
cudaMalloc(&d_y, N*sizeof(float));
cudaMalloc(&d_z, N*sizeof(float));

指针x、y和z分别指向使用 malloc 分配的主机内存空间，d_x、d_y和d_z指针分别指向使用CUDA运行时API cudaMalloc 函数分配的设备存储空间。 CUDA中的主机和设备有独立的内存空间，这两个空间都可以从主机代码进行管理。

为了初始化设备数组，使用 cudaMemcpy 将数据从x和y复制到相应的设备数组d_x和d_y，它的工作方式与标准的C memcpy 函数一样，只是增加了第四个参数，指定拷贝的方向。在这里，我们使用 cudaMemcpyHostToDevice 指定第一个（目标）参数是设备指针，第二个（源）参数是主机指针。

cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);

add内核由以下语句启动：

add<<<(N+256-1)/256, 256>>>(N, d_x, d_y, d_z);

<<< 和 >>> 符号之间的信息是执行配置，指示有多少设备线程并行执行内核。在CUDA中，软件中有一个线程层次结构，它模仿线程处理器在GPU上的分组方式。执行配置中的第一个参数指定网格中线程块的数量，第二个参数指定线程块中的线程数。线程块和网格可以通过为这些参数传递dim3（一个由CUDA用x、y和z成员定义的简单结构）值来生成一维、二维或三维的线程块和网格。对于add这个示例，只需要一维线程组，所以我们只传递整数。在本例中，我们使用包含256个线程的线程块启动内核，并使用“上整计算”来确定处理数组全部N个元素所需的线程块数（(N+256-1)/256）。

由于数组的元素数有不能被线程块大小整除的可能，内核代码必须检查内存访问是否越界。

在运行内核之后，使用 cudaMemcpy （拷贝方向： cudaMemcpyDeviceToHost ），从d_z指向的设备数组复制到z指向的主机数组，将结果返回给主机。

cudaMemcpy(z, d_z, N*sizeof(float), cudaMemcpyDeviceToHost);

程序的最后，使用 cudaFree() 和 free() 分别清理设备端和主机端申请的内存。

cudaFree(d_x);
cudaFree(d_y);
cudaFree(d_z);
free(x);
free(y);
free(z);

编译运行¶

接下来以瀚海22超级计算系统使用Slurm作业调度系统的交互式任务为例，演示CUDA程序在超算系统上编译运行的一般方法。

首先在 登录节点 ，利用 module 命令载入合适的CUDA版本。本例中，在 登录节点 运行命令 module avail 确认系统中已安装的CUDA版本，选择符合自身需求的版本，如cuda/11.7.1_515.65.01装载（ module load cuda/x.x.x , x.x.x指版本号）：

（可选操作）使用 nvcc --versio 命令可以查看确认当前环境载入的CUDA版本：

接着，使用 nvcc 命令编译CUDA程序，编译选项的含义可参考上节“NVCC编译引擎”：

然后，使用 salloc 命令向Slurm系统申请交互式任务。 salloc 命令选项及Slurm系统可参考本手册的“Slurm作业调度系统” [4]：

至此，已完成在 登录节点 的工作（CUDA程序的编译和Slurm交互式任务的提交）。我们转入 计算节点 。
使用SSH命令免密登入 计算节点 （计算节点的主机名称由 salloc 命令的输出确定）。

备注

由于瀚海22系统的登录节点和计算节点的 /home 目录都挂载共享存储对应目录，所以两者的 /home 目录内容相同。

在 计算节点 使用 module load 载入与登录节点一致的CUDA版本。然后运行已编译好的程序，完成计算。

小结¶

本小节的简单例子用于展示瀚海22超级计算系统上CUDA程序的编译运行（交互式）的一般流程。

备注

使用 salloc （交互式）提交任务是为了向初学者展示登录节点和计算节点的区别，并不是瀚海22的常规作业调度方式。瀚海22超级计算系统上的正式作业请通过编写计算脚本并以 sbatch 提交的方式实现资源的申请和计算的执行。

Naive加法计算没有考虑如何充分利用GPU的计算资源和数据带宽。对性能优化感兴趣，希望充分利用GPU的计算资源的读者可以进一步阅读NVIDIA关于CUDA编程的进阶读物 [5]。

One More Thing¶

瀚海22超级计算系统系统预装了NVIDIA HPC SDK（nvhpc），方便用户充分利用GPU计算资源，实现科学计算应用和性能优化等工作。

NVIDIA HPC SDK是NVIDIA公司提供的一个包含编译器、函数库和软件工具的软件包，包含了用于方便用户开发、增强程序性能和可移植性的一系列工具。NVIDIA HPC SDK C、C++和Fortran编译器支持使用标准C++和Fortran、OpenACC指令和CUDA对HPC建模和模拟应用程序进行GPU加速。GPU加速的数学库使普通HPC算法的性能最大化，优化的通信库使基于标准的多GPU和可扩展系统编程成为可能。性能剖析和调试工具简化了HPC应用程序的移植和优化，而容器化工具则使得在企业内部或在云端的部署变得容易。HPC SDK支持NVIDIA GPU和运行Linux的Arm、OpenPOWER或x86-64 CPU，为用户提供了构建NVIDIA GPU加速的HPC应用程序所需的工具。

编译器命令¶

nvc ：是一个用于NVIDIA GPU和AMD、Intel、OpenPOWER和Arm CPU的C语言编译命令。支持ISO C11，支持用OpenACC进行GPU编程，并支持用OpenACC和OpenMP进行多核CPU编程。
nvc++ ：是一个针对NVIDIA GPU和AMD、Intel、OpenPOWER和Arm CPU的C++语言编译命令。它为目标处理器调用C++编译器、汇编器和链接器，选项来自其命令行参数。支持ISO C++17，支持使用C++17并行算法、OpenACC和OpenMP的GPU和多核CPU编程。
nvfortran ：是一个用于NVIDIA GPU和AMD、Intel、OpenPOWER和Arm CPU的Fortran编译命令。支持ISO Fortran 2003/2008的许多特性，支持使用CUDA Fortran进行GPU编程，以及使用ISO Fortran并行语言特性、OpenACC和OpenMP进行GPU和多核CPU编程。
nvcc ：是用于NVIDIA GPU的CUDA C和CUDA C++编译器驱动程序。接受一系列传统的编译器选项，例如用于定义宏和include/library路径，以及用于引导编译过程。为NVIDIA GPU生成优化代码，并驱动支持AMD、Intel、OpenPOWER和Arm CPU的主机编译命令。

数学库¶

cuBLAS：cuBLAS库提供了基本线性代数子程序（BLAS，Basic Linear Algebra Subprograms）的GPU加速实现。cuBLAS利用针对NVIDIA GPU高度优化的行业标准BLAS API加速AI和HPC应用。cuBLAS库包含用于分批操作、跨多个GPU执行以及混合和低精度执行的扩展。
cuTENSOR：cuTENSOR库是第一个由GPU加速的张量线性代数库，提供张量收缩、还原和元素运算。cuTENSOR用于加速深度学习训练和推理、计算机视觉、量子化学和计算物理等领域的应用。
cuSPARSE：cuSPARSE库为稀疏矩阵提供了GPU加速的基本线性代数子程序，其功能可用于建立GPU加速的求解器。cuSPARSE被从事机器学习、计算流体力学、地震勘探和计算科学等应用的工程师和科学家广泛使用。
cuSOLVER：cuSOLVER库提供了针对NVIDIA GPU高度优化的密集和稀疏因式分解、线性求解器和eigensolvers。cuSOLVER用于加速科学计算和数据科学等不同领域的应用，并拥有针对混合精度张量加速和跨多个GPU执行的扩展。
cuFFT：cuFFT库提供了针对NVIDIA GPU高度优化的快速傅里叶变换（FFT，Fast Fourier Transform）实现。cuFFT被用于建立跨学科的商业和研究应用，如深度学习、计算机视觉、计算物理、分子动力学、量子化学以及地震和医学成像，并具有跨多个GPU执行的扩展。
cuRAND：cuRAND库是一个随机数发生器的GPU设备端实现。

常用工具¶

CUDA-GDB：用于调试CUDA应用程序的NVIDIA工具软件。
Nsight Compute：NVIDIA Nsight Compute是面向CUDA应用程序的下一代交互式内核分析器。它通过一个用户界面和命令行工具提供详细的性能指标和API调试。
Nsight System：NVIDIA Nsight System是一款全系统性能分析工具，旨在实现应用程序算法的可视化。有助于识别优化和调整机会，以便在CPU和GPU上高效扩展应用程序。

NVIDIA HPC SDK高性能计算工具包的完整介绍请参考 [6]。