diff --git a/fin b/fin
index 74382a6746..ebf9b355da 160000
--- a/fin
+++ b/fin
@@ -1 +1 @@
-Subproject commit 74382a6746d5aa0117f15d0a7cb779607caaa78d
+Subproject commit ebf9b355da09e10cb77e4b163a02e57be954095e
diff --git a/src/convolution_api.cpp b/src/convolution_api.cpp
index 02ae00d0aa..b428362fd0 100644
--- a/src/convolution_api.cpp
+++ b/src/convolution_api.cpp
@@ -66,13 +66,8 @@ static inline auto MakeFwdCtxAndProblem(miopenHandle_t handle,
                                                                        conv,
                                                                        direction};
 
-    auto ctx = [&] {
-        auto tmp = ExecutionContext{&miopen::deref(handle)};
-        tmp.DetectRocm();
-        problem.SetupFloats(tmp);
-        return tmp;
-    }();
-
+    auto ctx = ExecutionContext{&miopen::deref(handle)};
+    problem.SetupFloats(ctx);
     return std::make_tuple(std::move(ctx), std::move(problem));
 }
 
@@ -97,13 +92,8 @@ static inline auto MakeBwdCtxAndProblem(miopenHandle_t handle,
                                                                        conv,
                                                                        direction};
 
-    auto ctx = [&] {
-        auto tmp = ExecutionContext{&miopen::deref(handle)};
-        tmp.DetectRocm();
-        problem.SetupFloats(tmp);
-        return tmp;
-    }();
-
+    auto ctx = ExecutionContext{&miopen::deref(handle)};
+    problem.SetupFloats(ctx);
     return std::make_tuple(std::move(ctx), std::move(problem));
 }
 
@@ -127,13 +117,8 @@ static inline auto MakeWrWCtxAndProblem(miopenHandle_t handle,
                                                                        conv,
                                                                        direction};
 
-    auto ctx = [&] {
-        auto tmp = ExecutionContext{&miopen::deref(handle)};
-        tmp.DetectRocm();
-        problem.SetupFloats(tmp);
-        return tmp;
-    }();
-
+    auto ctx = ExecutionContext{&miopen::deref(handle)};
+    problem.SetupFloats(ctx);
     return std::make_tuple(std::move(ctx), std::move(problem));
 }
 
diff --git a/src/execution_context.cpp b/src/execution_context.cpp
index d6f8d1c74c..a9cd8806b2 100644
--- a/src/execution_context.cpp
+++ b/src/execution_context.cpp
@@ -206,7 +206,7 @@ bool IsHipKernelsEnabled()
 #endif
 }
 
-ExecutionContext& ExecutionContext::DetectRocm()
+void ExecutionContext::DetectRocm()
 {
     use_binaries            = false;
     use_asm_kernels         = false;
@@ -220,7 +220,6 @@ ExecutionContext& ExecutionContext::DetectRocm()
         use_binaries = !IsDisabled(MIOPEN_DEBUG_AMD_ROCM_PRECOMPILED_BINARIES{});
 #endif
     }
-    return *this;
 }
 
 } // namespace miopen
diff --git a/src/fusion.cpp b/src/fusion.cpp
index 1c1ab2e0ec..170e1b2d50 100644
--- a/src/fusion.cpp
+++ b/src/fusion.cpp
@@ -517,7 +517,6 @@ miopenStatus_t FusionPlanDescriptor::Compile(Handle& handle)
     const auto solvers    = GetFusedSolvers();
     auto fusion_ctx       = FusionContext{handle};
     auto fusion_problem   = FusionDescription{this};
-    fusion_ctx.DetectRocm();
     AnyInvokeParams invoke_params;
     miopen::OperatorArgs params;
     std::vector<Allocator::ManageDataPtr> invoke_bufs;
diff --git a/src/include/miopen/execution_context.hpp b/src/include/miopen/execution_context.hpp
index 4eba4b571d..83b9abdc1a 100644
--- a/src/include/miopen/execution_context.hpp
+++ b/src/include/miopen/execution_context.hpp
@@ -94,17 +94,15 @@ struct ExecutionContext
     inline Handle& GetStream() const { return *stream; }
     inline void SetStream(Handle* stream_) { stream = stream_; }
 
