特性

特性为用户提供关于使用 算子 构建的函数描述符的信息。它们分为 描述特性 和 执行特性。

---

描述特性

特性	描述
size_of<Description>	我们打算解决的问题规模。
type_of<Description>	使用的数据类型，可以是 type::real 或 type::complex。
precision_of<Description>	计算用于 A、B 和 C 的底层浮点值的精度。
function_of<Description>	要执行的函数。
arrangement_of<Description>	矩阵的排列模式 - arrangement::col_major 和 arrangement::row_major。
transpose_mode_of<Description>	矩阵的转置模式 - transpose_mode::non_transposed、transpose_mode::transposed 和 transpose_mode::conj_transposed。
alignment_of<Description>	矩阵 A、B 和 C 的对齐（以字节为单位）。
leading_dimension_of<Description>	提供由 Size 以及 Arrangement 或 TransposeMode 算子定义的引导维度。
sm_of<Description>	底层计算的目标架构。
is_blas<Description>	如果 Description 是使用 描述算子 形成的函数描述，则为 true。
is_blas_execution<Description>	如果 Description 是配置为使用 执行算子 执行的函数描述，则为 true。
is_complete_blas<Description>	如果 Description 是使用 描述算子 形成的有效且完整的函数描述，则为 true。
is_complete_blas_execution<Description>	如果 is_complete_blas<Description> 为 true 且 is_blas_execution<Description> 为 true，则为 true。

描述特性可以使用提供的辅助函数从函数描述符中检索。例如

#include <cublasdx.hpp>

using GEMM = decltype(cublasdx::Size<32, 32, 64>()
              + cublasdx::Precision<float, float, double>()
              + cublasdx::Type<cublasdx::type::real>()
              + cublasdx::TransposeMode<cublasdx::T, cublasdx::N>()
              + cublasdx::Function<cublasdx::function::MM>()
              + cublasdx::SM<700>()
              + cublasdx::Block());

if(cublasdx::is_complete_blas<GEMM>::value)
  std::cout << "GEMM (M x N x K): "
    << cublasdx::size_of<GEMM>::m << " x "
    << cublasdx::size_of<GEMM>::n << " x "
    << cublasdx::size_of<GEMM>::k << std::endl;

尺寸特性

// std::tuple<unsigned int, unsigned int, unsigned int>
cublasdx::size_of<BLAS>::value
cublasdx::size_of_v<BLAS>

// unsigned int
cublasdx::size_of<BLAS>::m
cublasdx::size_of_v_m<BLAS>

cublasdx::size_of<BLAS>::n
cublasdx::size_of_v_n<BLAS>

cublasdx::size_of<BLAS>::k
cublasdx::size_of_v_k<BLAS>

size_of 特性给出我们想要解决的问题的规模，由 尺寸算子 设置。如果描述符不是使用 尺寸算子 创建的，则编译将失败并显示错误消息。

类型特性

// cublasdx::type
cublasdx::type_of<BLAS>::value
cublasdx::type_of_v<BLAS>

函数中使用的数据类型（cublasdx::type::real 或 cublasdx::type::complex），由 类型算子 设置。

精度特性

// Precision type
cublasdx::precision_of<BLAS>::a_type
cublasdx::precision_of_a_t<BLAS>

cublasdx::precision_of<BLAS>::b_type
cublasdx::precision_of_b_t<BLAS>

cublasdx::precision_of<BLAS>::c_type
cublasdx::precision_of_c_t<BLAS>

计算输入数据 A 和 B 以及输出数据 C 的计算精度，由 精度算子 设置。

从 cuBLASDx 0.3.0 开始，输入精度已与计算精度解耦。有关更多详细信息，请参阅 精度算子。

函数特性

// cublasdx::function
cublasdx::function_of<BLAS>::value
cublasdx::function_of_v<BLAS>

要执行的函数，由 函数算子 设置。如果描述符不是使用 函数算子 创建的，则编译将失败并显示错误消息。

排列特性

// std::tuple<arrangement, arrangement, arrangement>
cublasdx::arrangement_of<BLAS>::value
cublasdx::arrangement_of_v<BLAS>

// cublasdx::arrangement
cublasdx::arrangement_of<BLAS>::a
cublasdx::arrangement_of_v_a<BLAS>

cublasdx::arrangement_of<BLAS>::b
cublasdx::arrangement_of_v_b<BLAS>

cublasdx::arrangement_of<BLAS>::c
cublasdx::arrangement_of_v_c<BLAS>

全局矩阵 A、B 和 C 的排列。

转置模式特性

警告

