GTC 大会火热进行中!别错过 NVIDIA CEO 黄仁勋的最新发布,以及 AI 和加速计算的必听会议。
高性能计算

在 CUDA 11 . 5 工具包中展示新功能

NVIDIA 宣布 CUDA 开发环境的最新版本 CUDA 11 . 5 。 CUDA 11 . 5 专注于增强您的 CUDA 应用程序的编程模型和性能。 CUDA 继续推动 GPU 加速的边界,并为 HPC 、可视化、 AI 、 ML 和 DL 中的新应用打下基础,和数据科学。

CUDA 11 . 5 有几个重要特性。这篇文章概述了关键功能:

  • CUDA 编程模型增强
    • 扫描合作小组中的集体
    • 标准化整数格式
    • 块压缩格式
    • C ++中可配置的缓存提示
  • MPS 增强功能(客户端内存限制)
  • WSL 上的 CUDA 驱动程序更新
  • CUDA Python GA
  • NVIDIA 开普勒驱动器的弃用
  • CUDA C ++(有关更多信息,请参阅使用 CUDA C ++编译辅助工具减少应用程序构建时间
  • NSight 计算/系统工具

CUDA 11 . 5 附带 R510 驱动程序,该驱动程序是一个长期支援科. CUDA 11 . 5 可供下载

CUDA 编程模型增强

此版本引入了关键的增强功能,以提高 CUDA 图形的可用性和性能,而无需对应用程序进行任何修改或任何其他用户干预。它还提高了多进程服务( MPS )的易用性。我们在 CUDA 编程指南中对异步编程模型进行了形式化。

扫描合作小组中的集体

reductions障碍,前缀和(也称为scans)一起,它们是并行计算的基石。扫描操作采用二进制运算符(通常为加法),并在输入数组上累积应用该运算符进行迭代。扫描可以是inclusive,包括所有元素x[0]…x[n],或在范围{0,x[0]…x[n-1]}上迭代的exclusive

例如,使用输入数组[3 1 7 0 4 1 6 3]的+运算符进行独占扫描将导致以下结果:

          [0 3 4 11 11 15 16 22]

CUDA 11 . 5 添加了一个新的头<cg/scan.h>,它在协作组名称空间中定义了四个新函数来实现这些操作。

Inclusive Scan Exclusive Scan Description
template<typename Group, typename T, typename OpType>
T inclusive_scan(const Group& g, T&& val, OpType&& op)
template<typename Group, typename T, typename OpType>
T exclusive_scan(const Group& g, T&& val, OpType&& op) ;
Perform scan using user supplied binary operator.
template<typename Group, typename T>
T inclusive_scan(const Group& g, T&& val);
template<typename Group, typenameT>
T exclusive_scan(const Group& g, T&& val);
Same as above with assumed op == plus<T>;
表 1 。扫描描述

在所有情况下,返回类型必须与输入值类型匹配。

规范化整数数据类型

标准化有符号和无符号 8 位和 16 位数据类型是 GPU 编程语言最广泛支持的一些纹理格式。 CUDA 一段时间以来一直支持将这些格式用于纹理对象,但在 11 . 5 版本中,我们扩展了对这些数据类型的现有支持,使与其他外部 API 的互操作更加直观。

我们在驱动程序和运行时 API 中引入了新的 CUDA 数组格式。驱动程序 API 公开了 12 种新的数组格式,如下所示:

CU_AD_FORMAT_UNORM_INT8X{1|2|4} 
CU_AD_FORMAT_UNORM_INT16X{1|2|4}
CU_AD_FORMAT_SNORM_INT8X{1|2|4}
CU_AD_FORMAT_SNORM_INT16X{1|2|4}

这些可用于创建 1 、 2 或 4 通道 CUDA 阵列。运行时 API 同样公开了 12 种新的等效通道格式:

cudaChannelFormatKindUnsignedNormalized8X{1|2|4}
cudaChannelFormatKindSignedNormalized8X{1|2|4}
cudaChannelFormatKindUnsignedNormalized16X{1|2|4}
cudaChannelFormatKindSignedNormalized16X{1|2|4}

这些还可用于创建 1 、 2 或 4 通道、 8 位或 16 位通道宽度 CUDA 阵列。此外,您现在可以从外部 API (如 DirectX12 / 11 或 Vulkan )导入匹配的格式化纹理,并将其映射为 CUDA 数组。使用资源视图创建纹理对象时,格式 texel 大小必须与数组 texel 大小匹配。

对于纹理对象,可以创建和访问它们,如以下代码示例所示:

cudaArray_t array;
cudaChannelFormatDesc formatDesc = {8, 8, 0, 0, cudaChannelFormatKindUnsignedNormalized8X2};
 
cudaMallocArray(&array, formatDesc, width, height);
 
cudaTextureDesc texDesc = {0};
texDesc.addressMode[0] = texDesc.addressMode[1] = cudaAddressModeClamp ;
 
// (3) Create CUDA texture object
cudaResourceDesc resDesc = {0};
resDesc.resType = cudaResourceTypeArray;
resDesc.res.array = array;
cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);