-    ExecutionContext(Handle* stream_) : stream(stream_) {}
+    ExecutionContext() { DetectRocm(); }
+    ExecutionContext(Handle* stream_) : stream(stream_) { DetectRocm(); }
 
-    ExecutionContext()                        = default;
     virtual ~ExecutionContext()               = default;
     ExecutionContext(const ExecutionContext&) = default;
     ExecutionContext(ExecutionContext&&)      = default;
     ExecutionContext& operator=(const ExecutionContext&) = default;
     ExecutionContext& operator=(ExecutionContext&&) = default;
 
-    ExecutionContext& DetectRocm();
-
 #if MIOPEN_EMBED_DB
     std::string GetPerfDbPathEmbed() const
     {
@@ -281,6 +279,8 @@ struct ExecutionContext
 
 private:
     Handle* stream = nullptr;
+
+    void DetectRocm();
 };
 
 bool IsHipKernelsEnabled();
diff --git a/src/include/miopen/find_solution.hpp b/src/include/miopen/find_solution.hpp
index 2cc119a7c8..fb55298cdb 100644
--- a/src/include/miopen/find_solution.hpp
+++ b/src/include/miopen/find_solution.hpp
@@ -361,8 +361,7 @@ struct SolverContainer
             return;
         }
 
-        auto ctx = ExecutionContext{&handle};
-        ctx.DetectRocm();
+        auto ctx        = ExecutionContext{&handle};
         const auto slns = SearchForSolutions(ctx, problem, 1);
 
         if(slns.empty())
diff --git a/src/include/miopen/fusion/context.hpp b/src/include/miopen/fusion/context.hpp
index 5076fc4d73..4cc948c091 100644
--- a/src/include/miopen/fusion/context.hpp
+++ b/src/include/miopen/fusion/context.hpp
@@ -35,7 +35,6 @@ struct FusionContext : miopen::ExecutionContext
     ConvolutionContext GetConvContext(const miopen::ProblemDescription& conv_problem) const
     {
         auto ctx = ConvolutionContext{*this};
-        ctx.DetectRocm();
         conv_problem.conv_problem.SetupFloats(ctx);
         return ctx;
     }
