运算符#
运算符用于描述属性并配置我们要解决问题的执行。它们分为描述运算符和执行运算符。
描述运算符#
运算符 |
默认值 |
描述 |
|---|---|---|
未设置 |
定义问题大小。 |
|
未设置 |
BLAS 函数。对于 |
|
|
全局内存中每个矩阵的排列以及共享内存中的默认排列。 |
|
|
每个矩阵的转置模式（自 0.2.0 版本起已弃用）。 |
|
|
GEMM |
|
|
输入和输出数据类型（ |
|
由 |
矩阵 |
|
|
矩阵 |
|
未设置 |
应为其生成 BLAS 函数的目标 CUDA 架构。 |
描述运算符定义了我们要解决的问题。与执行运算符结合使用，它们形成了一个完整的函数描述符，可以在 GPU 上执行。
添加运算符（以任意顺序）以构造操作描述符类型。例如，要描述用于非转置矩阵 A (m x k)、B (k x n)、C (m x n) 的矩阵乘法，其中复双精度值，其中 m = 8、n = 16、k = 32 用于在 Volta 架构上执行，可以编写
#include <cublasdx.hpp>
using GEMM = decltype(cublasdx::Size<8, 16, 32>()
+ cublasdx::Precision<double>()
+ cublasdx::Type<cublasdx::type::complex>()
+ cublasdx::Arrangement<cublasdx::col_major, cublasdx::col_major>()
+ cublasdx::Function<cublasdx::function::MM>()
+ cublasdx::SM<700>());
为了使函数描述符完整，需要以下内容
大小运算符#
cublasdx::Size<unsigned int M, unsigned int N, unsigned int K>()
设置要执行函数的的问题大小。
对于 GEMM
M- 矩阵op(A)和C中的逻辑行数。N- 矩阵op(B)和C中的逻辑列数。K- 矩阵op(A)中的逻辑列数和C中的行数。
例如，对于 GEMM，M、N 和 K 指定 A (M x K) 矩阵乘以 B (K x N) 矩阵，结果为 C (M x N) 矩阵（假设 A 和 B 未转置）。请参阅排列、转置模式和GEMM。
类型运算符#
cubladx::Type<cublasdx::type T>;
namespace cublasdx {
enum class type
{
real,
complex
};
}
设置计算中使用的输入和输出数据类型。对于实数据类型，请使用 type::real，对于复数据类型，请使用 type::complex。
精度运算符#
cublasdx::Precision<PA, PB=PA, PC=PA>;
设置 A、B 和 C 的计算精度。
精度类型可以是浮点型
__halffloatdouble__nv_fp8_e5m2__nv_fp8_e4m3__nv_bfloat16cublasdx::tfloat32_t
或整型
int8_tuint8_tint16_tuint16_tint32_tuint32_tint64_tuint64_t
它描述了用于输入和输出值的计算精度。这意味着这是输入值将在执行乘法指令之前转换成的类型。
注意
从 cuBLASDx 0.3.0 开始，计算精度已与数据精度分离，即每个矩阵的输入/输出数据可以是任意类型（甚至浮点 GEMM 的整型输入），前提是设置了对齐运算符并且满足以下条件之一
警告
如果使用与输入类型分离的计算精度，则必须显式设置对齐运算符。
排列运算符#
cublasdx::Arrangement<cublasdx::arrangement A, cublasdx::arrangement B = cublasdx::col_major, cublasdx::arrangement C = cublasdx::col_major>;
namespace cublasdx {
enum class arrangement
{
col_major,
row_major
};
inline constexpr auto col_major = arrangement::col_major;
inline constexpr auto left_layout = arrangement::col_major;
inline constexpr auto row_major = arrangement::row_major;
inline constexpr auto right_major = arrangement::row_major;
}
设置函数中使用的全局 A、B 和 C 矩阵的顺序。顺序可以是列优先或行优先。排列运算符直接影响获取内存布局和建议的共享内存布局方法，通过显式设置默认的获取内存布局值并选择建议的共享内存布局以优化全局共享传输。
为矩阵选择特定顺序并不意味着该函数不接受具有不同布局的矩阵，但它可能会影响整体性能。
警告
Arrangement 和 TransposeMode 运算符不能同时定义。
转置模式运算符#
警告
TransposeMode 运算符自 0.2.0 版本起已弃用，并且可能会在未来版本中删除。
cublasdx::TransposeMode<cublasdx::transpose_mode ATransposeMode, cublasdx::transpose_mode BTransposeMode>;
namespace cublasdx {
enum class transpose_mode
{
non_transposed,
transposed,
conj_transposed,
};
inline constexpr auto N = transpose_mode::non_transposed;
inline constexpr auto T = transpose_mode::transposed;
inline constexpr auto C = transpose_mode::conj_transposed;
}
设置函数中使用的 A 和 B 矩阵的转置模式。例如，TransposeMode<N, N>() 将 A 和 B 矩阵的转置模式设置为 GEMM 的 non-transposed。转置模式的可能值是
transpose_mode::non_transposed,transpose_mode::transposed, 和transpose_mode::conj_transposed(共轭转置)。
警告
排列和
TransposeMode运算符不能同时定义。对于
TransposeMode运算符中的矩阵使用transpose_mode::non_transposed对应于 排列运算符中的arrangement::col_major。对于
TransposeMode运算符中的矩阵使用transpose_mode::transposed对应于 排列运算符中的arrangement::row_major。对于
TransposeMode运算符中的矩阵使用transpose_mode::conj_transposed对应于 排列运算符中的arrangement::row_major和在 execute() 方法中作为该矩阵的变换运算符传递的cublasdx::conjugate。
主导维度运算符#
cublasdx::LeadingDimension<unsigned int LDA, unsigned int LDB, unsigned int LDC>()
定义矩阵 A、B 和 C 的主导维度。矩阵的主导维度是列优先矩阵到下一列开头或行优先矩阵到下一行开头的步幅（以元素为单位）。
如果在通用矩阵乘法运算的描述中使用了排列，则可以用以下方式描述 A、B、C 矩阵
矩阵
A的实际维度是 \(LDA\times K\)，其中如果A是列优先矩阵，则LDA >= M，否则为 \(LDA\times M\)，其中LDA >= K。矩阵
B的实际维度是 \(LDB\times N\)，其中如果B是列优先矩阵，则LDA >= K，否则为 \(LDB\times K\)，其中LDB >= N。矩阵
C的实际维度是 \(LDC\times N\)，其中如果C是列优先矩阵，则LDC >= M，否则为 \(LDC\times M\)，其中LDB >= N。
矩阵也可以使用布局来描述（请参阅 CuTe：布局），即整数元组对：形状和步幅（元素之间的距离）。简而言之，形状表示每个维度中元素的数量，而步幅表示每个维度中元素之间的距离。
A- 形状为 \(M\times K\) 的矩阵，如果A是列优先矩阵，则在第 1 维度中步幅为1，在第 2 维度中步幅为LDA，否则在第 1 维度中为LDA，在第 2 维度中为1。B- 形状为 \(K\times N\) 的矩阵，如果B是列优先矩阵，则在第 1 维度中步幅为1，在第 2 维度中步幅为LDB，否则在第 1 维度中为LDB，在第 2 维度中为1。C- 形状为 \(M\times N\) 的矩阵，如果C是列优先矩阵，则在第 1 维度中步幅为1，在第 2 维度中步幅为LDC，否则在第 1 维度中为LDC，在第 2 维度中为1。
警告
TransposeMode 运算符自 0.2.0 版本起已弃用，并且可能会在未来版本中删除。
如果在通用矩阵乘法运算的描述中使用了转置模式（已弃用），则在 BLAS 命名法中，可以用以下方式描述 A、B、C 矩阵的维度
矩阵
A的实际维度是 \(LDA\times K\)，其中如果A未转置，则LDA >= M，否则为 \(LDA\times M\)，其中LDA >= K。矩阵
B的实际维度是 \(LDB\times N\)，其中如果B未转置，则LDA >= K，否则为 \(LDB\times K\)，其中LDB >= N。矩阵
C的实际维度是 \(LDC\times N\)，其中LDC >= M。
对齐运算符#
cublasdx::Alignment<unsigned int AAlignment, unsigned int BAlignment, unsigned int CAlignment>()
cublasdx::MaxAlignment = Alignment<16, 16, 16> // alias of maximum supported alignements
定义传递给 execute(...) 方法的矩阵 A、B 和 C 的指针（原始指针或封装在 CuTe 张量中）的对齐方式（以字节为单位）。
请注意，对于给定的描述运算符，对齐方式直接影响所需的共享内存量。此外，它还可能影响性能。
要求
AAlignments、BAlignments、CAlignments应该是 2 的幂，并且小于或等于最大支持的对齐方式，即 16 字节。AAlignments、BAlignments、CAlignments是所选计算值类型或输入值类型（如果使用分离精度）的对齐方式的倍数。
警告
如果使用与输入类型分离的计算精度，则必须显式设置对齐运算符。
函数运算符#
cublasdx::Function<cublasdx::function F>()
namespace cublasdx {
enum class function
{
MM
};
}
设置要执行的 BLAS 函数。
通用矩阵乘法#
Function<function::MM> 将运算设置为通用矩阵乘法，定义为以下之一
\(\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}\)
其中 \({\alpha}\) 和 \({\beta}\) 是标量（实数或复数），A、B 和 C 是维度分别为 \(A: m\times k\)、\(B: k\times n\) 和 \(C: m\times n\) 的矩阵。
SM 运算符#
cublasdx::SM<unsigned int CC>()
设置底层 BLAS 函数要使用的目标架构 CC。支持的架构有
Volta：
700和720(sm_70, sm_72)。Turing：
750(sm_75)。Ampere:
800,860和870(sm_80, sm_86, sm_87)。Ada:
890(sm_89)。Hopper:
900(sm_90, sm_90a)。
注意
当为 9.0a 计算能力编译 cuBLASDx 时，在 SM 运算符中使用 900（另请参阅 CUDA C++ 编程指南：功能可用性）。
警告
不保证在不同 CUDA 架构的 GPU 上使用完全相同的输入执行完全相同的 BLAS 函数会产生位相同的结果。
执行运算符#
执行运算符配置函数在 GPU 上的运行方式。与 描述运算符 结合使用，它们构成了一个完整的函数描述符，可以在 GPU 上执行。
Block 运算符#
cublasdx::Block()
生成在单个 CUDA 块中运行的集体操作。线程将协同工作以计算集体操作。可以使用 块配置运算符 配置布局和参与执行的线程数。
例如，以下代码示例为 GEMM 函数创建一个函数描述符，该函数将在单个 CUDA 块中运行
#include <cublasdx.hpp>
using GEMM = decltype(cublasdx::Size<32, 32, 64>()
+ cublasdx::Precision<double, __half, double>()
+ cublasdx::Type<cublasdx::type::real>()
+ cublasdx::TransposeMode<cublasdx::T, cublasdx::N>()
+ cublasdx::Function<cublasdx::function::MM>()
+ cublasdx::SM<700>()
+ cublasdx::Block());
块配置运算符#
块配置运算符允许用户配置单个 CUDA 块的块大小。
运算符 |
默认值 |
描述 |
|---|---|---|
基于启发式方法 |
用于执行 BLAS 函数的线程数。 |
注意
块配置运算符只能与 Block 运算符 一起使用。
警告
不保证使用完全相同的输入，但在不同
前导维度 (LeadingDimension),
CUDA 架构 (SM), 或
线程数 (BlockDim)
的情况下执行完全相同的 BLAS 函数会产生位相同的结果。
BlockDim 运算符#
struct cublasdx::BlockDim<unsigned int X, unsigned int Y, unsigned int Z>()
将 CUDA 块大小设置为 (X, Y, Z) 以配置执行，这意味着它设置参与执行的线程数及其布局。使用此运算符，用户可以在具有不同线程数的 1D、2D 或 3D 块中运行 BLAS 函数。可以通过 BLAS::block_dim 特征访问设置的块维度。
将 BlockDim<X, Y, Z> 添加到描述中对 BLAS 函数的执行提出了以下要求
内核必须使用 3D 块维度
dim3(X1, Y1, Z1)启动，其中X1 >= X，Y1 >= Y，并且Z1 >= Z，并且对于 1D
BlockDim<X>，内核必须使用dim3(X1, Y1, Z1)启动，其中X1 >= X。对于 2D
BlockDim<X, Y>，内核必须使用dim3(X, Y1, Z1)启动，其中Y1 >= Y。对于 3D
BlockDim<X, Y, Z>，内核必须使用dim3(X, Y, Z1)启动，其中Z1 >= Z。
X * Y * Z个线程必须参与执行。参与线程必须是连续（相邻）的线程。
cuBLASDx 的未来版本可能会解除或放宽列出的要求。
注意
cuBLASDx 无法在运行时验证所有内核启动配置并检查是否满足所有要求，因此用户有责任遵守上面列出的规则。违反这些规则被认为是未定义的行为，并可能导致不正确的结果和/或失败。
示例
BlockDim<64>，内核使用块维度 dim3(128, 1, 1) 启动 - OK（正确）BlockDim<64>，内核使用块维度 dim3(64, 4, 1) 启动 - OK（正确）BlockDim<64>，内核使用块维度 dim3(64, 2, 2) 启动 - OK（正确）BlockDim<16, 16>，内核使用块维度 dim3(16, 32, 1) 启动 - OK（正确）BlockDim<16, 16>，内核使用块维度 dim3(16, 16, 2) 启动 - OK（正确）BlockDim<8, 8, 8>，内核使用块维度 dim3(8, 8, 16) 启动 - OK（正确）BlockDim<64>，内核使用块维度 dim3(32, 1, 1) 启动 - INCORRECT（不正确）BlockDim<64>，内核使用块维度 dim3(32, 2, 1) 启动 - INCORRECT（不正确）BlockDim<16, 16>，内核使用块维度 dim3(256, 1, 1) 启动 - INCORRECT（不正确）BlockDim<8, 8, 8>，内核使用块维度 dim3(512, 2, 1) 启动 - INCORRECT（不正确）BlockDim 的值可以从 BLAS 描述中通过 BLAS::block_dim 特征访问。当未设置 BlockDim 时，将使用默认块维度（默认值为 BLAS::suggested_block_dim）。
如果 cuBLASDx 提供的默认块维度小于内核的最佳维度，那么在增加参与计算的线程数之前，尝试默认值仍然可能是最佳选择。
限制
X * Y * Z必须大于或等于 32。
注意
建议
X * Y * Z为 32、64、128、256、512 或 1024。建议
X * Y * Z是 32 的倍数。