发行说明#

本节包括重大更改、新功能、性能改进和各种问题。除非另有说明,否则列出的问题不应影响功能。当功能受到影响时,我们会提供解决方法以避免该问题(如果可用)。

0.3.0#

cuBLASDx 库的第三个早期访问 (EA) 版本带来了对选定整数类型的支持、结果存储在寄存器中的矩阵乘法,以及计算类型与输入/输出类型的解耦。

新功能#

  • 寄存器片段张量支持
    • 添加了分区、谓词和转换工具。

    • 添加了复制和分区实用程序,用于将数据从寄存器文件缓冲区移入/移出。

  • 新的 GEMM 寄存器 API 允许更好的性能。
    • 使用寄存器存储 C 矩阵(结果矩阵)可以通过在寄存器中执行累加和输入/输出,来节省共享内存大小和传输。

  • 整数类型支持,包括对整数类型的 MMA 指令支持。

  • 计算精度与输入精度解耦。
    • GEMM 函数现在可以接受任何输入并在寄存器中转换它,从而节省内存使用。

  • 添加了更强大的共享内存管理工具,从而可以构建高效的流水线执行。

  • 由于以下原因,该库不再静态断言 GEMM 问题是否适合共享内存
    • AB 可以互为别名或重叠,

    • 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> 始终有效。

0.2.0#

cuBLASDx 库的第二个早期访问 (EA) 版本带来了张量 API、混合精度和性能改进。

新功能#

  • 所有设备扩展库都捆绑在一个名为 nvidia-mathdx-24.08.0.tar.gz 的软件包中。

  • 添加了新的基于张量的 execute(…) API

  • 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

已知问题#

0.1.0#

cuBLASDx 库的第一个早期访问 (EA) 版本。

新功能#

  • 支持通用矩阵乘法

    • 张量核心支持 fp16、fp64、复数 fp64 计算。

  • 支持 SM70 - SM90 CUDA 架构。

  • 包含多个示例

已知问题#

  • 由于 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>;