| 8 | using namespace cuda; |
| 9 | |
| 10 | const void* MatrixMulForwardImpl::AlgoFloat32SIMT::get_available_op( |
| 11 | const SizeArgs& args) const { |
| 12 | using namespace cutlass::library; |
| 13 | auto&& param = args.opr->param(); |
| 14 | auto layoutA = |
| 15 | param.transposeA ? LayoutTypeID::kColumnMajor : LayoutTypeID::kRowMajor; |
| 16 | auto layoutB = |
| 17 | param.transposeB ? LayoutTypeID::kColumnMajor : LayoutTypeID::kRowMajor; |
| 18 | |
| 19 | int alignment = min_alignment_requirement(); |
| 20 | GemmKey key{ |
| 21 | NumericTypeID::kF32, |
| 22 | layoutA, |
| 23 | NumericTypeID::kF32, |
| 24 | layoutB, |
| 25 | NumericTypeID::kF32, |
| 26 | LayoutTypeID::kRowMajor, |
| 27 | NumericTypeID::kF32, |
| 28 | m_algo_param.threadblock_m, |
| 29 | m_algo_param.threadblock_n, |
| 30 | m_algo_param.threadblock_k, |
| 31 | m_algo_param.warp_m, |
| 32 | m_algo_param.warp_n, |
| 33 | m_algo_param.warp_k, |
| 34 | 1, |
| 35 | 1, |
| 36 | 1, |
| 37 | 2, |
| 38 | alignment, |
| 39 | alignment, |
| 40 | SplitKMode::kNone}; |
| 41 | return (void*)Singleton::get().operation_table.find_op(key); |
| 42 | } |
| 43 | |
| 44 | bool MatrixMulForwardImpl::AlgoFloat32SIMT::is_available(const SizeArgs& args) const { |
| 45 | bool available = args.opr->param().format == param::MatrixMul::Format::DEFAULT && |