transpose_mode_of 特性自 0.2.0 版本起已弃用，并可能在未来版本中移除。

// cublasdx::transpose_mode
cublasdx::transpose_mode_of<BLAS>::a_transpose_mode
cublasdx::transpose_mode_of_a<BLAS>

cublasdx::transpose_mode_of<BLAS>::b_transpose_mode
cublasdx::transpose_mode_of_b<BLAS>

矩阵 A 和 B 的转置模式。

对齐特性

// std::tuple<unsigned int, unsigned int, unsigned int>
cublasdx::alignment_of<BLAS>::value
cublasdx::alignment_of_v<BLAS>

// unsigned int
cublasdx::alignment_of<BLAS>::a
cublasdx::alignment_of_v_a<BLAS>

cublasdx::alignment_of<BLAS>::b
cublasdx::alignment_of_v_b<BLAS>

cublasdx::alignment_of<BLAS>::c
cublasdx::alignment_of_v_c<BLAS>

输入矩阵 A、B 和 C 的对齐（以字节为单位）。

leading_dimension_of

// std::tuple<unsigned int, unsigned int, unsigned int>
cublasdx::leading_dimension_of<BLAS>::value
cublasdx::leading_dimension_of_v<BLAS>

// unsigned int
cublasdx::leading_dimension_of_v_a<BLAS>
cublasdx::leading_dimension_of_v_b<BLAS>
cublasdx::leading_dimension_of_v_c<BLAS>

类型 cublasdx::leading_dimension_of 为矩阵 A、B 和 C 提供引导维度。返回的值由 尺寸 以及 排列 或 转置模式 算子定义。

要求

• BLAS 必须定义尺寸、函数和转置模式。请参阅 描述算子 部分。

• BLAS 必须包含 Block 算子。

SM 特性

// unsigned int
cublasdx::sm_of<BLAS>::value
cublasdx::sm_of_v<BLAS>

用于运行函数的 GPU 架构。例如，Volta (sm_70) 的值为 700。

is_blas 特性

// bool
cublasdx::is_blas<BLAS>::value
cublasdx::is_blas_v<BLAS>

如果描述符是使用 描述算子 形成的函数描述，则特性为 true。

is_blas_execution 特性

// bool
cublasdx::is_blas_execution<BLAS>::value
cublasdx::is_blas_execution_v<BLAS>

如果描述符是配置为执行的函数描述，并使用 描述算子 和 执行算子 形成，则特性为 true。

is_complete_blas 特性

// bool
cublasdx::is_complete_blas<BLAS>::value
cublasdx::is_complete_blas_v<BLAS>

如果描述符是使用 描述算子 形成的完整函数描述，则特性为 true。

注意

在此上下文中，“完整”意味着描述符已使用所有必要的 描述算子 形成，并且仅缺少 执行算子 即可运行。

为了使函数描述符完整，需要满足以下条件

• 一个且仅一个 尺寸算子。

• 一个且仅一个 函数算子。

• 一个且仅一个 SM 算子。

is_complete_blas_execution 特性

// bool
cublasdx::is_complete_blas_execution<BLAS>::value
cublasdx::is_complete_blas_execution_v<BLAS>

如果 is_blas_execution 特性 和 is_complete_blas 特性 均为 true，则特性为 true。

注意

如果描述符 descriptor 的 is_complete_blas_execution 特性 为 true，那么我们可以使用 执行方法 来执行该函数。

执行特性

执行特性可以直接从已使用 执行算子 配置的 BLAS 描述符中检索。可用的执行特性可能取决于用于构建描述符的算子。目前，Block 算子 是唯一可用的执行算子。

Block 特性

特性	默认值	描述
Description::<a/b/c>_value_type	float, float, float	计算数据 A、B、C 的类型。
Description::<a/b/c>_dim	由问题规模和转置模式确定。	由问题规模和转置模式确定的矩阵 A、B、C 的逻辑维度。
Description::ld<a/b/c>	由问题规模确定或通过 LeadingDimension 算子 设置	矩阵 A、B、C 的引导维度。
Description::<a/b/c>_alignment	由矩阵 A、B、C 的类型确定或通过 Alignment 算子 设置	输入矩阵 A、B、C 的对齐（以字节为单位）。
Description::<a/b/c>_size	由问题规模和设置的引导维度确定。	矩阵 A、B、C 中的元素数量。包括由设置的引导维度确定的填充。
Description::block_dim	基于启发式或通过 BlockDim 算子 设置	类型为 dim3 的值，表示计算 BLAS 函数的线程数。设置或推荐的 CUDA 块维度。
Description::suggested_block_dim	基于启发式。	用于计算 BLAS 函数的推荐线程数，表示为 CUDA 块维度（dim3 值）。
Description::max_threads_per_block	X * Y * Z 其中 X, Y, Z = Description::block_dim	Description::block_dim 中的线程数。