// Read from texture object in a kernel as follows:
// float4 texel = tex2D<float4>(texture, x, y);
 
// (6) Release all resources
cudaDestroyTextureObject(texObj); 
cudaFreeArray(array);

同样,对于曲面对象:

cudaArray_t array;
cudaChannelFormatDesc formatDesc = {16, 16, 16, 16, cudaChannelFormatKindSignedNormalized16X4};
 
cudaMallocArray(&array, formatDesc, width, height);
 
// (3) Create CUDA surface object
cudaResourceDesc resDesc = {0};
resDesc.resType = cudaResourceTypeArray;
resDesc.res.array = array;
cudaCreateSurfaceObject(&surfObj, &resDesc);
 
// Read/Write to/from surface object in a kernel as follows
// Read:
// short4 texel = surf2DRead(surface, xInBytes, y);
// Unformatted stores:
// surf2DWrite<short4>(texel, surface, xInBytes, y);
// Formatted stores: (Formatted surface stores are currently not exposed in CUDA runtime device functions)
// sust.p.2d.v4.b32
 
// (6) Release all resources
cudaDestroySurfaceObject(surfObj); 
cudaFreeMipmappedArray(array);

对块压缩数据类型的支持

在所有图形编程语言和框架中,用于减小纹理大小的最常见有损压缩技术之一是使用块压缩( BC )纹理格式。使用这些格式可以显著节省纹理的内存占用。有几种 BC 格式,每种格式都有其独特的优点和缺点,通常称为 BCn 格式。

NVIDIA GPU 体系结构本机支持 BCn 格式,并且通过纹理资源视图在 CUDA 中的支持有限。现在,我们在驱动程序和运行时 API 中引入新的 BC CUDA 数组格式。

Driver API Format Runtime API Format Channel Count Channel Width in Bits
CU_AD_FORMAT_BC1_UNORM
CU_AD_FORMAT_BC1_UNORM_SRGB
cudaChannelFormatKindUnsignedBlockCompressed1
cudaChannelFormatKindUnsignedBlockCompressed1SRGB
4 [8,8,8,8]
CU_AD_FORMAT_BC2_UNORM
CU_AD_FORMAT_BC2_UNORM_SRGB
cudaChannelFormatKindUnsignedBlockCompressed2
cudaChannelFormatKindUnsignedBlockCompressed2SRGB
4 [8,8,8,8]
CU_AD_FORMAT_BC3_UNORM
CU_AD_FORMAT_BC3_UNORM_SRGB
cudaChannelFormatKindUnsignedBlockCompressed3
cudaChannelFormatKindUnsignedBlockCompressed3SRGB
4 [8,8,8,8]
CU_AD_FORMAT_BC4_UNORM
CU_AD_FORMAT_BC4_SNORM
cudaChannelFormatKindUnsignedBlockCompressed4
cudaChannelFormatKindSignedBlockCompressed4
1 [8,0,0,0]
CU_AD_FORMAT_BC5_UNORM
CU_AD_FORMAT_BC5_SNORM
cudaChannelFormatKindUnsignedBlockCompressed5
cudaChannelFormatKindSignedBlockCompressed5
2 [8,8,0,0]
CU_AD_FORMAT_BC6H_UF16
CU_AD_FORMAT_BC6H_SF16
cudaChannelFormatKindUnsignedBlockCompressed6H
cudaChannelFormatKindSignedBlockCompressed6H
3 [16,16,16,0]
CU_AD_FORMAT_BC7_UNORM
CU_AD_FORMAT_BC7_UNORM_SRGB
cudaChannelFormatKindUnsignedBlockCompressed7
cudaChannelFormatKindUnsignedBlockCompressed7SRGB
4 [8,8,8,8]
表 2 。 BC 数据类型格式

