发行说明#
本节包括重大更改、新功能、性能改进和各种问题。除非另有说明,否则列出的问题不应影响功能。当功能受到影响时,我们会提供解决方法以避免该问题(如果可用)。
0.3.0#
cuBLASDx 库的第三个早期访问 (EA) 版本带来了对选定整数类型的支持、结果存储在寄存器中的矩阵乘法,以及计算类型与输入/输出类型的解耦。
新功能#
- 寄存器片段张量支持
添加了分区、谓词和转换工具。
添加了复制和分区实用程序,用于将数据从寄存器文件缓冲区移入/移出。
- 新的 GEMM 寄存器 API 允许更好的性能。
使用寄存器存储
C
矩阵(结果矩阵)可以通过在寄存器中执行累加和输入/输出,来节省共享内存大小和传输。
整数类型支持,包括对整数类型的 MMA 指令支持。
- 计算精度与输入精度解耦。
GEMM 函数现在可以接受任何输入并在寄存器中转换它,从而节省内存使用。
添加了更强大的共享内存管理工具,从而可以构建高效的流水线执行。
- 由于以下原因,该库不再静态断言 GEMM 问题是否适合共享内存
A
和B
可以互为别名或重叠,C
可以完全在寄存器文件中,以及输入精度可以是任意的,并且不取决于计算精度。
重大更改#
共享内存切片和共享内存大小实用程序不再作为
BLAS
方法提供。共享内存大小特性已从
BLAS
类型中删除。is_supported
特性已被删除,并添加了is_supported_smem_restrict
,is_supported_rmem_restrict
来代替它。该库不再断言大小是否适合共享内存,现在是用户的责任。
已解决问题#
cuBLASDx 内部版本控制定义已修复: *
CUBLASDX_VERSION_MAJOR
现在在CUBLASDX_VERSION
中定义为 10000 的倍数,而不是 1000,即CUBLASDX_VERSION_MAJOR = CUBLASDX_VERSION / 10000
。 *CUBLASDX_VERSION_MINOR
现在最多可以有两位数字。 * 由于没有主要版本发布,因此此版本与先前版本之间没有额外的间隔。 * 所有定义都可以在文件cublasdx_version.hpp
中检查。已添加缺失的特性以遵循
C++
元编程约定
已知问题#
- 建议使用最新的 CUDA 工具包和
NVCC
编译器。 CUDA 12.4
在 SM90 上使用寄存器 API 时,存在已知边缘情况,会导致生成不正确的 FP8 MMA 模拟代码。CUDA 12.1
存在已知边缘情况,当将先前命名的值传递给 3 值运算符时会崩溃。例如,
Alignment<varn_name_1, var_name_2, var_name_3>
可能会导致编译挂起,而Alignment<8, 8, 8>
始终有效。
- 建议使用最新的 CUDA 工具包和
0.2.0#
cuBLASDx 库的第二个早期访问 (EA) 版本带来了张量 API、混合精度和性能改进。
新功能#
所有设备扩展库都捆绑在一个名为 nvidia-mathdx-24.08.0.tar.gz 的软件包中。
添加了新的基于张量的 execute(…) API
由于支持 CuTe 张量 (cute::Tensor),矩阵的性能和用户友好界面得到改进。
用于在矩阵之间切片共享内存的辅助方法。
由于 get_layout_*() 方法和 cublasdx::make_tensor,可以轻松创建张量。
关于共享内存中矩阵的最佳布局的建议,以提高矩阵乘法和全局共享内存 I/O 操作的性能。
cublasdx::copy 用于复制共享和全局内存张量。
在 Precision 中添加了对混合精度的支持,特别是长期以来要求的
TensorFloat-32:
Precision<tfloat32_t, tfloat32_t, float>
,和Precision<__half, __half, float>
.
支持 8 位浮点矩阵(
__nv_fp8_e4m3
和__nv_fp8_e5m2
)和 GEMM。添加了 Alignment 算子以提供对齐信息。建议使用 16 字节(128 位)对齐以获得更好的性能。
TransposeMode 算子已弃用,并替换为 Arrangement 算子。
TransposeMode 算子(或替换它的 Arrangement)不再明确需要定义完整的 BLAS 执行。
A 矩阵的默认排列为
row_major
,B 矩阵为col_major
,C 矩阵为col_major
。A 矩阵的默认转置模式现在为
transposed
,B 矩阵为non-transposed
。
已知问题#
建议使用最新的 CUDA 工具包和
NVCC
编译器。
0.1.0#
cuBLASDx 库的第一个早期访问 (EA) 版本。
新功能#
已知问题#
由于 CUDA Toolkit 12.2 NVCC 编译器在某些情况下会报告不正确的编译错误,当使用 GEMM 描述类型的
value_type
类型时。下面介绍了有问题的代码以及可能的解决方法// Any GEMM description type using GEMM = decltype(Size<32 /* M */, 32 /* N */, 32 /* K */>() + Precision<double>() + Type<type::real>() + TransposeMode<t_mode /* A */, t_mode /* B */>() + Function<function::MM>() + SM<700>() + Block()); using type = typename GEMM::value_type; // compilation error // Workaround #1 using type = typename decltype(GEMM())::value_type; // Workaround #2 (used in cuBLASDx examples) template <typename T> using value_type_t = typename T::value_type; using type = value_type_t<GEMM>;