Block 特性可以从使用 Block 算子 构建的描述符中检索。

例如

#include <cublasdx.hpp>

using BLAS = decltype(... + Block());

template<class BLAS>
__launch_bounds__(BLAS::max_threads_per_block)
__global__ void kernel(typename BLAS::a_value_type* input_a, ...) {
  ...
}

值类型特性

BLAS::a_value_type
BLAS::b_value_type
BLAS::c_value_type

矩阵 A、B 和 C 类型的计算类型。请参阅 支持的 MMA 数据类型。

由 精度 和 类型 确定。

默认类型为 float。它可以是实数或复数类型，具体取决于设置的 类型 算子。

注意

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

1. 它可以隐式转换为使用 精度算子 和 类型算子 选择的数据类型。

2. 对于输入：适当的转换加载操作作为参数之一提供。它接受输入类型值。其结果必须至少可以隐式转换为计算类型。

3. 对于输出：适当的转换存储操作作为参数之一提供。它接受结果计算类型（通常是 C 类型，由 精度算子 和 类型算子 定义）。其结果必须至少可以隐式转换为输出类型。

下表列出了矩阵 A、B 和 C 的值类型，具体取决于在 精度 算子中设置的相应精度和通过 类型 算子设置的类型。有关支持的精度的完整列表，请参阅这些地方。

标记为“替代项”的列表示可以安全地代替 <a/b/c>_value_type 类型使用的类型。只要其他类型的对齐和大小与 <a/b/c>_value_type 相同，也可以使用其他类型。

精度 \ 类型	实数
	<a/b/c>_value_type	替代项#1
__nv_fp8_e5m2	__nv_fp8_e5m2	cublasdx::float_e5m2_t
__nv_fp8_e4m3	__nv_fp8_e4m3	cublasdx::float_e4m3_t
__half	__half	cublasdx::half_t
__nv_bfloat16	__nv_bfloat16	cublasdx::bfloat16_t
cublasdx::tfloat32_t	cublasdx::tfloat32_t	cublasdx::tfloat32_t
其他	其他	不适用

精度 \ 类型	复数
	<a/b/c>_value_type	替代项#1	替代项#2
__nv_fp8_e5m2	cublasdx::complex<__nv_fp8_e5m2>	cublasdx::complex<cublasdx::float_e5m2_t>	__nv_fp8x2_e5m2
__nv_fp8_e4m3	cublasdx::complex<__nv_fp8_e4m3>	cublasdx::complex<cublasdx::float_e4m3_t>	__nv_fp8x2_e4m3
__half	cublasdx::complex<__half>	cublasdx::complex<cublasdx::half_t>	__half2
__nv_bfloat16	cublasdx::complex<__nv_bfloat16>	cublasdx::complex<cublasdx::bfloat16_t>	__nv_bfloat162
其他	cublasdx::complex<other>	不适用	不适用

注意

类型 cublasdx::float_e5m2_t、cublasdx::float_e4m3_t、cublasdx::half_t、cublasdx::bfloat16_t 和 cublasdx::tfloat32_t 是 CUTLASS 类型 cutlass::float_e5m2_t、cutlass::float_e4m3_t、cutlass::half_t、cutlass::bfloat16_t 和 cutlass::tfloat32_t 的别名。

矩阵维度特性

警告

矩阵维度特性自 0.2.0 版本起已弃用，并可能在未来版本中移除。

// tuple<unsigned int, unsigned int>
BLAS::a_dim
BLAS::b_dim
BLAS::c_dim

矩阵 A、B、C 的逻辑维度，以 (rows, columns) 元组的形式表示。维度由问题规模、排列和转置模式确定。

请参阅 GEMM、尺寸算子、排列算子 和 转置模式算子。

引导维度特性

// unsigned int
BLAS::lda
BLAS::ldb
BLAS::ldc

矩阵 A、B、C 的引导维度。

请参阅 GEMM、尺寸算子 和 LeadingDimension 算子。

矩阵尺寸特性

// unsigned int
BLAS::a_size
BLAS::b_size
BLAS::c_size

矩阵 A、B、C 中的元素数量。它包括由设置的引导维度确定的填充。

请参阅 GEMM、尺寸算子 和 LeadingDimension 算子。

对齐特性