这些格式可用于使用cuArray[3D]Create运行时 API 或cuArray[3D]Create驱动程序 API 创建 BCn 格式的cuMipmappedArrayCreate数组。类似地,可以使用 CUDA 运行时 API 或cuMipmappedArrayCreate驱动程序 API 创建 CUDA mipmapped 数组。使用这些格式创建 CUDA 阵列时,阵列范围必须是压缩块大小的倍数( 2D 为 4×4 , 3D 为 4x4x1 )。这些阵列还可用于创建纹理对象。

C ++中可配置的缓存提示

在 CUDA 11 . 4 中发现新功能中,我们引入了一个 PTX ISA 扩展,为驻留在 GPU 上的数据向编译器和运行时提供缓存提示。在 CUDA 11 . 5 中,我们将此能力扩展到带有注释指针的 C ++。它们充当普通指针,并应用附加属性。

注释指针是使用<cuda/annotated_ptr>中定义的函数创建的,缓存驻留提示定义为以下内容之一:

cuda::access_property::normal		(evict_normal_demote)
cuda::access_property::streaming        (evict_first)
cuda::access_property::persisting       (evict_last)
cuda::access_property::shared           (shared memory)
cuda::access_property::global           (evict_normal)

例如,在内核代码中,声明和使用带注释的指针可能类似于以下代码:

static __device__
void my_kernel(int * in, int * out) { 
cuda::access_property ap(cuda::access_property::persisting{}); 

// Retrieve global id 
int i = blockIdx.x * blockDim.x + threadIdx.x; 

cuda::annotated_ptr<int, cuda::access_property> in_ann{in, ap}; 
cuda::annotated_ptr<int, cuda::access_property> out_ann{out, ap};
 
...
}

MPS 增强功能(客户端内存限制)

当 GPU 的计算能力超过任何单个应用程序时,运行共享同一 GPU 硬件的多个应用程序进程可能会很有吸引力。多进程服务( MPS )运行时体系结构控制多个独立进程同时使用单个 GPU 。

但是,当多个独立进程共享 GPU 时,设置总体内存分配限制通常很有用,以避免任何单个进程占用过多的可用 GPU 内存。

在 CUDA 11 . 5 中,我们引入了一组新的控制机制,使您能够限制 MPS 客户端进程的固定内存分配。您可以通过默认的全局限制层次结构控制内存分配。

默认全局限制

可以使用设备的set_default_device_pinned_mem_limit控制命令显式启用默认全局内存限制。设置此命令将在将来生成的所有 MPS 服务器的所有 MPS 客户端上强制执行设备固定内存限制。

$nvidia-cuda-mps-control set_default_device_pinned_mem_limit 0 2G

每服务器限制:对于内存资源限制的细粒度控制,可以使用set_device_pinned_mem_limit控制命令在特定 MPS 服务器上选择性地设置限制。设置此命令将在特定 MPS 服务器的所有 MPS 客户端上强制执行设备固定内存限制。

$nvidia-cuda-mps-control set_device_pinned_mem_limit <pid> 1 1G

每个客户限额:前面两种控制机制为特定 MPS 服务器的所有 MPS 客户端设置了一个总体限制。希望更好地控制资源限制的用户;也就是说,在每个 MPS 客户端的基础上,可以通过为每个客户端进程分别设置CUDA_MPS_PINNED_DEVICE_MEM_LIMIT环境变量来实现。

此环境变量与 CUDA \ u VISIBLE \ u 设备具有相同的语义。值字符串可以包含逗号分隔的设备序号和设备 UUID ,每个设备的内存限制由等号(=)分隔。

$export CUDA_MPS_PINNED_DEVICE_MEM_LIMIT="0=1G,1=2G,GPU-7ce23cd8-5c91-34a1-9e1b-28bd1419ce90=1024M"

WSL 驱动程序上的 CUDA

用于英特尔 x86 体系结构的 NVIDIA Windows GPU 驱动程序将支持 WSL2 ,并可在 Windows 11 的 Windows Insider Preview ( WIP )程序之外访问。有关 WSL 支持详细信息的 CUDA ,请参阅《 WSL 用户指南》上的 CUDA 的支持矩阵和限制部分。

CUDA Python

Python CUDA 为现有工具包和库的驱动程序和运行时 API 提供Cython绑定和 Python 包装,以简化基于 GPU 的加速处理。 Python 是科学、工程、数据分析和深度学习应用程序中最流行的编程语言之一。 CUDA Python 的目标是将 Python 生态系统统一为一组接口,提供 Python 对 CUDA 主机 API 的全面覆盖和访问。

