要求和功能#

要求#

cuBLASDx 是一个 CUDA C++ 仅头文件库。因此，使用该库所需的软件列表相对较小

CUDA Toolkit 11.4 或更新版本
支持的 CUDA 编译器（需要 C++17）
支持的主机编译器（需要 C++17）
（可选）CMake (版本 3.18 或更高版本)

依赖项

commonDx（随 MathDx 软件包一起提供）
CUTLASS 3.6.0 或更新版本（CUTLASS 3.6.0 随 MathDx 软件包一起提供）

支持的编译器#

CUDA 编译器

NVCC 11.4.152+ (CUDA Toolkit 11.4 或更新版本)
（实验性支持）NVRTC 11.4.152+ (CUDA Toolkit 11.4 或更新版本)

主机 / C++ 编译器

GCC 7+
Clang 9+（仅限 Linux/WSL2）
HPC SDK nvc++ 23.1+

注意

我们建议使用 GCC 9+ 作为主机编译器，以及最新 CUDA Toolkit 附带的 NVCC 作为 CUDA 编译器。

警告

尚未在 Windows 上使用 MSVC 测试编译 cuBLASDx，目前尚不支持。但是，可以使用 NVRTC 在 Windows 上编译带有 cuBLASDx 的内核，如示例所示。

注意

cuBLASDx 会为不支持的 C++ 标准版本发出错误，可以通过在编译期间定义 CUBLASDX_IGNORE_DEPRECATED_DIALECT 来消除这些错误。cuBLASDx 不保证在使用 cuBLASDx 中不支持的 C++ 标准版本时能够正常工作。

警告

由于已知的 GCC 问题，如果主机编译器是 GCC 11+，则 cuBLASDx 仅支持 NVCC 11.6+。此外，我们建议在 Volta 和 Turing 架构上运行内核时使用 NVCC 12.3+，以避免已知的编译器错误，该错误会导致某些用例的错误结果。

支持的功能#

这是 cuBLASDx 的早期访问 (EA) 版本。该库的当前功能是 cuBLASDx 在首次发布时将拥有的功能的子集。

支持的功能包括

创建运行 GEMM 的块描述符 - 通用矩阵乘法例程：\(\mathbf{C}_{m\times n} = {\alpha} \times \mathbf{A}_{m\times k} \times \mathbf{B}_{k\times n} + {\beta} \times \mathbf{C}_{m\times n}\)（参见函数运算符）。
自动使用 Tensor Cores 和自动数据布局以获得最佳内存访问模式。
使用寄存器片段或共享内存作为累加的输入/输出内存空间。
双向信息流，从用户到描述符通过运算符，以及从描述符到用户通过特性。
使用 SM 运算符定位特定 GPU 架构。这使用户能够使用建议的参数配置描述符以实现目标性能。

支持的内存空间#

cuBLASDx 支持由 m、n、k 维度定义的所有 GEMM 大小，这些维度可以容纳在寄存器文件 (RF) 和共享内存中。矩阵 A、B 必须适合共享内存才能执行计算。这些输入矩阵可能会重叠或互为别名。每个 CUDA 线程块的最大共享内存量可以在CUDA C 编程指南中找到。

输入/输出 C 矩阵可以是

在共享内存中提供（它不能与任何输入元素互为别名）用于 \(\mathbf{C}_{m\times n} = {\alpha} \times \mathbf{A}_{m\times k} \times \mathbf{B}_{k\times n} + {\beta} \times \mathbf{C}_{m\times n}\)
作为寄存器片段提供用于累加到 \(\mathbf{C}_{m\times n} = \mathbf{A}_{m\times k} \times \mathbf{B}_{k\times n} + \mathbf{C}_{m\times n}\)
通过值作为返回寄存器片段从 \(\mathbf{C}_{m\times n} = \mathbf{A}_{m\times k} \times \mathbf{B}_{k\times n}\) 返回

支持的计算类型#

cuBLASDx 支持 2 个域的计算

实数
复数，

在 7 种浮点精度中

半精度 (__half)
单精度 (float)
双精度 (double)
fp8_e4m3 (__nv_fp8_e4m3)
fp8_e5m2 (__nv_fp8_e5m2)
bf16 (__nv_bfloat16)
tf32 (cublasdx::tfloat32_t)

在 8 种整数精度中

