| 192 | |
| 193 | template <typename Opr> |
| 194 | typename TimedProfiler<Opr>::TResult TimedProfiler<Opr>::prof_impl( |
| 195 | const TParam& raw_param) { |
| 196 | MIDOUT_B(Opr, midout_iv(MGB_HASH_STR("TimedProfiler::prof_impl"))) |
| 197 | #if MGB_ROCM |
| 198 | bool miopen_algo_search_enabled; |
| 199 | megcore::getMIOpenAlgoSearchStatus(&miopen_algo_search_enabled); |
| 200 | mgb_assert(miopen_algo_search_enabled, "MIOpen algo search not enabled"); |
| 201 | #endif |
| 202 | auto&& param = raw_param.as_single_pod<Param>(); |
| 203 | CompNode cn = CompNode::load(param.comp_node_physical, param.comp_node_logical); |
| 204 | auto megdnn_opr = opr::intl::create_megdnn_opr<Opr>(cn); |
| 205 | std::array<TensorLayout, arity> layouts; |
| 206 | |
| 207 | auto from_enum = [&](DTypeEnum enumv) -> DType { |
| 208 | switch (enumv) { |
| 209 | |
| 210 | #define cb(_dt) \ |
| 211 | case DTypeTrait<_dt>::enumv: \ |
| 212 | return _dt(1.0f, static_cast<uint8_t>(0)) |
| 213 | cb(dtype::Quantized8Asymm); |
| 214 | cb(dtype::Quantized4Asymm); |
| 215 | #undef cb |
| 216 | |
| 217 | #define cb(_dt) \ |
| 218 | case DTypeTrait<_dt>::enumv: \ |
| 219 | return _dt(1.0f) |
| 220 | |
| 221 | cb(dtype::QuantizedS8); |
| 222 | cb(dtype::QuantizedS16); |
| 223 | cb(dtype::QuantizedS32); |
| 224 | cb(dtype::QuantizedS4); |
| 225 | default: |
| 226 | return DType::from_enum(enumv); |
| 227 | #undef cb |
| 228 | } |
| 229 | }; |
| 230 | for (int i = 0; i < arity; ++i) { |
| 231 | layouts[i] = {param.shapes[i], from_enum(param.dtypes[i])}; |
| 232 | } |
| 233 | |
| 234 | megdnn_opr->param() = param.opr_param; |
| 235 | megdnn_opr->execution_policy() = param.execution_policy.deserialize(); |
| 236 | |
| 237 | // Allocate preprocessed weight buffers. |
| 238 | TensorLayoutArray preprocessed_layout; |
| 239 | if_constexpr<opr_supports_preprocess<Opr>()>([&](auto _) { |
| 240 | if (param.allow_weight_preprocess) { |
| 241 | preprocessed_layout = APPLY( |
| 242 | _(megdnn_opr)->deduce_preprocessed_filter_layout(args...), layouts); |
| 243 | } |
| 244 | }); |
| 245 | |
| 246 | megdnn::Algorithm* algo = |
| 247 | megdnn_opr->get_algorithm_from_desc(megdnn_opr->execution_policy().algo); |
| 248 | mgb_assert(algo); |
| 249 | |
| 250 | #if !MGB_BUILD_SLIM_SERVING |
| 251 | #if MGB_CUDA || MGB_ROCM |
nothing calls this directly
no test coverage detected