// unsigned int
BLAS::a_alignment
BLAS::b_alignment
BLAS::c_alignment

矩阵 A、B、C 的对齐（以字节为单位）。请参阅 对齐算子。

Block 维度特性

// dim3
BLAS::block_dim

类型为 dim3 的值，表示将用于执行请求的 BLAS 函数的线程数。它等于指定的（或默认的，如果未指定）CUDA 块维度。请参阅 BlockDim 算子。

如果在 BLAS 描述中未使用 BlockDim，则 BLAS::block_dim 的默认值等于 BLAS::suggested_block_dim。

建议 Block 维度特性

// dim3
BLAS::suggested_block_dim

BLAS 描述的推荐线程数。

每个 Block 的最大线程数特性

BLAS::max_threads_per_block

线程数是推荐的或设置的块维度，X * Y * Z 其中 X, Y, Z = Description::block_dim。

其他特性

特性	描述
is_supported_smem_restrict<Description, Arch>	验证在提供的 CUDA 架构 (Architecture) 上，给定的 BLAS 操作是否受共享内存 API 支持，且不使用矩阵别名。
is_supported_rmem_restrict<Description, Arch>	验证在提供的 CUDA 架构 (Architecture) 上，给定的 BLAS 操作是否受寄存器 API 支持，且不使用矩阵别名。
suggested_leading_dimension_of<Description, Arch>	提供可能提高 BLAS 操作性能的引导维度。
suggested_alignment_of<Description>	矩阵 A、B 和 C 的建议对齐（以字节为单位）。

is_supported_smem_restrict

namespace cublasdx {
  // BLAS - BLAS description without CUDA architecture defined using SM operator
  // Architecture - unsigned integer representing CUDA architecture (SM)
  template<class BLAS, unsigned int Architecture>
  struct is_supported_smem_restrict : std::bool_constant<...> { };

  // Helper variable template
  template<class BLAS, unsigned int Architecture>
  inline constexpr bool is_supported_smem_restrict_v<BLAS, Architecture> = is_supported_smem_restrict<BLAS, Architecture>::value;
}

// true if BLAS is supported on the provided CUDA architecture
cublasdx::is_supported_smem_restrict<BLAS, Architecture>::value;

cublasdx::is_supported_smem_restrict 检查 BLAS 操作是否在 Architecture CUDA 架构上受支持，假设

1. 使用共享内存 API（所有矩阵 A、B 和 C 都必须适合共享内存）。

2. 输入精度与计算精度相同，并且

3. 矩阵不使用别名或相互重叠其范围，

   • 即，矩阵 B 的第一个元素在矩阵 A 的最后一个元素之后开始，等等。

// true if BLAS is supported on the provided CUDA architecture
cublasdx::is_supported_smem_restrict<BLAS, Architecture>::value;

要求

• BLAS 必须定义尺寸、函数和转置模式。请参阅 描述算子 部分。

• BLAS 必须包含 Block 算子。

• BLAS 不能通过 SM 算子定义目标 CUDA 架构。

示例

using namespace cublasdx;

using BLAS = decltype(Size<128, 128, 128>() + Type<type::real>() + Block() + Precision<float>());
cublasdx::is_supported_smem_restrict<BLAS, 900>::value; // true
cublasdx::is_supported_smem_restrict<BLAS, 800>::value; // false
cublasdx::is_supported_smem_restrict<BLAS, 700>::value; // false

using BLAS = decltype(Size<96, 96, 96>() + Type<type::real>() + Block() + Precision<float>());
cublasdx::is_supported_smem_restrict<BLAS, 900>::value; // true
cublasdx::is_supported_smem_restrict<BLAS, 800>::value; // true
cublasdx::is_supported_smem_restrict<BLAS, 700>::value; // false

is_supported_rmem_restrict

namespace cublasdx {
  // BLAS - BLAS description without CUDA architecture defined using SM operator
  // Architecture - unsigned integer representing CUDA architecture (SM)
  template<class BLAS, unsigned int Architecture>
  struct is_supported_rmem_restrict : std::bool_constant<...> { };

  // Helper variable template
  template<class BLAS, unsigned int Architecture>
  inline constexpr bool is_supported_rmem_restrict_v<BLAS, Architecture> = is_supported_rmem_restrict<BLAS, Architecture>::value;
}

// true if BLAS is supported on the provided CUDA architecture
cublasdx::is_supported_rmem_restrict<BLAS, Architecture>::value;

cublasdx::is_supported_rmem_restrict 检查 BLAS 操作是否在 Architecture CUDA 架构上受支持，假设