有符号 8 位 (int8_t)
无符号 8 位 (uint8_t)
有符号 16 位 (int16_t)
无符号 16 位 (uint16_t)
有符号 32 位 (int32_t)
无符号 32 位 (uint32_t)
有符号 64 位 (int64_t)
无符号 64 位 (uint64_t)

从 cuBLASDx 0.2.0 开始，支持 A、B 和 C 不同精度的矩阵乘法。

支持任何 3 种精度组合，只要满足以下任一条件

它们都是浮点精度。
它们都是整数精度，并且
1. 累加器至少比任何输入宽 4 倍，
2. 输入符号性暗示累加器符号性。

不支持混合浮点/整数 GEMM，但可以将寄存器内输入转换应用于该效果。

支持的输入类型#

注意

从 cuBLASDx 0.3.0 开始，计算精度已与数据精度解耦，即，每个矩阵的输入/输出数据可以是任意类型（甚至浮点 GEMM 的整数输入），前提是设置了对齐运算符并且满足以下至少一个条件

它可以隐式转换为使用精度运算符和类型运算符选择的数据类型。
对于输入：提供适当的转换加载操作作为参数之一。它采用输入类型值。其结果必须至少可以隐式转换为计算类型。
对于输出：提供适当的转换存储操作作为参数之一。它采用结果计算类型（通常是精度运算符和类型运算符定义的 C 类型）。其结果必须至少可以隐式转换为输出类型。

支持的输入布局#

数据可以以 cublasdx::tensor（或底层 cute::Tensor）描述的任何布局提供，只要它是 2 维张量。其模式可以是分层的，但只需要有两个模式。

非重叠 A 和 B 的支持的最大尺寸#

下面您可以找到一个表格，其中列出了三种常用浮点精度（半精度、单精度和双精度）和类型（实数或复数）的最大支持尺寸，假设 m、n 和 k 维度相等，并且 A、B 和 C 的精度相同。

如果满足以下条件，则有效支持的维度会大得多

维度不相等（长而宽的矩阵）
A 和 B 是别名并共享元素（例如，A 与其转置相乘只需要拟合 A）

功能	类型，A/B/C 的精度	架构	最大尺寸
功能	类型，A/B/C 的精度	架构	受限 AB，C 在共享内存中	受限 AB，C 在寄存器中
GEMM	实数，半精度	70, 72	128	156
		75	104	127
		80, 87	166	203
		86, 89	129	157
		90	196	240
	实数，单精度复数，半精度	70, 72	90	110
		75	73	89
		80, 87	117	143
		86, 89	91	111
		90	139	170
	实数，双精度复数，单精度	70, 72	64	78
		75	52	63
		80, 87	83	101
		86, 89	64	78
		90	98	120
	复数，双精度	70, 72	45	55
		75	36	44
		80, 87	58	71
		86, 89	45	55
		90	69	84

警告

从 cuBLASDx 0.3.0 开始，对于选择的尺寸是否适合设备没有静态断言。这是允许输入重叠以及提供基于寄存器的累加 API 以及共享内存 API 的结果。

支持的 MMA 数据类型#

下表列出了 A、B 和 C 的精度，对于这些精度，可以使用专门的 Tensor Core 操作。

缩放的类型和精度，即 \({\alpha}\) 和 \({\beta}\)，预计与矩阵 C 的类型和精度相同，并且 A/B/C 必须是全实数或全复数数据类型。如果 IO 精度与计算精度解耦，则缩放类型必须与计算类型兼容。

精度 A	精度 B	精度 C	注意
fp8_e4m3	fp8_e4m3	float	MMA，SM89+
fp8_e4m3	fp8_e5m2	float	MMA，SM89+
fp8_e5m2	fp8_e5m2	float	MMA，SM89+
fp8_e5m2	fp8_e4m3	float	MMA，SM89+
半精度	半精度	半精度	MMA，SM70+
半精度	半精度	float	MMA，SM70+
bf16	bf16	float	MMA，SM80+
tf32	tf32	float	MMA，SM80+
双精度	双精度	双精度	MMA，SM80+
int8_t	int8_t	int32_t	MMA，SM75+
uint8_t	int8_t	int32_t	MMA，SM75+
int8_t	uint8_t	int32_t	MMA，SM75+
uint8_t	uint8_t	int32_t	MMA，SM75+

注意

如果 MMA 存在于 A、B 和 C 的精度组合中，则 cuBLASDx 将在支持的 GPU 架构上自动使用 MMA 指令。否则，cuBLASDx 将使用 FMA 指令，并且不保证性能。