diff --git a/src/ocl/convolutionocl.cpp b/src/ocl/convolutionocl.cpp
index fa50bca9d4..182d33b4e1 100644
--- a/src/ocl/convolutionocl.cpp
+++ b/src/ocl/convolutionocl.cpp
@@ -110,7 +110,6 @@ static Invoker PrepareInvoker(ExecutionContext ctx,
                               const NetworkConfig& config,
                               solver::Id solver_id)
 {
-    ctx.DetectRocm();
     problem.SetupFloats(ctx);
     ctx.do_search = false;
 
@@ -255,7 +254,6 @@ void ConvolutionDescriptor::FindConvFwdAlgorithm(Handle& handle,
         conv::ProblemDescription(xDesc, wDesc, yDesc, *this, conv::Direction::Forward);
     const auto ctx = [&] {
         auto tmp = ExecutionContext{&handle};
-        tmp.DetectRocm();
         problem.SetupFloats(tmp);
         tmp.do_search = exhaustiveSearch;
         return tmp;
@@ -649,7 +647,6 @@ std::vector<miopenConvSolution_t> GetSolutions(const ExecutionContext& exec_ctx,
     // All the above can be found by calling IsApplicable().
     // We need fully initialized context for this, see below.
     auto ctx = ConvolutionContext{exec_ctx};
-    ctx.DetectRocm();
 
     for(const auto& pair : fdb_record)
     {
@@ -724,7 +721,6 @@ std::size_t ConvolutionDescriptor::GetForwardSolutionWorkspaceSize(Handle& handl
         conv::ProblemDescription{xDesc, wDesc, yDesc, *this, conv::Direction::Forward};
     auto ctx = ConvolutionContext{};
     ctx.SetStream(&handle);
-    ctx.DetectRocm();
     if(sol.IsApplicable(ctx, problem))
         return sol.GetWorkspaceSize(ctx, problem);
     MIOPEN_THROW(miopenStatusBadParm,
@@ -804,7 +800,6 @@ void ConvolutionDescriptor::FindConvBwdDataAlgorithm(Handle& handle,
 
     const auto ctx = [&] {
         auto tmp = ExecutionContext{&handle};
-        tmp.DetectRocm();
         problem.SetupFloats(tmp);
         tmp.do_search = exhaustiveSearch;
         return tmp;
@@ -935,7 +930,6 @@ std::size_t ConvolutionDescriptor::GetBackwardSolutionWorkspaceSize(Handle& hand
         conv::ProblemDescription{dyDesc, wDesc, dxDesc, *this, conv::Direction::BackwardData};
     auto ctx = ConvolutionContext{};
     ctx.SetStream(&handle);
-    ctx.DetectRocm();
     if(sol.IsApplicable(ctx, problem))
         return sol.GetWorkspaceSize(ctx, problem);
     else
@@ -1013,7 +1007,6 @@ void ConvolutionDescriptor::FindConvBwdWeightsAlgorithm(Handle& handle,
         conv::ProblemDescription{dyDesc, dwDesc, xDesc, *this, conv::Direction::BackwardWeights};
     const auto ctx = [&] {
         auto tmp = ExecutionContext{&handle};
-        tmp.DetectRocm();
         problem.SetupFloats(tmp);
         tmp.do_search = exhaustiveSearch;
         return tmp;
@@ -1135,7 +1128,6 @@ std::size_t ConvolutionDescriptor::GetWrwSolutionWorkspaceSize(Handle& handle,
         conv::ProblemDescription{dyDesc, dwDesc, xDesc, *this, conv::Direction::BackwardWeights};
     auto ctx = ConvolutionContext{};
     ctx.SetStream(&handle);
-    ctx.DetectRocm();
     if(sol.IsApplicable(ctx, problem))
         return sol.GetWorkspaceSize(ctx, problem);
     else
diff --git a/src/problem.cpp b/src/problem.cpp
index 3cd811e70a..35a85074fe 100644
--- a/src/problem.cpp
+++ b/src/problem.cpp
@@ -259,8 +259,7 @@ std::vector<Solution> Problem::FindSolutionsImpl(Handle& handle,
     }
     else
     {
-        auto tmp_ctx = ExecutionContext{&handle};
-        tmp_ctx.DetectRocm();
+        auto tmp_ctx             = ExecutionContext{&handle};
         const auto workspace_max = conv_desc.GetWorkSpaceSize(tmp_ctx, conv_problem);
         workspace_size           = std::min(options.workspace_limit, workspace_max);
         owned_workspace          = workspace_size != 0 ? handle.Create(workspace_size) : nullptr;
@@ -349,7 +348,6 @@ std::vector<Solution> Problem::FindSolutionsImpl(Handle& handle,
     const auto legacy_problem = ProblemDescription{conv_problem};
     const auto netcfg         = conv_problem.BuildConfKey();
     auto conv_ctx             = ConvolutionContext{{&handle}};
-    conv_ctx.DetectRocm();
     conv_problem.SetupFloats(conv_ctx);
 
     decltype(auto) db = GetDb(conv_ctx);
diff --git a/src/solution.cpp b/src/solution.cpp
index 22d38546bf..5f5fa18512 100644
--- a/src/solution.cpp
+++ b/src/solution.cpp
@@ -178,7 +178,6 @@ void Solution::RunImpl(Handle& handle,
 
     const auto legacy_problem = ProblemDescription{conv_problem};
     auto conv_ctx             = ConvolutionContext{{&handle}};
-    conv_ctx.DetectRocm();
     conv_problem.SetupFloats(conv_ctx);
 
     decltype(auto) db        = GetDb(conv_ctx);
diff --git a/test/conv_common.hpp b/test/conv_common.hpp
index 4193b5c75d..f79964d5da 100644
--- a/test/conv_common.hpp
+++ b/test/conv_common.hpp
@@ -93,7 +93,6 @@ static inline bool is_direct_fwd_bwd_data_supported(miopen::Handle& handle,
         ctx.general_compile_options = "";
         ctx.SetStream(&handle);
         problem.SetupFloats(ctx);
-        ctx.DetectRocm();
         if(FindAllDirectSolutions(ctx, problem, {}).empty())
             return false;
     }
@@ -119,7 +118,6 @@ static inline bool is_direct_bwd_wrw_supported(miopen::Handle& handle,
     ctx.disable_perfdb_access   = true;
     ctx.SetStream(&handle);
     problem.SetupFloats(ctx);
-    ctx.DetectRocm();
 
     return !FindAllBwdWrW2DSolutions(ctx, problem, {}).empty();
 }
@@ -146,7 +144,6 @@ static inline bool skip_config(miopen::Handle& handle,
     ctx.disable_perfdb_access   = true;
     ctx.SetStream(&handle);
     problem.conv_problem.SetupFloats(ctx);
-    ctx.DetectRocm();
 
     return ctx.GetStream().GetDeviceName() == "gfx908" && problem.Is2d() && problem.IsFp16() &&
            problem.IsLayoutDefault() && ctx.use_hip_kernels && problem.GetGroupCount() == 1 &&
@@ -547,7 +544,7 @@ struct verify_forward_conv : conv_base<T, Tout>
         std::vector<char> ws;
         miopen::Allocator::ManageDataPtr ws_dev = nullptr;
 
-        const auto ctx     = ExecutionContext{&handle}.DetectRocm();
+        const auto ctx     = ExecutionContext{&handle};
         const auto problem = ConvProblemDescription{
             input.desc,
             weights.desc,
@@ -1035,7 +1032,7 @@ struct verify_backward_conv : conv_base<T>
         bool fallback_path_taken = false;
         std::size_t count        = 0;
 
-        const auto ctx     = ExecutionContext{&handle}.DetectRocm();
+        const auto ctx     = ExecutionContext{&handle};
         const auto problem = ConvProblemDescription{
             out.desc,
             weights.desc,
@@ -1405,7 +1402,7 @@ struct verify_backward_weights_conv : conv_base<T>
         bool fallback_path_taken = false;
         std::size_t count        = 0;
 
-        const auto ctx = ExecutionContext{&handle}.DetectRocm();
+        const auto ctx = ExecutionContext{&handle};
         const auto problem =
             ConvProblemDescription{filter.mode != miopenTranspose ? out.desc : input.desc,
                                    rweights.desc,
@@ -1666,7 +1663,7 @@ struct verify_forward_conv_int8 : conv_base<T>
         auto in_vpad_dev  = handle.Write(input_vpad.data);
         auto wei_vpad_dev = handle.Write(weights_vpad.data);
 
-        const auto ctx     = ExecutionContext{&handle}.DetectRocm();
+        const auto ctx     = ExecutionContext{&handle};
         const auto problem = ConvProblemDescription{
             is_transform ? weight_vpad_desc : weights.desc,
             is_transform ? input_vpad_desc : input.desc,
@@ -2275,7 +2272,6 @@ struct conv_driver : test_driver
                 };
 
                 auto ctx = miopen::ExecutionContext{&get_handle()};
-                ctx.DetectRocm();
 
                 bool skip_forward = false;
 
diff --git a/test/embed_sqlite.cpp b/test/embed_sqlite.cpp
index 8d23f7eef0..dfb663802e 100644
--- a/test/embed_sqlite.cpp
+++ b/test/embed_sqlite.cpp
@@ -69,7 +69,6 @@ struct EmbedSQLite : test_driver
         const auto problem = miopen::ProblemDescription{conv_problem};
         miopen::ConvolutionContext ctx{};
         ctx.SetStream(&handle);
-        ctx.DetectRocm();
         // Check PerfDb
         {
             // Get filename for the sys db
diff --git a/test/find_db.cpp b/test/find_db.cpp
index d67f0206b5..b5dd4b83d6 100644
--- a/test/find_db.cpp
+++ b/test/find_db.cpp
@@ -103,7 +103,7 @@ struct FindDbTest : test_driver
     {
         MIOPEN_LOG_I("Starting backward find-db test.");
 
-        const auto ctx = ExecutionContext{&handle}.DetectRocm();
+        const auto ctx = ExecutionContext{&handle};
         const auto problem =
             conv::ProblemDescription{y.desc, w.desc, x.desc, filter, conv::Direction::BackwardData};
         const auto workspace_size = filter.GetWorkSpaceSize(ctx, problem);
@@ -137,7 +137,7 @@ struct FindDbTest : test_driver
     {
         std::cout << "Starting forward find-db test." << std::endl;
 
-        const auto ctx = ExecutionContext{&handle}.DetectRocm();
+        const auto ctx = ExecutionContext{&handle};
         const auto problem =
             conv::ProblemDescription{x.desc, w.desc, y.desc, filter, conv::Direction::Forward};
         const auto workspace_size = filter.GetWorkSpaceSize(ctx, problem);
@@ -171,7 +171,7 @@ struct FindDbTest : test_driver
     {
         MIOPEN_LOG_I("Starting wrw find-db test.");
 
-        const auto ctx     = ExecutionContext{&handle}.DetectRocm();
+        const auto ctx     = ExecutionContext{&handle};
         const auto problem = conv::ProblemDescription{
             y.desc, w.desc, x.desc, filter, conv::Direction::BackwardWeights};
         const auto workspace_size = filter.GetWorkSpaceSize(ctx, problem);
diff --git a/test/gpu_conv.hpp b/test/gpu_conv.hpp
index 6be07d842d..240b191028 100644
--- a/test/gpu_conv.hpp
+++ b/test/gpu_conv.hpp
@@ -91,7 +91,6 @@ bool gpu_ref_convolution_fwd(const tensor<Tin>& input,
             input.desc, weights.desc, rout.desc, filter, miopen::conv::Direction::Forward};
         auto ctx = miopen::ConvolutionContext{};
         ctx.SetStream(&handle);
-        ctx.DetectRocm();
         if(naive_solver.IsApplicable(ctx, problem))
         {
             gpu_ref_used          = true;
@@ -128,7 +127,6 @@ bool gpu_ref_convolution_bwd(tensor<Tin>& input,
             output.desc, weights.desc, input.desc, filter, miopen::conv::Direction::BackwardData};
         auto ctx = miopen::ConvolutionContext{};
         ctx.SetStream(&handle);
-        ctx.DetectRocm();
         if(naive_solver.IsApplicable(ctx, problem))
         {
             gpu_ref_used          = true;
@@ -169,7 +167,6 @@ bool gpu_ref_convolution_wrw(const tensor<Tin>& input,
                                              miopen::conv::Direction::BackwardWeights};
         auto ctx = miopen::ConvolutionContext{};
         ctx.SetStream(&handle);
-        ctx.DetectRocm();
         if(naive_solver.IsApplicable(ctx, problem))
         {
             gpu_ref_used          = true;
diff --git a/test/gpu_nchw_nhwc_transpose.cpp b/test/gpu_nchw_nhwc_transpose.cpp
index 5c68e0887d..01898888d5 100644
--- a/test/gpu_nchw_nhwc_transpose.cpp
+++ b/test/gpu_nchw_nhwc_transpose.cpp
@@ -350,7 +350,6 @@ struct transpose_test : transpose_base
 
             miopen::ExecutionContext ctx;
             ctx.SetStream(&miopen::deref(this->handle));
-            ctx.DetectRocm();
             // ctx.SetupFloats();
 
             TRANSPOSE_SOL transpose_sol(ctx, to_miopen_data_type<T>::get(), n, c, h, w);
diff --git a/test/gtest/bad_fusion_plan.cpp b/test/gtest/bad_fusion_plan.cpp
index 67ae57aa13..80cd6bdd10 100644
--- a/test/gtest/bad_fusion_plan.cpp
+++ b/test/gtest/bad_fusion_plan.cpp
@@ -142,7 +142,6 @@ class TestFusionPlan
         Solver solv{};
         const auto fusion_problem = miopen::FusionDescription{&fusePlanDesc};
         auto fusion_ctx           = miopen::FusionContext{handle};
-        fusion_ctx.DetectRocm();
 
         return solv.IsApplicable(fusion_ctx, fusion_problem);
     }
diff --git a/test/gtest/cba_infer.cpp b/test/gtest/cba_infer.cpp
index b8c3af0573..8690824ff4 100644
--- a/test/gtest/cba_infer.cpp
+++ b/test/gtest/cba_infer.cpp
@@ -71,7 +71,6 @@ void RunSolver(miopen::FusionPlanDescriptor& fusePlanDesc,
     Solver solv{};
     const auto fusion_problem = miopen::FusionDescription{&fusePlanDesc};
     auto fusion_ctx           = miopen::FusionContext{handle};
-    fusion_ctx.DetectRocm();
     if(!solv.IsApplicable(fusion_ctx, fusion_problem))
     {
         test_skipped = true;
@@ -95,7 +94,6 @@ void RunTunableSolver(miopen::FusionPlanDescriptor& fusePlanDesc,
     Solver solv{};
     const auto fusion_problem = miopen::FusionDescription{&fusePlanDesc};
     auto fusion_ctx           = miopen::FusionContext{handle};
-    fusion_ctx.DetectRocm();
     if(!solv.IsApplicable(fusion_ctx, fusion_problem))
     {
         test_skipped = true;
diff --git a/test/gtest/group_conv_fwd.cpp b/test/gtest/group_conv_fwd.cpp
index 75a8ffa3fc..e5a1dbbea9 100644
--- a/test/gtest/group_conv_fwd.cpp
+++ b/test/gtest/group_conv_fwd.cpp
@@ -59,7 +59,6 @@ void SolverFwd(const miopen::TensorDescriptor& inputDesc,
     auto ctx = miopen::ConvolutionContext{};
 
     ctx.SetStream(&handle);
-    ctx.DetectRocm();
 
     if(!solv.IsApplicable(ctx, problem))
     {
diff --git a/test/gtest/kernel_tuning_net.cpp b/test/gtest/kernel_tuning_net.cpp
index 78b96ba403..0a8885486d 100644
--- a/test/gtest/kernel_tuning_net.cpp
+++ b/test/gtest/kernel_tuning_net.cpp
@@ -95,7 +95,6 @@ void TestParameterPredictionModel(miopen::ProblemDescription problem,
         GTEST_SKIP();
     miopen::ConvolutionContext ctx;
     ctx.SetStream(&handle);
-    ctx.DetectRocm();
     T perf_config;
     bool valid = false;
     perf_config.RunParmeterPredictionModel(ctx, problem, valid);
diff --git a/test/gtest/na_infer.cpp b/test/gtest/na_infer.cpp
index 8716a03348..f0d3407b2b 100644
--- a/test/gtest/na_infer.cpp
+++ b/test/gtest/na_infer.cpp
@@ -47,7 +47,6 @@ void RunSolver(miopen::FusionPlanDescriptor& fusePlanDesc,
     Solver solv{};
     const auto fusion_problem = miopen::FusionDescription{&fusePlanDesc};
     auto fusion_ctx           = miopen::FusionContext{handle};
-    fusion_ctx.DetectRocm();
     if(!solv.IsApplicable(fusion_ctx, fusion_problem))
     {
         test_skipped = true;
diff --git a/test/gtest/solver_convasm3x3u.cpp b/test/gtest/solver_convasm3x3u.cpp
index 87b75a6e4c..3f174039f2 100644
--- a/test/gtest/solver_convasm3x3u.cpp
+++ b/test/gtest/solver_convasm3x3u.cpp
@@ -52,7 +52,6 @@ void SolverFwd(const miopen::TensorDescriptor& inputDesc,
     auto ctx = miopen::ConvolutionContext{};
 
     ctx.SetStream(&handle);
-    ctx.DetectRocm();
 
     if(!solv.IsApplicable(ctx, problem))
     {
diff --git a/test/gtest/tuna_net.cpp b/test/gtest/tuna_net.cpp
index 527f44b6b9..288c4f0068 100644
--- a/test/gtest/tuna_net.cpp
+++ b/test/gtest/tuna_net.cpp
@@ -93,7 +93,6 @@ void TestSolverPredictionModel(miopen::ProblemDescription& problem, std::size_t
         GTEST_SKIP();
     miopen::ConvolutionContext ctx;
     ctx.SetStream(&handle);
-    ctx.DetectRocm();
     std::vector<std::size_t> solvers = miopen::ai::immed_mode::PredictSolver(problem, ctx, device);
     std::size_t solver =
         std::distance(solvers.begin(), std::max_element(solvers.begin(), solvers.end()));
diff --git a/test/tensor_reorder.cpp b/test/tensor_reorder.cpp
index 25bf2e4102..8d43cee49a 100644
--- a/test/tensor_reorder.cpp
+++ b/test/tensor_reorder.cpp
@@ -408,7 +408,6 @@ struct tensor_reorder_driver : tensor_reorder_base_driver
 
             miopen::ExecutionContext ctx;
             ctx.SetStream(&miopen::deref(this->handle));
-            ctx.DetectRocm();
             // ctx.SetupFloats();
             auto reorder_sol = MakeTensorReorderAttributes(ctx,
                                                            to_miopen_data_type<T>::get(),