库开发人员可以直接从 Python 使用 CUDA Python 与 CUDA 的低级接口。我们很高兴地宣布,从 11 . 5 版开始, CUDA Python 已普遍提供,可以使用 PIP 或 Conda 安装。 CUDA 支持的所有平台都支持该库。

其他增强功能

除了前面讨论的关键功能外, CUDA 11 . 5 中还有一些增强功能,有助于改进基线功能和可用性。

多线程提交吞吐量

在 11 . 5 中,我们减少了 CPU 线程之间 CUDA API 的序列化开销。默认情况下,将启用这些更改。但是,为了帮助对潜在更改可能导致的问题进行分类,我们提供了一个环境变量CUDA_REDUCE_API_SERIALIZATION,以关闭这些更改。这是前面讨论的基础更改之一,有助于 CUDA 图的性能改进。

NVIDIA 开普勒驱动器的弃用

NVIDIA 开普勒微体系结构于 2012 年首次引入,并已逐步淘汰。对于所有基于开普勒 NVIDIA 的 SKU ,我们已经从 R495 版本开始取消了对驱动程序的支持。但是, CUDA 工具包开发工具和对 NVIDIA 开普勒 SKU 的支持将在未来的 CUDA 11 . x 版本中继续。

C ++语言对 CUDA 的支持

作为本版本的一部分,有一些 CUDA 11 . 5 支持的关键 C ++语言增强。

  • CUDA C ++编译器支持 NVRTC 和 PTX 中的并发编译,以提高编译时间。编译器现在还可以检测并消除未使用的 CUDA 内核,通过更好的代码优化减少编译时间、二进制大小和总体性能。
  • 对 128 位整数值的有限支持作为用户反馈的预览功能发布,同时引入了 NVRTC 的静态库版本,并扩展了主机编译器支持以包括 Clang 12 . 0 。
  • CUDA C++编译器具有我们在[VZX57 ]帖子中深入研究的特性。
  • NVIDIA C ++标准库( LIbCu +++) 1 . 5 . 0 被 CUDA 11 . 4 发布。
  • 推力 1 . 12 . 0 具有新的推力::通用_矢量 API ,使您能够使用推力的 CUDA 统一内存。

NSight 开发人员工具

NVIDIA NSight 开发工具: NSight System 2021 . 4 、 NSight Compute 2021 . 3 和 NSight Graphics 2021 . 4 . 2 现在提供了新版本,用于通过分析和调试 CUDA 代码提高性能。

新发布的 NSight 系统 2021 . 4改进了对 Windows 、 Direct3D12 和 Vulkan 支持的评测。此版本添加了一些功能,以帮助更好地理解操作系统中断时的进程执行,并添加了数据捕获以识别数据包队列瓶颈。特色包括:

  • Windows ISR 和 DPC 跟踪
  • GPU 基于硬件的调度跟踪
  • Windows Direct3D12
  • Vulkan 与 WDDM 事件的相关性
  • NVTX 事件分类支持
  • 多种系统环境的多报告加载

有关更多信息,请参阅下载中心

NSight Compute 2021 . 3 添加了一些功能,以帮助用户了解其 CUDA 内核的性能。新的占用率计算器活动对 CUDA 内核的资源利用率进行建模,以便您可以交互式地调整模型参数,以查看它们如何影响占用率。屋顶折线图现在支持分层脱机,它表示除设备内存之外的内存层次中的其他级别。您可以查看内核是否存在与缓存访问请求相关的瓶颈。

还有其他改进,包括更多可配置的基线比较、从 CLI 查看源代码级信息以及其他 SSH 功能。有关 NSight 开发工具新功能的更多信息,请参阅发行说明download页面。

Screen shot of Nsight Compute software Occupancy Calculator feature showing GPU memory occupancy levels and areas for improvement.
图 1 。占用率计算器

最新的 NSight 图形 2021 . 4 . 2 现在包括对 Windows 11 的支持。这意味着您现在可以使用download Direct3D 和 Vulkan 的 NVIDIA 图形调试器和探查器,在最前沿的 Windows 版本上创建惊人的 3D 图形。

有关更多信息,请参阅以下参考资料:

在 GTC 上,我们在几个会议中免费向 NSight 开发人员工具展示了高级优化功能。注册

有关 CUDA 11 系列工具包功能以及现有功能概述的更多信息,请参阅 CUDA 11 项特色展示文章和过去的 CUDA – 相关职位

 

Tags