执行方法#
执行方法用于运行用户使用 cuBLASDx 运算符定义的 BLAS 函数。
注意
目前，cuBLASDx 仅支持在 CUDA 线程块级别（块执行）上执行。
块执行方法#
如果描述符是使用 块运算符 构建的，并且 is_complete_blas_execution 特性 为 true，则可以使用块执行方法。
寄存器 API#
方法 execute(...) 运行由 BLAS 描述符定义的计算，接受两种类型的参数。
using BLAS = decltype(cublasdx::Size<M, N, K>() + ...);
// #1 - Registers with accumulator API
template<class AEngine, class ALayout, // Types derived from pointer and layout used to create tensor_a
class BEngine, class BLayout, // Types derived from pointer and layout used to create tensor_b
class CEngine, class CLayout, // Types derived from pointer and layout used to create tensor_c
class ALoadOp = cublasdx::identity, // Transform operation applied when data is loaded from matrix A
class BLoadOp = cublasdx::identity> // Transform operation applied when data is loaded from matrix B
inline __device__ void execute(const cublasdx::tensor<AEngine, ALayout>& tensor_a,
const cublasdx::tensor<BEngine, BLayout>& tensor_b,
cublasdx::tensor<CEngine, CLayout> & tensor_c,
const ALoadOp& a_load_op = {},
const BLoadOp& b_load_op = {})
// #2 - Registers without accumulator API
template<class AEngine, class ALayout,
class BEngine, class BLayout,
class ALoadOp = cublasdx::identity,
class BLoadOp = cublasdx::identity>
inline __device__ auto execute(const cublasdx::tensor<AEngine, ALayout>& tensor_a,
const cublasdx::tensor<BEngine, BLayout>& tensor_b,
const ALoadOp& a_load_op = {},
const BLoadOp& b_load_op = {})
方法 #1 接受用于 A 和 B 矩阵的共享内存张量，但接受用于 C 矩阵的寄存器片段。它不返回任何内容，因为它会将 A 与 B 相乘的结果添加到 C，从而得到
\(\mathbf{C}_{m\times n} = \mathbf{A}_{m\times k} \times \mathbf{B}_{k\times n} + \mathbf{C}_{m\times n}\)
寄存器片段必须预先存在，可以来自方法 #2 的先前执行，或者从分区器对象创建（请参阅 数据分区器）。它必须与它所用于的 GEMM 的精度和分区完全匹配。
方法 #2 仅接受用于 A 和 B 矩阵的共享内存张量。它返回一个不透明的元组，其中包含结果寄存器片段及其对应的分区器。结果对应于
\(\mathbf{C}_{m\times n} = \mathbf{A}_{m\times k} \times \mathbf{B}_{k\times n}\)
C++17 结构化绑定可用于清晰地检索这两个值
auto [c_register_fragment, partitioner] = BLAS().execute(a_shared_tensor, b_shared_tensor, a_load_op, b_load_op);
下面的代码示例演示了如何使用两种 execute(...) 方法。
#include <cublasdx.hpp>
using GEMM = decltype(cublasdx::Size<32, 32, 32>()
+ cublasdx::Precision<cublasdx::tfloat32_t, cublasdx::tfloat32_t, float>()
+ cublasdx::Type<cublasdx::type::real>()
+ cublasdx::Arrangement<cublasdx::row_major, cublasdx::col_major>()
+ cublasdx::Function<cublasdx::function::MM>()
+ cublasdx::MaxAlignment() // max alignment (16, 16, 16) is the default
+ cublasdx::SM<800>()
+ cublasdx::Block());
using a_data_type = typename GEMM::a_value_type;
using b_data_type = typename GEMM::b_value_type;
using c_data_type = typename GEMM::c_value_type;
extern __shared__ __align__(16) char smem[];
// smem_<a/b> are aligned to cublasdx::alignment_of<GEMM>::<a/b>
auto [smem_a, smem_b] = cublasdx::slice_shared_memory_ab<GEMM>(smem);
//*********** Method #1, register API with accumulator
{
// Make global memory tensor
auto a_global_tensor = cublasdx::make_tensor(a, GEMM::get_layout_gmem_a());
auto b_global_tensor = cublasdx::make_tensor(b, GEMM::get_layout_gmem_b());
auto c_global_tensor = cublasdx::make_tensor(c, GEMM::get_layout_gmem_c());
// Make shared memory tensor
auto a_shared_tensor = cublasdx::make_tensor(smem_a, GEMM::get_layout_smem_a());
auto b_shared_tensor = cublasdx::make_tensor(smem_b, GEMM::get_layout_smem_b());
// Load data from global to shared memory using cublasdx::copy API
using alignment = cublasdx::alignment_of<GEMM>;
cublasdx::copy<GEMM, alignment::a>(a_global_tensor, a_shared_tensor);
cublasdx::copy<GEMM, alignment::b>(b_global_tensor, b_shared_tensor);
cublasdx::copy_wait();
// Execute
auto partitioner = GEMM::get_partitioner();
auto c_register_fragment = partitioner.make_accumulator_fragment();
cublasdx::copy_fragment<alignment::c>(c_global_tensor, c_register_fragment, partitioner);
GEMM().execute(a_shared_tensor, b_shared_tensor, c_register_fragment);
// Store back to global memory using cublasdx::copy_fragment API
cublasdx::copy_fragment<alignment::c>(c_register_fragment, c_global_tensor, partitioner);
}
//*********** Method #2, cublasdx tensor APIs, without accumulator
{
// Make global memory tensor
auto a_global_tensor = cublasdx::make_tensor(a, GEMM::get_layout_gmem_a());
auto b_global_tensor = cublasdx::make_tensor(b, GEMM::get_layout_gmem_b());
auto c_global_tensor = cublasdx::make_tensor(c, GEMM::get_layout_gmem_c());
// Make shared memory tensor
auto a_shared_tensor = cublasdx::make_tensor(smem_a, GEMM::get_layout_smem_a());
auto b_shared_tensor = cublasdx::make_tensor(smem_b, GEMM::get_layout_smem_b());
// Load data from global to shared memory using cublasdx::copy API
using alignment = cublasdx::alignment_of<GEMM>;
cublasdx::copy<GEMM, alignment::a>(a_global_tensor, a_shared_tensor);
cublasdx::copy<GEMM, alignment::b>(b_global_tensor, b_shared_tensor);
cublasdx::copy_wait();
// Execute
auto [c_register_fragment, partitioner] = GEMM().execute(a_shared_tensor, b_shared_tensor);
// Store back to global memory using cublasdx::copy_fragment API
cublasdx::copy_fragment<alignment::c>(c_register_fragment, c_global_tensor, partitioner);
}
输入数据属性#
注意
从 cuBLASDx 0.3.0 开始，计算精度已与数据精度解耦，即，每个矩阵的输入/输出数据可以是任意类型（即使是浮点 GEMM 的整数输入），前提是设置了 对齐运算符 并且满足以下至少一个条件
警告
如果使用与输入类型解耦的计算精度，则必须显式设置 对齐运算符。
标量（alpha 和 beta）的底层元素类型默认假定为 BLAS::c_value_type，但它们可以是任何类型，只要
它们的对齐方式和大小与
BLAS::c_value_type相同，并且它们可以转换为
BLAS::c_value_type
转换操作输入#
所有方法都接受转换仿函数。a_load_op、b_load_op、c_load_op 在从每个矩阵读取元素时应用，c_store_op 在矩阵乘法的结果存储在 C 矩阵之前应用。每个仿函数都必须表示一个元素级转换，它
对于加载转换：接受
execute(...)方法的相应输入类型，并返回可隐式转换为BLAS::<a/b/c>_value_type类型的值。对于存储转换：接受
BLAS::<a/b/c>_value_type并返回execute(...)方法的相应输入类型。
示例
using GEMM = decltype(Size<128, 128, 128>() + Type<type::real>() + Precision<float, float, double>() + Block() + ...);
struct multiple_by_2 {
template<class T>
__device__ constexpr T operator()(const T arg) const {
return arg * static_cast<T>(2.0f);
}
};
struct negate {
template <class T>
__device__ constexpr T operator()(const T arg) const {
return -arg;
}
};
GEMM().execute(..., multiple_by_2{}, cublasdx::conjugate{}, cublasdx::identity{}, negate{});
警告
不保证在不同 CUDA 架构的 GPU 上使用完全相同的输入执行完全相同的 BLAS 函数将产生位相同的結果。
值格式#
BLAS::a_value_type
BLAS::b_value_type
BLAS::c_value_type
对于每种精度的复数，复数中的第一个值是实部，第二个值是虚部。对于实数，BLAS::<a/b/c>_value_type 与用于描述 BLAS 的 Precision<PA, PB, PC> 中的 P 相同（或默认精度）。
输入/输出数据格式#
本节介绍正确计算所需的输入和输出数据格式（布局）。
GEMM (function::MM)#
用于通用矩阵乘法的张量 API（execute() 方法，它期望使用 cublasdx::tensor 表示的矩阵）接受由具有任意布局的张量表示的矩阵。由于张量对象携带有关矩阵的维度、内存位置和布局的所有信息，因此不需要其他隐式假设。矩阵的维度必须与 大小 运算符定义的维度匹配。另请参阅 获取内存布局 和 建议的共享内存布局 部分。
用于通用矩阵乘法的指针 API（execute() 的 #2 和 #3 重载）假定输入矩阵 matrix_a、matrix_b、matrix_c 中的值按照添加到描述中的 排列 运算符定义的方式存储（默认情况下，matrix_a 为行优先格式，matrix_b 为列优先格式，matrix_c 为列优先格式）。