1. 使用寄存器 API（只有矩阵 A 和 B 必须适合共享内存）。

2. 输入精度与计算精度相同。

3. 矩阵不使用别名或相互重叠其范围，

   • 即，矩阵 B 的第一个元素在矩阵 A 的最后一个元素之后开始

// true if BLAS is supported on the provided CUDA architecture
cublasdx::is_supported_rmem_restrict<BLAS, Architecture>::value;

要求

• BLAS 必须定义尺寸、函数和转置模式。请参阅 描述算子 部分。

• BLAS 必须包含 Block 算子。

• BLAS 不能通过 SM 算子定义目标 CUDA 架构。

示例

using namespace cublasdx;

using BLAS = decltype(Size<128, 128, 128>() + Type<type::real>() + Block() + Precision<float>());
cublasdx::is_supported_smem_restrict<BLAS, 800>::value; // false
cublasdx::is_supported_rmem_restrict<BLAS, 800>::value; // true

suggested_leading_dimension_of

namespace cublasdx {
  // BLAS - BLAS description without CUDA architecture defined using SM operator
  // Architecture - unsigned integer representing CUDA architecture (SM)
  template<class BLAS, unsigned int Architecture>
  struct suggested_leading_dimension_of {
      static constexpr unsigned int lda;
      static constexpr unsigned int ldb;
      static constexpr unsigned int ldc;

      using type = LeadingDimension<lda, ldb, ldc>;

      using value_type                  = COMMONDX_STL_NAMESPACE::tuple<unsigned int, unsigned int, unsigned int>;
      static constexpr value_type value = value_type {lda, ldb, ldc};
      constexpr                   operator value_type() const noexcept { return value; }
  }

  // LeadingDimension operator with suggested leading dimensions
  template<class BLAS, unsigned int Architecture>
  using suggested_leading_dimension_of_t = typename suggested_leading_dimension_of<BLAS, Architecture>::type;

  template<class Description, unsigned int Architecture>
  inline constexpr COMMONDX_STL_NAMESPACE::tuple<unsigned int, unsigned int, unsigned int> suggested_leading_dimension_of_v = suggested_leading_dimension_of<Description, Architecture>::value;

  template<class Description, unsigned int Architecture>
  inline constexpr unsigned int suggested_leading_dimension_of_v_a = suggested_leading_dimension_of<Description, Architecture>::lda;
  template<class Description, unsigned int Architecture>
  inline constexpr unsigned int suggested_leading_dimension_of_v_b = suggested_leading_dimension_of<Description, Architecture>::ldb;
  template<class Description, unsigned int Architecture>
  inline constexpr unsigned int suggested_leading_dimension_of_v_c = suggested_leading_dimension_of<Description, Architecture>::ldc;
}

类型 cublasdx::suggested_leading_dimension_of 为矩阵 A、B 和 C 提供建议的引导维度。建议尝试使用它们，因为在许多情况下，这将提高 BLAS 操作的性能。当 M、N、K 是 2 的幂或 16 的倍数时，它可能特别有用。您可以查看 cuBLASDx 性能示例。

要求

• BLAS 必须定义尺寸、函数和转置模式。请参阅 描述算子 部分。

• BLAS 必须包含 Block 算子。

• 用于架构 A1 的建议引导维度不得与架构 A2 一起使用。

示例

using namespace cublasdx;

using BLAS1 = decltype(Size<128, 128, 128>() + Type<type::real>() + Block() + Precision<float>());
using SuggestedLD = cublasdx::suggested_leading_dimension_of_t<BLAS, 900>;
using BLAS2 = decltype(BLAS1() + SuggestedLD() + SM<900>());

建议对齐特性

// cublasdx::Alignment
cublasdx::suggested_alignment_of_t<BLAS>

// std::tuple<unsigned int, unsigned int, unsigned int>
cublasdx::suggested_alignment_of_v<BLAS>

// unsigned int
cublasdx::suggested_alignment_of<BLAS>::a
cublasdx::suggested_alignment_of_v_a<BLAS>

cublasdx::suggested_alignment_of<BLAS>::b
cublasdx::suggested_alignment_of_v_b<BLAS>

cublasdx::suggested_alignment_of<BLAS>::c
cublasdx::suggested_alignment_of_v_c<BLAS>

输入矩阵 A、B 和 C 的建议对齐（以字节为单位）。

它们设置为最大支持的对齐，即 cublasdx::MaxAlignment。请参阅 对齐算子。建议使用它们，因为这些对齐更可能导致向量化共享内存访问，从而加快执行速度。