From c567b2e964c01c576a6e670ee85a1ca3920ef4ce Mon Sep 17 00:00:00 2001
From: Jun Liu <Liu.Jun@amd.com>
Date: Mon, 18 Dec 2023 15:54:32 -0800
Subject: [PATCH 1/5] [HotFix] revert #2580 and re-enable smoke tests (#2616)

* Revert "[gtest] conversion for code coverage tests (#2580)"

This reverts commit c5a2384dc0f29682ed51aeccf9b981dbdf7e058f.

* re-enable smoke tests in CI

* remove problematic github action
---
 .github/workflows/update_develop_nightly.yml |  20 --
 Jenkinsfile                                  |   6 +-
 test/CMakeLists.txt                          |  22 ++
 test/driver.hpp                              |   7 +-
 test/gtest/CMakeLists.txt                    |   2 +-
 test/gtest/conv3d_codecov.cpp                | 172 ----------------
 test/gtest/conv_2d.hpp                       |  66 ++++++
 test/gtest/conv_igemm_mlir_xdlops.cpp        |   2 +-
 test/gtest/immed_conv2d_codecov.cpp          | 206 -------------------
 test/gtest/immed_conv3d_codecov.cpp          | 206 -------------------
 test/gtest/pooling2d_asymmetric.cpp          | 136 ------------
 test/gtest/pooling2d_codecov.cpp             | 136 ------------
 test/gtest/pooling2d_wide.cpp                | 136 ------------
 test/immed_conv2d.cpp                        |  62 ++++++
 test/{conv3d.hpp => immed_conv3d.cpp}        |  12 +-
 test/{pooling2d.hpp => pooling2d.cpp}        |   2 +
 16 files changed, 165 insertions(+), 1028 deletions(-)
 delete mode 100644 .github/workflows/update_develop_nightly.yml
 delete mode 100644 test/gtest/conv3d_codecov.cpp
 create mode 100644 test/gtest/conv_2d.hpp
 delete mode 100644 test/gtest/immed_conv2d_codecov.cpp
 delete mode 100644 test/gtest/immed_conv3d_codecov.cpp
 delete mode 100644 test/gtest/pooling2d_asymmetric.cpp
 delete mode 100644 test/gtest/pooling2d_codecov.cpp
 delete mode 100644 test/gtest/pooling2d_wide.cpp
 create mode 100644 test/immed_conv2d.cpp
 rename test/{conv3d.hpp => immed_conv3d.cpp} (91%)
 rename test/{pooling2d.hpp => pooling2d.cpp} (98%)

diff --git a/.github/workflows/update_develop_nightly.yml b/.github/workflows/update_develop_nightly.yml
deleted file mode 100644
index a8b5aabdb9..0000000000
--- a/.github/workflows/update_develop_nightly.yml
+++ /dev/null
@@ -1,20 +0,0 @@
-name: Sync branch
-
-on:
-  pull_request:
-  workflow_dispatch:
-  push:
-    branches:
-      - develop
-jobs:
-  sync-branch:
-    name: Update nightly branch
-    runs-on: ubuntu-latest
-    steps:
-      - name: Checkout repository
-        uses: actions/checkout@main
-      - uses: connor-baer/action-sync-branch@main
-        with:
-          branch: develop_nightly
-          token: ${{ secrets.GITHUB_TOKEN }}
-          force: false
\ No newline at end of file
diff --git a/Jenkinsfile b/Jenkinsfile
index 7e07d3deba..a103509f23 100644
--- a/Jenkinsfile
+++ b/Jenkinsfile
@@ -436,15 +436,15 @@ pipeline {
             description: "")
         booleanParam(
             name: "BUILD_SMOKE_FP32",
-            defaultValue: env.BRANCH_NAME == env.NIGHTLY_BRANCH ? true : false,
+            defaultValue: true,
             description: "")
         booleanParam(
             name: "BUILD_SMOKE_AUX1",
-            defaultValue: env.BRANCH_NAME == env.NIGHTLY_BRANCH ? true : false,
+            defaultValue: true,
             description: "")
         booleanParam(
             name: "BUILD_SMOKE_FP16_BF16_INT8",
-            defaultValue: env.BRANCH_NAME == env.NIGHTLY_BRANCH ? true : false,
+            defaultValue: true,
             description: "")
         booleanParam(
             name: "BUILD_FULL_TESTS",
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index ffaeb93228..d35ed0c43f 100644
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -673,6 +673,20 @@ function(add_custom_test NAME)
         set_tests_properties(${NAME} PROPERTIES DISABLED On)
     endif()
 endfunction()
+if(${CODECOV_TEST})
+    add_custom_test(test_conv3d_codecov
+        COMMAND $<TARGET_FILE:test_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 2 4 4 4 4 --weights 2 4 1 1 1 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
+    )
+    add_custom_test(test_immed_conv2d_codecov
+        COMMAND $<TARGET_FILE:test_immed_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --input  2 2 14 14 --weights 8 2 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
+    )
+    add_custom_test(test_immed_conv3d_codecov
+        COMMAND $<TARGET_FILE:test_immed_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 1 4 4 4 4 --weights 2 4 3 3 3 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
+    )
+    add_custom_test(test_pooling2d_codecov
+        COMMAND $<TARGET_FILE:test_pooling2d> ${MIOPEN_TEST_FLOAT_ARG} --input 1, 192, 28, 28 --lens 2 2 --strides 2 2 --pads 0 0 ${MIOPEN_TEST_FLAGS_ARGS}
+    )
+endif()
 
 if(${MIOPEN_TEST_WITH_MIOPENDRIVER})
     add_custom_test(test_miopendriver_regression_issue_1576 SKIP_UNLESS_ALL GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED
@@ -762,6 +776,14 @@ set(TEST_CONV_VERBOSE_F ${MIOPEN_TEST_FLOAT_ARG} --verbose --disable-backward-da
 set(TEST_CONV_VERBOSE_B ${MIOPEN_TEST_FLOAT_ARG} --verbose --disable-forward --disable-backward-weights)
 set(TEST_CONV_VERBOSE_W ${MIOPEN_TEST_FLOAT_ARG} --verbose --disable-forward --disable-backward-data)
 
+add_custom_test(test_pooling2d_asymmetric SKIP_UNLESS_ALL HALF_ENABLED GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED
+    COMMAND $<TARGET_FILE:test_pooling2d> ${MIOPEN_TEST_FLOAT_ARG} --all --dataset 1 --limit 0 ${MIOPEN_TEST_FLAGS_ARGS}
+)
+
+add_custom_test(test_pooling2d_wide SKIP_UNLESS_ALL HALF_ENABLED GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED
+    COMMAND $<TARGET_FILE:test_pooling2d> ${MIOPEN_TEST_FLOAT_ARG} --all --dataset 2 --limit 0 ${MIOPEN_TEST_FLAGS_ARGS}
+)
+
 set(IMPLICITGEMM_MLIR_ENV_F_XDLOPS ${IMPLICITGEMM_MLIR_ENV_BASE} MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmFwdXdlops)
 set(IMPLICITGEMM_MLIR_ENV_B_XDLOPS ${IMPLICITGEMM_MLIR_ENV_BASE} MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmBwdXdlops)
 set(IMPLICITGEMM_MLIR_ENV_W_XDLOPS ${IMPLICITGEMM_MLIR_ENV_BASE} MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmWrWXdlops)
diff --git a/test/driver.hpp b/test/driver.hpp
index 97896d6c5d..825a0d4511 100644
--- a/test/driver.hpp
+++ b/test/driver.hpp
@@ -1242,11 +1242,6 @@ void test_drive_impl_1(std::string program_name, std::vector<std::string> as)
     Driver d{};
     d.program_name = program_name;
 
-    std::cout << program_name << " ";
-    for(const auto& str : as)
-        std::cout << str << " ";
-    std::cout << std::endl;
-
     std::set<std::string> keywords{
         "--help", "-h", "--half", "--float", "--double", "--int8", "--bfloat16"};
     d.parse(keyword_set{keywords});
@@ -1386,7 +1381,7 @@ template <template <class...> class Driver>
 void test_drive(int argc, const char* argv[])
 {
     std::vector<std::string> as(argv + 1, argv + argc);
-    // as.emplace_back("--float");
+    as.emplace_back("--float");
     for(auto&& arg : as)
     {
         if(arg == "--half")
diff --git a/test/gtest/CMakeLists.txt b/test/gtest/CMakeLists.txt
index 1e3fb206f3..2e80fb5eac 100644
--- a/test/gtest/CMakeLists.txt
+++ b/test/gtest/CMakeLists.txt
@@ -30,7 +30,7 @@ function(add_gtest TEST_NAME)
     endif()
     if(NOT WIN32) # TODO: cannot run on Windows due to missing DLL dependencies
       # Enable CMake to discover the test binary
-      gtest_discover_tests(test_${TEST_NAME} PROPERTIES ENVIRONMENT "MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR};MIOPEN_TEST_FLOAT_ARG=${MIOPEN_TEST_FLOAT_ARG};MIOPEN_TEST_ALL=${MIOPEN_TEST_ALL};MIOPEN_TEST_MLIR=${MIOPEN_TEST_MLIR};MIOPEN_TEST_COMPOSABLEKERNEL=${MIOPEN_TEST_COMPOSABLEKERNEL};CODECOV_TEST=${CODECOV_TEST}")
+      gtest_discover_tests(test_${TEST_NAME} PROPERTIES ENVIRONMENT "MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR};MIOPEN_TEST_FLOAT_ARG=${MIOPEN_TEST_FLOAT_ARG};MIOPEN_TEST_ALL=${MIOPEN_TEST_ALL};MIOPEN_TEST_MLIR=${MIOPEN_TEST_MLIR};MIOPEN_TEST_COMPOSABLEKERNEL=${MIOPEN_TEST_COMPOSABLEKERNEL}")
     endif()
 
     if(WIN32)
diff --git a/test/gtest/conv3d_codecov.cpp b/test/gtest/conv3d_codecov.cpp
deleted file mode 100644
index f85b072722..0000000000
--- a/test/gtest/conv3d_codecov.cpp
+++ /dev/null
@@ -1,172 +0,0 @@
-/*******************************************************************************
- *
- * MIT License
- *
- * Copyright (c) 2023 Advanced Micro Devices, Inc.
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- *
- *******************************************************************************/
-#include "conv3d.hpp"
-#include "get_handle.hpp"
-#include <miopen/env.hpp>
-#include <gtest/gtest.h>
-
-MIOPEN_DECLARE_ENV_VAR_BOOL(CODECOV_TEST)
-MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLAGS_ARGS)
-
-static bool SkipTest(void) { return !miopen::IsEnabled(ENV(CODECOV_TEST)); }
-
-void GetArgs(const std::string& param, std::vector<std::string>& tokens)
-{
-    std::stringstream ss(param);
-    std::istream_iterator<std::string> begin(ss);
-    std::istream_iterator<std::string> end;
-    while(begin != end)
-        tokens.push_back(*begin++);
-}
-
-class Conv3dFloat : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-class Conv3dHalf : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-class Conv3dBFloat16 : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-class Conv3dInt8 : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-void Run3dDriver(miopenDataType_t prec)
-{
-
-    std::vector<std::string> params;
-    switch(prec)
-    {
-    case miopenHalf: params = Conv3dHalf::GetParam(); break;
-    case miopenBFloat16: params = Conv3dBFloat16::GetParam(); break;
-    case miopenFloat: params = Conv3dFloat::GetParam(); break;
-    case miopenInt8: params = Conv3dInt8::GetParam(); break;
-    case miopenFloat8:
-    case miopenBFloat8:
-    case miopenInt32:
-    case miopenDouble:
-        FAIL() << "miopenInt32, miopenDouble, miopenFloat8, miopenBFloat8 "
-                  "data type not supported by "
-                  "conv3d_codecov test";
-
-    default: params = Conv3dFloat::GetParam();
-    }
-
-    for(const auto& test_value : params)
-    {
-        std::vector<std::string> tokens;
-        GetArgs(test_value, tokens);
-        std::vector<const char*> ptrs;
-
-        std::transform(tokens.begin(), tokens.end(), std::back_inserter(ptrs), [](const auto& str) {
-            return str.data();
-        });
-
-        testing::internal::CaptureStderr();
-        test_drive<conv3d_driver>(ptrs.size(), ptrs.data());
-        auto capture = testing::internal::GetCapturedStderr();
-        std::cout << capture;
-    }
-};
-
-bool IsTestSupportedForDevice(const miopen::Handle& handle) { return true; }
-
-TEST_P(Conv3dFloat, FloatTest)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run3dDriver(miopenFloat);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-TEST_P(Conv3dHalf, HalfTest)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run3dDriver(miopenHalf);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-TEST_P(Conv3dBFloat16, BFloat16Test)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run3dDriver(miopenBFloat16);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-TEST_P(Conv3dInt8, Int8Test)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run3dDriver(miopenInt8);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-std::vector<std::string> GetTestCases(const std::string& precision)
-{
-    const auto& flag_arg = miopen::GetStringEnv(ENV(MIOPEN_TEST_FLAGS_ARGS));
-
-    const std::vector<std::string> test_cases = {
-        // clang-format off
-    {"test_conv3d " + precision + " --input 2 4 4 4 4 --weights 2 4 1 1 1 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 "+flag_arg}
-        // clang-format on
-    };
-
-    return test_cases;
-}
-
-INSTANTIATE_TEST_SUITE_P(Conv3D, Conv3dFloat, testing::Values(GetTestCases("--float")));
-
-INSTANTIATE_TEST_SUITE_P(Conv3D, Conv3dHalf, testing::Values(GetTestCases("--half")));
-
-INSTANTIATE_TEST_SUITE_P(Conv3D, Conv3dBFloat16, testing::Values(GetTestCases("--bfloat16")));
-
-INSTANTIATE_TEST_SUITE_P(Conv3D, Conv3dInt8, testing::Values(GetTestCases("--int8")));
diff --git a/test/gtest/conv_2d.hpp b/test/gtest/conv_2d.hpp
new file mode 100644
index 0000000000..df73453ef2
--- /dev/null
+++ b/test/gtest/conv_2d.hpp
@@ -0,0 +1,66 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2023 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+#include "conv_common.hpp"
+
+template <class T>
+struct conv2d_driver : conv_driver<T>
+{
+    conv2d_driver() : conv_driver<T>()
+    {
+        this->add(this->input_dims, "input");
+        this->add(this->weight_tensor_dims, "weights");
+        this->add(this->batch_size,
+                  "batch_size",
+                  this->generate_data_limited(this->get_batch_sizes(), 1));
+        this->add(this->input_channels,
+                  "input_channels",
+                  this->generate_data_limited(this->get_input_channels(), 1, {32}));
+        this->add(this->output_channels,
+                  "output_channels",
+                  this->generate_data_limited(this->get_output_channels(), 1, {64}));
+        this->add(this->spatial_dim_elements,
+                  "spatial_dim_elements",
+                  this->generate_data_limited(this->get_2d_spatial_dims(), 1, {28, 28}));
+        this->add(this->filter_dims,
+                  "filter_dims",
+                  this->generate_data_limited(this->get_2d_filter_dims(), 2, {3, 3}));
+        this->add(this->pads_strides_dilations,
+                  "pads_strides_dilations",
+                  this->generate_data_limited(this->get_2d_pads_strides_dilations(), 2));
+        this->add(this->trans_output_pads,
+                  "trans_output_pads",
+                  this->generate_data(this->get_2d_trans_output_pads()));
+        this->add(this->in_layout, "in_layout", this->generate_data({"NCHW"}));
+        this->add(this->fil_layout, "fil_layout", this->generate_data({"NCHW"}));
+        this->add(this->out_layout, "out_layout", this->generate_data({"NCHW"}));
+        this->add(this->deterministic, "deterministic", this->generate_data({false}));
+        this->add(this->tensor_vect, "tensor_vect", this->generate_data({0}));
+        this->add(this->vector_length, "vector_length", this->generate_data({1}));
+        // Only valid for int8 input and weights
+        this->add(this->output_type, "output_type", this->generate_data({"int32"}));
+        this->add(this->int8_vectorize, "int8_vectorize", this->generate_data({false}));
+    }
+};
diff --git a/test/gtest/conv_igemm_mlir_xdlops.cpp b/test/gtest/conv_igemm_mlir_xdlops.cpp
index 3d339d9c07..bcc9495204 100644
--- a/test/gtest/conv_igemm_mlir_xdlops.cpp
+++ b/test/gtest/conv_igemm_mlir_xdlops.cpp
@@ -4,7 +4,7 @@
 #include <gtest/gtest.h>
 #include <miopen/miopen.h>
 #include <miopen/env.hpp>
-#include "conv2d.hpp"
+#include "conv_2d.hpp"
 #include "get_handle.hpp"
 
 MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLOAT_ARG)
diff --git a/test/gtest/immed_conv2d_codecov.cpp b/test/gtest/immed_conv2d_codecov.cpp
deleted file mode 100644
index ba2a9de09c..0000000000
--- a/test/gtest/immed_conv2d_codecov.cpp
+++ /dev/null
@@ -1,206 +0,0 @@
-/*******************************************************************************
- *
- * MIT License
- *
- * Copyright (c) 2023 Advanced Micro Devices, Inc.
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- *
- *******************************************************************************/
-#include "conv_common.hpp"
-#include "get_handle.hpp"
-#include <miopen/env.hpp>
-#include <gtest/gtest.h>
-
-MIOPEN_DECLARE_ENV_VAR_BOOL(CODECOV_TEST)
-MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLAGS_ARGS)
-
-template <class T>
-struct conv2d_driver : conv_driver<T, ConvApi::Immediate>
-{
-    conv2d_driver() : conv_driver<T, ConvApi::Immediate>()
-    {
-        this->add(this->input_dims, "input");
-        this->add(this->weight_tensor_dims, "weights");
-        this->add(this->batch_size,
-                  "batch_size",
-                  this->generate_data_limited(this->get_batch_sizes(), 1, {16}));
-        this->add(this->input_channels,
-                  "input_channels",
-                  this->generate_data_limited(this->get_input_channels(), 1, {32}));
-        this->add(this->output_channels,
-                  "output_channels",
-                  this->generate_data_limited(this->get_output_channels(), 1, {32}));
-        this->add(this->spatial_dim_elements,
-                  "spatial_dim_elements",
-                  this->generate_data_limited(this->get_2d_spatial_dims(), 1, {56, 56}));
-        this->add(this->filter_dims,
-                  "filter_dims",
-                  this->generate_data_limited(this->get_2d_filter_dims(), 2, {3, 3}));
-        this->add(this->pads_strides_dilations,
-                  "pads_strides_dilations",
-                  this->generate_data_limited(this->get_2d_pads_strides_dilations(), 2));
-        this->add(this->trans_output_pads,
-                  "trans_output_pads",
-                  this->generate_data_limited(this->get_2d_trans_output_pads(), 1));
-        this->add(this->in_layout, "in_layout", this->generate_data({"NCHW"}));
-        this->add(this->fil_layout, "fil_layout", this->generate_data({"NCHW"}));
-        this->add(this->out_layout, "out_layout", this->generate_data({"NCHW"}));
-    }
-};
-
-static bool SkipTest(void) { return !miopen::IsEnabled(ENV(CODECOV_TEST)); }
-
-void GetArgs(const std::string& param, std::vector<std::string>& tokens)
-{
-    std::stringstream ss(param);
-    std::istream_iterator<std::string> begin(ss);
-    std::istream_iterator<std::string> end;
-    while(begin != end)
-        tokens.push_back(*begin++);
-}
-
-class Conv2dFloat : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-class Conv2dHalf : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-class Conv2dBFloat16 : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-class Conv2dInt8 : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-void Run2dDriver(miopenDataType_t prec)
-{
-
-    std::vector<std::string> params;
-    switch(prec)
-    {
-    case miopenHalf: params = Conv2dHalf::GetParam(); break;
-    case miopenBFloat16: params = Conv2dBFloat16::GetParam(); break;
-    case miopenFloat: params = Conv2dFloat::GetParam(); break;
-    case miopenInt8: params = Conv2dInt8::GetParam(); break;
-    case miopenFloat8:
-    case miopenBFloat8:
-    case miopenInt32:
-    case miopenDouble:
-        FAIL() << "miopenInt32, miopenDouble, miopenFloat8, miopenBFloat8 "
-                  "data type not supported by "
-                  "immed_conv2d_codecov test";
-
-    default: params = Conv2dFloat::GetParam();
-    }
-
-    for(const auto& test_value : params)
-    {
-        std::vector<std::string> tokens;
-        GetArgs(test_value, tokens);
-        std::vector<const char*> ptrs;
-
-        std::transform(tokens.begin(), tokens.end(), std::back_inserter(ptrs), [](const auto& str) {
-            return str.data();
-        });
-
-        testing::internal::CaptureStderr();
-        test_drive<conv2d_driver>(ptrs.size(), ptrs.data());
-        auto capture = testing::internal::GetCapturedStderr();
-        std::cout << capture;
-    }
-};
-
-bool IsTestSupportedForDevice(const miopen::Handle& handle) { return true; }
-
-TEST_P(Conv2dFloat, FloatTest)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run2dDriver(miopenFloat);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-TEST_P(Conv2dHalf, HalfTest)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run2dDriver(miopenHalf);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-TEST_P(Conv2dBFloat16, BFloat16Test)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run2dDriver(miopenBFloat16);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-TEST_P(Conv2dInt8, Int8Test)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run2dDriver(miopenInt8);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-std::vector<std::string> GetTestCases(const std::string& precision)
-{
-    const auto& flag_arg = miopen::GetStringEnv(ENV(MIOPEN_TEST_FLAGS_ARGS));
-
-    const std::vector<std::string> test_cases = {
-        // clang-format off
-    {"test_immed_conv2d " + precision + " --input  2 2 14 14 --weights 8 2 3 3 --pads_strides_dilations 0 0 1 1 1 1 "+flag_arg}
-        // clang-format on
-    };
-
-    return test_cases;
-}
-
-INSTANTIATE_TEST_SUITE_P(ImmedConv2D, Conv2dFloat, testing::Values(GetTestCases("--float")));
-
-INSTANTIATE_TEST_SUITE_P(ImmedConv2D, Conv2dHalf, testing::Values(GetTestCases("--half")));
-
-INSTANTIATE_TEST_SUITE_P(ImmedConv2D, Conv2dBFloat16, testing::Values(GetTestCases("--bfloat16")));
-
-INSTANTIATE_TEST_SUITE_P(ImmedConv2D, Conv2dInt8, testing::Values(GetTestCases("--int8")));
diff --git a/test/gtest/immed_conv3d_codecov.cpp b/test/gtest/immed_conv3d_codecov.cpp
deleted file mode 100644
index 3f22479d91..0000000000
--- a/test/gtest/immed_conv3d_codecov.cpp
+++ /dev/null
@@ -1,206 +0,0 @@
-/*******************************************************************************
- *
- * MIT License
- *
- * Copyright (c) 2023 Advanced Micro Devices, Inc.
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- *
- *******************************************************************************/
-#include "conv_common.hpp"
-#include "get_handle.hpp"
-#include <miopen/env.hpp>
-#include <gtest/gtest.h>
-
-MIOPEN_DECLARE_ENV_VAR_BOOL(CODECOV_TEST)
-MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLAGS_ARGS)
-
-template <class T>
-struct conv3d_driver : conv_driver<T, ConvApi::Immediate>
-{
-    conv3d_driver() : conv_driver<T, ConvApi::Immediate>()
-    {
-        this->add(this->input_dims, "input");
-        this->add(this->weight_tensor_dims, "weights");
-        this->add(this->batch_size,
-                  "batch_size",
-                  this->generate_data_limited(this->get_batch_sizes(), 1, {8}));
-        this->add(this->input_channels,
-                  "input_channels",
-                  this->generate_data_limited(this->get_input_channels(), 1, {2}));
-        this->add(this->output_channels,
-                  "output_channels",
-                  this->generate_data_limited(this->get_output_channels(), 1, {16}));
-        this->add(this->spatial_dim_elements,
-                  "spatial_dim_elements",
-                  this->generate_data_limited(this->get_3d_spatial_dims(), 1, {16, 16, 16}));
-        this->add(this->filter_dims,
-                  "filter_dims",
-                  this->generate_data_limited(this->get_3d_filter_dims(), 2, {5, 5, 5}));
-        this->add(this->pads_strides_dilations,
-                  "pads_strides_dilations",
-                  this->generate_data_limited(this->get_3d_pads_strides_dilations(), 2));
-        this->add(this->trans_output_pads,
-                  "trans_output_pads",
-                  this->generate_data_limited(this->get_3d_trans_output_pads(), 1));
-        this->add(this->in_layout, "in_layout", this->generate_data({"NCDHW"}));
-        this->add(this->fil_layout, "fil_layout", this->generate_data({"NCDHW"}));
-        this->add(this->out_layout, "out_layout", this->generate_data({"NCDHW"}));
-    }
-};
-
-static bool SkipTest(void) { return !miopen::IsEnabled(ENV(CODECOV_TEST)); }
-
-void GetArgs(const std::string& param, std::vector<std::string>& tokens)
-{
-    std::stringstream ss(param);
-    std::istream_iterator<std::string> begin(ss);
-    std::istream_iterator<std::string> end;
-    while(begin != end)
-        tokens.push_back(*begin++);
-}
-
-class Conv3dFloat : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-class Conv3dHalf : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-class Conv3dBFloat16 : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-class Conv3dInt8 : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-void Run3dDriver(miopenDataType_t prec)
-{
-
-    std::vector<std::string> params;
-    switch(prec)
-    {
-    case miopenHalf: params = Conv3dHalf::GetParam(); break;
-    case miopenBFloat16: params = Conv3dBFloat16::GetParam(); break;
-    case miopenFloat: params = Conv3dFloat::GetParam(); break;
-    case miopenInt8: params = Conv3dInt8::GetParam(); break;
-    case miopenFloat8:
-    case miopenBFloat8:
-    case miopenInt32:
-    case miopenDouble:
-        FAIL() << "miopenInt32, miopenDouble, miopenFloat8, miopenBFloat8 "
-                  "data type not supported by "
-                  "immed_conv3d_codecov test";
-
-    default: params = Conv3dFloat::GetParam();
-    }
-
-    for(const auto& test_value : params)
-    {
-        std::vector<std::string> tokens;
-        GetArgs(test_value, tokens);
-        std::vector<const char*> ptrs;
-
-        std::transform(tokens.begin(), tokens.end(), std::back_inserter(ptrs), [](const auto& str) {
-            return str.data();
-        });
-
-        testing::internal::CaptureStderr();
-        test_drive<conv3d_driver>(ptrs.size(), ptrs.data());
-        auto capture = testing::internal::GetCapturedStderr();
-        std::cout << capture;
-    }
-};
-
-bool IsTestSupportedForDevice(const miopen::Handle& handle) { return true; }
-
-TEST_P(Conv3dFloat, FloatTest)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run3dDriver(miopenFloat);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-TEST_P(Conv3dHalf, HalfTest)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run3dDriver(miopenHalf);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-TEST_P(Conv3dBFloat16, BFloat16Test)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run3dDriver(miopenBFloat16);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-TEST_P(Conv3dInt8, Int8Test)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run3dDriver(miopenInt8);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-std::vector<std::string> GetTestCases(const std::string& precision)
-{
-    const auto& flag_arg = miopen::GetStringEnv(ENV(MIOPEN_TEST_FLAGS_ARGS));
-
-    const std::vector<std::string> test_cases = {
-        // clang-format off
-    {"test_immed_conv3d " + precision + " --input 1 4 4 4 4 --weights 2 4 3 3 3 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 "+flag_arg}
-        // clang-format on
-    };
-
-    return test_cases;
-}
-
-INSTANTIATE_TEST_SUITE_P(ImmedConv3D, Conv3dFloat, testing::Values(GetTestCases("--float")));
-
-INSTANTIATE_TEST_SUITE_P(ImmedConv3D, Conv3dHalf, testing::Values(GetTestCases("--half")));
-
-INSTANTIATE_TEST_SUITE_P(ImmedConv3D, Conv3dBFloat16, testing::Values(GetTestCases("--bfloat16")));
-
-INSTANTIATE_TEST_SUITE_P(ImmedConv3D, Conv3dInt8, testing::Values(GetTestCases("--int8")));
diff --git a/test/gtest/pooling2d_asymmetric.cpp b/test/gtest/pooling2d_asymmetric.cpp
deleted file mode 100644
index 894210b813..0000000000
--- a/test/gtest/pooling2d_asymmetric.cpp
+++ /dev/null
@@ -1,136 +0,0 @@
-/*******************************************************************************
- *
- * MIT License
- *
- * Copyright (c) 2019 Advanced Micro Devices, Inc.
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- *
- *******************************************************************************/
-
-#include "pooling2d.hpp"
-#include "get_handle.hpp"
-#include <miopen/env.hpp>
-#include <gtest/gtest.h>
-
-MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL)
-MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLAGS_ARGS)
-
-static bool SkipTest(void) { return miopen::IsDisabled(ENV(MIOPEN_TEST_ALL)); }
-
-void GetArgs(const std::string& param, std::vector<std::string>& tokens)
-{
-    std::stringstream ss(param);
-    std::istream_iterator<std::string> begin(ss);
-    std::istream_iterator<std::string> end;
-    while(begin != end)
-        tokens.push_back(*begin++);
-}
-
-class Pooling2dFloat : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-class Pooling2dHalf : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-void Run2dDriver(miopenDataType_t prec)
-{
-
-    std::vector<std::string> params;
-    switch(prec)
-    {
-    case miopenFloat: params = Pooling2dFloat::GetParam(); break;
-    case miopenHalf: params = Pooling2dHalf::GetParam(); break;
-    case miopenBFloat16:
-    case miopenInt8:
-    case miopenFloat8:
-    case miopenBFloat8:
-    case miopenInt32:
-    case miopenDouble:
-        FAIL()
-            << "miopenBFloat16, miopenInt8, miopenInt32, miopenDouble, miopenFloat8, miopenBFloat8 "
-               "data type not supported by "
-               "immed_conv2d_codecov test";
-
-    default: params = Pooling2dFloat::GetParam();
-    }
-
-    for(const auto& test_value : params)
-    {
-        std::vector<std::string> tokens;
-        GetArgs(test_value, tokens);
-        std::vector<const char*> ptrs;
-
-        std::transform(tokens.begin(), tokens.end(), std::back_inserter(ptrs), [](const auto& str) {
-            return str.data();
-        });
-
-        testing::internal::CaptureStderr();
-        test_drive<pooling2d_driver>(ptrs.size(), ptrs.data());
-        auto capture = testing::internal::GetCapturedStderr();
-        std::cout << capture;
-    }
-};
-
-bool IsTestSupportedForDevice(const miopen::Handle& handle) { return true; }
-
-TEST_P(Pooling2dFloat, FloatTest)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run2dDriver(miopenFloat);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-TEST_P(Pooling2dHalf, HalfTest)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run2dDriver(miopenHalf);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-std::vector<std::string> GetTestCases(const std::string& precision)
-{
-    const auto& flag_arg = miopen::GetStringEnv(ENV(MIOPEN_TEST_FLAGS_ARGS));
-
-    const std::vector<std::string> test_cases = {
-        // clang-format off
-    {"test_pooling2d " + precision + " --all --dataset 1 --limit 0 "+flag_arg}
-        // clang-format on
-    };
-
-    return test_cases;
-}
-
-INSTANTIATE_TEST_SUITE_P(Pooling2D, Pooling2dFloat, testing::Values(GetTestCases("--float")));
-
-INSTANTIATE_TEST_SUITE_P(Pooling2D, Pooling2dHalf, testing::Values(GetTestCases("--half")));
diff --git a/test/gtest/pooling2d_codecov.cpp b/test/gtest/pooling2d_codecov.cpp
deleted file mode 100644
index 52daa50f3f..0000000000
--- a/test/gtest/pooling2d_codecov.cpp
+++ /dev/null
@@ -1,136 +0,0 @@
-/*******************************************************************************
- *
- * MIT License
- *
- * Copyright (c) 2019 Advanced Micro Devices, Inc.
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- *
- *******************************************************************************/
-
-#include "pooling2d.hpp"
-#include "get_handle.hpp"
-#include <miopen/env.hpp>
-#include <gtest/gtest.h>
-
-MIOPEN_DECLARE_ENV_VAR_BOOL(CODECOV_TEST)
-MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLAGS_ARGS)
-
-static bool SkipTest(void) { return !miopen::IsEnabled(ENV(CODECOV_TEST)); }
-
-void GetArgs(const std::string& param, std::vector<std::string>& tokens)
-{
-    std::stringstream ss(param);
-    std::istream_iterator<std::string> begin(ss);
-    std::istream_iterator<std::string> end;
-    while(begin != end)
-        tokens.push_back(*begin++);
-}
-
-class Pooling2dFloat : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-class Pooling2dHalf : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-void Run2dDriver(miopenDataType_t prec)
-{
-
-    std::vector<std::string> params;
-    switch(prec)
-    {
-    case miopenHalf: params = Pooling2dHalf::GetParam(); break;
-    case miopenFloat: params = Pooling2dFloat::GetParam(); break;
-    case miopenBFloat16:
-    case miopenInt8:
-    case miopenFloat8:
-    case miopenBFloat8:
-    case miopenInt32:
-    case miopenDouble:
-        FAIL()
-            << "miopenBFloat16, miopenInt8, miopenInt32, miopenDouble, miopenFloat8, miopenBFloat8 "
-               "data type not supported by "
-               "immed_conv2d_codecov test";
-
-    default: params = Pooling2dFloat::GetParam();
-    }
-
-    for(const auto& test_value : params)
-    {
-        std::vector<std::string> tokens;
-        GetArgs(test_value, tokens);
-        std::vector<const char*> ptrs;
-
-        std::transform(tokens.begin(), tokens.end(), std::back_inserter(ptrs), [](const auto& str) {
-            return str.data();
-        });
-
-        testing::internal::CaptureStderr();
-        test_drive<pooling2d_driver>(ptrs.size(), ptrs.data());
-        auto capture = testing::internal::GetCapturedStderr();
-        std::cout << capture;
-    }
-};
-
-bool IsTestSupportedForDevice(const miopen::Handle& handle) { return true; }
-
-TEST_P(Pooling2dFloat, FloatTest)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run2dDriver(miopenFloat);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-TEST_P(Pooling2dHalf, HalfTest)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run2dDriver(miopenHalf);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-std::vector<std::string> GetTestCases(const std::string& precision)
-{
-    const auto& flag_arg = miopen::GetStringEnv(ENV(MIOPEN_TEST_FLAGS_ARGS));
-
-    const std::vector<std::string> test_cases = {
-        // clang-format off
-    {"test_pooling2d " + precision + "--input 1, 192, 28, 28 --lens 2 2 --strides 2 2 --pads 0 0 "+flag_arg}
-        // clang-format on
-    };
-
-    return test_cases;
-}
-
-INSTANTIATE_TEST_SUITE_P(Pooling2D, Pooling2dFloat, testing::Values(GetTestCases("--float")));
-
-INSTANTIATE_TEST_SUITE_P(Pooling2D, Pooling2dHalf, testing::Values(GetTestCases("--half")));
diff --git a/test/gtest/pooling2d_wide.cpp b/test/gtest/pooling2d_wide.cpp
deleted file mode 100644
index 159da6f608..0000000000
--- a/test/gtest/pooling2d_wide.cpp
+++ /dev/null
@@ -1,136 +0,0 @@
-/*******************************************************************************
- *
- * MIT License
- *
- * Copyright (c) 2019 Advanced Micro Devices, Inc.
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- *
- *******************************************************************************/
-
-#include "pooling2d.hpp"
-#include "get_handle.hpp"
-#include <miopen/env.hpp>
-#include <gtest/gtest.h>
-
-MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL)
-MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLAGS_ARGS)
-
-static bool SkipTest(void) { return miopen::IsDisabled(ENV(MIOPEN_TEST_ALL)); }
-
-void GetArgs(const std::string& param, std::vector<std::string>& tokens)
-{
-    std::stringstream ss(param);
-    std::istream_iterator<std::string> begin(ss);
-    std::istream_iterator<std::string> end;
-    while(begin != end)
-        tokens.push_back(*begin++);
-}
-
-class Pooling2dFloat : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-class Pooling2dHalf : public testing::TestWithParam<std::vector<std::string>>
-{
-};
-
-void Run2dDriver(miopenDataType_t prec)
-{
-
-    std::vector<std::string> params;
-    switch(prec)
-    {
-    case miopenFloat: params = Pooling2dFloat::GetParam(); break;
-    case miopenHalf: params = Pooling2dHalf::GetParam(); break;
-    case miopenBFloat16:
-    case miopenInt8:
-    case miopenFloat8:
-    case miopenBFloat8:
-    case miopenInt32:
-    case miopenDouble:
-        FAIL()
-            << "miopenBFloat16, miopenInt8, miopenInt32, miopenDouble, miopenFloat8, miopenBFloat8 "
-               "data type not supported by "
-               "immed_conv2d_codecov test";
-
-    default: params = Pooling2dFloat::GetParam();
-    }
-
-    for(const auto& test_value : params)
-    {
-        std::vector<std::string> tokens;
-        GetArgs(test_value, tokens);
-        std::vector<const char*> ptrs;
-
-        std::transform(tokens.begin(), tokens.end(), std::back_inserter(ptrs), [](const auto& str) {
-            return str.data();
-        });
-
-        testing::internal::CaptureStderr();
-        test_drive<pooling2d_driver>(ptrs.size(), ptrs.data());
-        auto capture = testing::internal::GetCapturedStderr();
-        std::cout << capture;
-    }
-};
-
-bool IsTestSupportedForDevice(const miopen::Handle& handle) { return true; }
-
-TEST_P(Pooling2dFloat, FloatTest)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run2dDriver(miopenFloat);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-TEST_P(Pooling2dHalf, HalfTest)
-{
-    const auto& handle = get_handle();
-    if(IsTestSupportedForDevice(handle) && !SkipTest())
-    {
-        Run2dDriver(miopenHalf);
-    }
-    else
-    {
-        GTEST_SKIP();
-    }
-};
-
-std::vector<std::string> GetTestCases(const std::string precision)
-{
-    const auto& flag_arg = miopen::GetStringEnv(ENV(MIOPEN_TEST_FLAGS_ARGS));
-
-    const std::vector<std::string> test_cases = {
-        // clang-format off
-    {"test_pooling2d " + precision + " --all --dataset 2 --limit 0 "+flag_arg}
-        // clang-format on
-    };
-
-    return test_cases;
-}
-
-INSTANTIATE_TEST_SUITE_P(Pooling2D, Pooling2dFloat, testing::Values(GetTestCases("--float")));
-
-INSTANTIATE_TEST_SUITE_P(Pooling2D, Pooling2dHalf, testing::Values(GetTestCases("--half")));
diff --git a/test/immed_conv2d.cpp b/test/immed_conv2d.cpp
new file mode 100644
index 0000000000..6649c69d65
--- /dev/null
+++ b/test/immed_conv2d.cpp
@@ -0,0 +1,62 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2019 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+#include "conv_common.hpp"
+
+template <class T>
+struct conv2d_driver : conv_driver<T, ConvApi::Immediate>
+{
+    conv2d_driver() : conv_driver<T, ConvApi::Immediate>()
+    {
+        this->add(this->input_dims, "input");
+        this->add(this->weight_tensor_dims, "weights");
+        this->add(this->batch_size,
+                  "batch_size",
+                  this->generate_data_limited(this->get_batch_sizes(), 1, {16}));
+        this->add(this->input_channels,
+                  "input_channels",
+                  this->generate_data_limited(this->get_input_channels(), 1, {32}));
+        this->add(this->output_channels,
+                  "output_channels",
+                  this->generate_data_limited(this->get_output_channels(), 1, {32}));
+        this->add(this->spatial_dim_elements,
+                  "spatial_dim_elements",
+                  this->generate_data_limited(this->get_2d_spatial_dims(), 1, {56, 56}));
+        this->add(this->filter_dims,
+                  "filter_dims",
+                  this->generate_data_limited(this->get_2d_filter_dims(), 2, {3, 3}));
+        this->add(this->pads_strides_dilations,
+                  "pads_strides_dilations",
+                  this->generate_data_limited(this->get_2d_pads_strides_dilations(), 2));
+        this->add(this->trans_output_pads,
+                  "trans_output_pads",
+                  this->generate_data_limited(this->get_2d_trans_output_pads(), 1));
+        this->add(this->in_layout, "in_layout", this->generate_data({"NCHW"}));
+        this->add(this->fil_layout, "fil_layout", this->generate_data({"NCHW"}));
+        this->add(this->out_layout, "out_layout", this->generate_data({"NCHW"}));
+    }
+};
+
+int main(int argc, const char* argv[]) { test_drive<conv2d_driver>(argc, argv); }
diff --git a/test/conv3d.hpp b/test/immed_conv3d.cpp
similarity index 91%
rename from test/conv3d.hpp
rename to test/immed_conv3d.cpp
index 2b49c0f10e..a968cd6100 100644
--- a/test/conv3d.hpp
+++ b/test/immed_conv3d.cpp
@@ -26,9 +26,9 @@
 #include "conv_common.hpp"
 
 template <class T>
-struct conv3d_driver : conv_driver<T>
+struct conv3d_driver : conv_driver<T, ConvApi::Immediate>
 {
-    conv3d_driver() : conv_driver<T>()
+    conv3d_driver() : conv_driver<T, ConvApi::Immediate>()
     {
         this->add(this->input_dims, "input");
         this->add(this->weight_tensor_dims, "weights");
@@ -37,16 +37,16 @@ struct conv3d_driver : conv_driver<T>
                   this->generate_data_limited(this->get_batch_sizes(), 1, {8}));
         this->add(this->input_channels,
                   "input_channels",
-                  this->generate_data_limited(this->get_input_channels(), 1, {32}));
+                  this->generate_data_limited(this->get_input_channels(), 1, {2}));
         this->add(this->output_channels,
                   "output_channels",
-                  this->generate_data_limited(this->get_output_channels(), 1, {32}));
+                  this->generate_data_limited(this->get_output_channels(), 1, {16}));
         this->add(this->spatial_dim_elements,
                   "spatial_dim_elements",
                   this->generate_data_limited(this->get_3d_spatial_dims(), 1, {16, 16, 16}));
         this->add(this->filter_dims,
                   "filter_dims",
-                  this->generate_data_limited(this->get_3d_filter_dims(), 2, {3, 3, 3}));
+                  this->generate_data_limited(this->get_3d_filter_dims(), 2, {5, 5, 5}));
         this->add(this->pads_strides_dilations,
                   "pads_strides_dilations",
                   this->generate_data_limited(this->get_3d_pads_strides_dilations(), 2));
@@ -58,3 +58,5 @@ struct conv3d_driver : conv_driver<T>
         this->add(this->out_layout, "out_layout", this->generate_data({"NCDHW"}));
     }
 };
+
+int main(int argc, const char* argv[]) { test_drive<conv3d_driver>(argc, argv); }
diff --git a/test/pooling2d.hpp b/test/pooling2d.cpp
similarity index 98%
rename from test/pooling2d.hpp
rename to test/pooling2d.cpp
index 128e81cce2..c0df88c481 100644
--- a/test/pooling2d.hpp
+++ b/test/pooling2d.cpp
@@ -105,3 +105,5 @@ struct pooling2d_driver : pooling_driver<T>
         this->add(this->wsidx, "wsidx", this->generate_data({0, 1}));
     }
 };
+
+int main(int argc, const char* argv[]) { test_drive<pooling2d_driver>(argc, argv); }

From d0382e969e58a32a01afeffec841f9cb7bf058bc Mon Sep 17 00:00:00 2001
From: Artur Wojcik <artur.wojcik@outlook.com>
Date: Tue, 19 Dec 2023 00:58:02 +0100
Subject: [PATCH 2/5] [Windows] use find_package() for SQLite3 (#2564)

---
 CMakeLists.txt            | 3 +--
 src/CMakeLists.txt        | 7 +------
 src/sqlite/CMakeLists.txt | 2 +-
 3 files changed, 3 insertions(+), 9 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 2fd18fce13..3a15bf13e4 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -91,8 +91,7 @@ option(BUILD_SHARED_LIBS "Create shared libraries" ON)
 
 if(MIOPEN_ENABLE_SQLITE)
     # MIOpen now depends on SQLite as well
-    find_package(PkgConfig)
-    pkg_check_modules(SQLITE3 REQUIRED sqlite3)
+    find_package(SQLite3 REQUIRED)
 endif()
 find_package(BZip2)
 find_package(nlohmann_json 3.9.1 REQUIRED)
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 1d5548db7e..535ead20a6 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -829,12 +829,7 @@ endif()
 #######################################
 if(MIOPEN_ENABLE_SQLITE)
     # MIOpen depends on SQLite
-    target_include_directories(MIOpen SYSTEM PRIVATE ${SQLITE3_STATIC_INCLUDE_DIRS})
-    target_include_directories(MIOpen SYSTEM INTERFACE $<BUILD_INTERFACE:${SQLITE3_STATIC_INCLUDE_DIRS}>)
-    target_compile_options(MIOpen PRIVATE ${SQLITE3_STATIC_CFLAGS})
-    target_compile_options(MIOpen INTERFACE $<BUILD_INTERFACE:${SQLITE3_STATIC_CFLAGS}>)
-    target_link_libraries(MIOpen PRIVATE ${SQLITE3_STATIC_LDFLAGS})
-    target_link_libraries(MIOpen INTERFACE $<BUILD_INTERFACE:${SQLITE3_STATIC_LDFLAGS}>)
+    target_link_libraries(MIOpen PRIVATE SQLite::SQLite3)
 endif()
 ############################################################
 # MIOpen depends on librt for Boost.Interprocess
diff --git a/src/sqlite/CMakeLists.txt b/src/sqlite/CMakeLists.txt
index 823a44cb7a..2e3d69b55c 100644
--- a/src/sqlite/CMakeLists.txt
+++ b/src/sqlite/CMakeLists.txt
@@ -34,7 +34,7 @@ target_compile_options(sqlite_memvfs PRIVATE
         $<$<CXX_COMPILER_ID:Clang>:-Wno-everything>
         $<$<CXX_COMPILER_ID:GNU>:-w>)
 
-target_include_directories(sqlite_memvfs SYSTEM PRIVATE ${SQLITE3_STATIC_INCLUDE_DIRS})
+target_link_libraries(sqlite_memvfs PRIVATE SQLite::SQLite3)
 # set_target_properties(
 #     sqlite_memvfs
 #     PROPERTIES

From f9cfc56b42dedc732da28ad50680a712a0e4256d Mon Sep 17 00:00:00 2001
From: "dependabot[bot]" <49699333+dependabot[bot]@users.noreply.github.com>
Date: Mon, 18 Dec 2023 15:58:26 -0800
Subject: [PATCH 3/5] [Doc] Bump rocm-docs-core from 0.30.1 to 0.30.2 in
 /docs/sphinx (#2620)

Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.30.1 to 0.30.2.
- [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases)
- [Changelog](https://github.com/RadeonOpenCompute/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](https://github.com/RadeonOpenCompute/rocm-docs-core/compare/v0.30.1...v0.30.2)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  dependency-type: direct:production
  update-type: version-update:semver-patch
...

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
---
 docs/sphinx/requirements.in  | 2 +-
 docs/sphinx/requirements.txt | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/docs/sphinx/requirements.in b/docs/sphinx/requirements.in
index ee7a8d0d9e..5ab9e9085e 100644
--- a/docs/sphinx/requirements.in
+++ b/docs/sphinx/requirements.in
@@ -1 +1 @@
-rocm-docs-core==0.30.1
+rocm-docs-core==0.30.2
diff --git a/docs/sphinx/requirements.txt b/docs/sphinx/requirements.txt
index ce13fde525..454c8157d0 100644
--- a/docs/sphinx/requirements.txt
+++ b/docs/sphinx/requirements.txt
@@ -100,7 +100,7 @@ requests==2.31.0
     # via
     #   pygithub
     #   sphinx
-rocm-docs-core==0.30.1
+rocm-docs-core==0.30.2
     # via -r requirements.in
 smmap==5.0.0
     # via gitdb

From 700e2e7033c9ed3b37e75a006cc7f4f7d31f29e2 Mon Sep 17 00:00:00 2001
From: Artur Wojcik <artur.wojcik@outlook.com>
Date: Tue, 19 Dec 2023 00:59:58 +0100
Subject: [PATCH 4/5] [Windows] use official ZStd package from Facebook (#2565)

---
 cmake/Findzstd.cmake | 91 --------------------------------------------
 requirements.txt     |  1 +
 2 files changed, 1 insertion(+), 91 deletions(-)
 delete mode 100644 cmake/Findzstd.cmake

diff --git a/cmake/Findzstd.cmake b/cmake/Findzstd.cmake
deleted file mode 100644
index 43ea6f9b40..0000000000
--- a/cmake/Findzstd.cmake
+++ /dev/null
@@ -1,91 +0,0 @@
-################################################################################
-#
-# MIT License
-#
-# Copyright (c) 2023 Advanced Micro Devices, Inc.
-#
-# Permission is hereby granted, free of charge, to any person obtaining a copy
-# of this software and associated documentation files (the "Software"), to deal
-# in the Software without restriction, including without limitation the rights
-# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
-# copies of the Software, and to permit persons to whom the Software is
-# furnished to do so, subject to the following conditions:
-#
-# The above copyright notice and this permission notice shall be included in all
-# copies or substantial portions of the Software.
-#
-# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
-# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
-# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
-# SOFTWARE.
-#
-################################################################################
-
-# Try to find the zstd library
-#
-# If successful, the following variables will be defined:
-# zstd_INCLUDE_DIR
-# zstd_LIBRARY
-# zstd_STATIC_LIBRARY
-# zstd_FOUND
-#
-# Additionally, one of the following import targets will be defined:
-# zstd::libzstd_shared
-# zstd::libzstd_static
-
-if(MSVC)
-  set(zstd_STATIC_LIBRARY_SUFFIX "_static\\${CMAKE_STATIC_LIBRARY_SUFFIX}$")
-else()
-  set(zstd_STATIC_LIBRARY_SUFFIX "\\${CMAKE_STATIC_LIBRARY_SUFFIX}$")
-endif()
-
-find_path(zstd_INCLUDE_DIR NAMES zstd.h)
-find_library(zstd_LIBRARY NAMES zstd zstd_static)
-find_library(zstd_STATIC_LIBRARY NAMES
-  zstd_static
-  "${CMAKE_STATIC_LIBRARY_PREFIX}zstd${CMAKE_STATIC_LIBRARY_SUFFIX}")
-
-include(FindPackageHandleStandardArgs)
-find_package_handle_standard_args(
-    zstd DEFAULT_MSG
-    zstd_LIBRARY zstd_INCLUDE_DIR
-)
-
-if(zstd_FOUND)
-  if(zstd_LIBRARY MATCHES "${zstd_STATIC_LIBRARY_SUFFIX}$")
-    set(zstd_STATIC_LIBRARY "${zstd_LIBRARY}")
-  elseif (NOT TARGET zstd::libzstd_shared)
-    add_library(zstd::libzstd_shared SHARED IMPORTED)
-    if(MSVC)
-      # IMPORTED_LOCATION is the path to the DLL and IMPORTED_IMPLIB is the "library".
-      get_filename_component(zstd_DIRNAME "${zstd_LIBRARY}" DIRECTORY)
-      string(REGEX REPLACE "${CMAKE_INSTALL_LIBDIR}$" "${CMAKE_INSTALL_BINDIR}" zstd_DIRNAME "${zstd_DIRNAME}")
-      get_filename_component(zstd_BASENAME "${zstd_LIBRARY}" NAME)
-      string(REGEX REPLACE "\\${CMAKE_LINK_LIBRARY_SUFFIX}$" "${CMAKE_SHARED_LIBRARY_SUFFIX}" zstd_BASENAME "${zstd_BASENAME}")
-      set_target_properties(zstd::libzstd_shared PROPERTIES
-          INTERFACE_INCLUDE_DIRECTORIES "${zstd_INCLUDE_DIR}"
-          IMPORTED_LOCATION "${zstd_DIRNAME}/${zstd_BASENAME}"
-          IMPORTED_IMPLIB "${zstd_LIBRARY}")
-      unset(zstd_DIRNAME)
-      unset(zstd_BASENAME)
-    else()
-      set_target_properties(zstd::libzstd_shared PROPERTIES
-          INTERFACE_INCLUDE_DIRECTORIES "${zstd_INCLUDE_DIR}"
-          IMPORTED_LOCATION "${zstd_LIBRARY}")
-    endif()
-  endif()
-  if(zstd_STATIC_LIBRARY MATCHES "${zstd_STATIC_LIBRARY_SUFFIX}$" AND
-     NOT TARGET zstd::libzstd_static)
-    add_library(zstd::libzstd_static STATIC IMPORTED)
-    set_target_properties(zstd::libzstd_static PROPERTIES
-        INTERFACE_INCLUDE_DIRECTORIES "${zstd_INCLUDE_DIR}"
-        IMPORTED_LOCATION "${zstd_STATIC_LIBRARY}")
-  endif()
-endif()
-
-unset(zstd_STATIC_LIBRARY_SUFFIX)
-
-mark_as_advanced(zstd_INCLUDE_DIR zstd_LIBRARY zstd_STATIC_LIBRARY)
diff --git a/requirements.txt b/requirements.txt
index 2c81433131..8a81242667 100755
--- a/requirements.txt
+++ b/requirements.txt
@@ -1,5 +1,6 @@
 sqlite3@3.43.2 -DCMAKE_POSITION_INDEPENDENT_CODE=On
 boost@1.83 -DCMAKE_POSITION_INDEPENDENT_CODE=On --build -DCMAKE_CXX_FLAGS=" -std=c++14 -Wno-enum-constexpr-conversion -Wno-deprecated-builtins -Wno-deprecated-declarations "
+facebook/zstd@v1.4.5 -X subdir -DCMAKE_DIR=build/cmake
 ROCmSoftwarePlatform/half@10abd99e7815f0ca5d892f58dd7d15a23b7cf92c --build
 ROCmSoftwarePlatform/rocMLIR@rocm-5.5.0 -H sha256:a5f62769d28a73e60bc8d61022820f050e97c977c8f6f6275488db31512e1f42 -DBUILD_FAT_LIBROCKCOMPILER=1 -DCMAKE_IGNORE_PATH=/opt/conda/envs/py_3.9 -DCMAKE_IGNORE_PREFIX_PATH=/opt/conda
 nlohmann/json@v3.11.2 -DJSON_MultipleHeaders=ON -DJSON_BuildTests=Off

From ed002f5877789f135883dcf13c935fa29de22f4c Mon Sep 17 00:00:00 2001
From: Christopher Erb <Christopher.Erb@amd.com>
Date: Tue, 19 Dec 2023 10:26:27 -0600
Subject: [PATCH 5/5] Revert "[HotFix] revert #2580 and re-enable smoke tests
 (#2616)"

This reverts commit c567b2e964c01c576a6e670ee85a1ca3920ef4ce.
---
 .github/workflows/update_develop_nightly.yml |  20 ++
 Jenkinsfile                                  |   6 +-
 test/CMakeLists.txt                          |  22 --
 test/{immed_conv3d.cpp => conv3d.hpp}        |  12 +-
 test/driver.hpp                              |   7 +-
 test/gtest/CMakeLists.txt                    |   2 +-
 test/gtest/conv3d_codecov.cpp                | 172 ++++++++++++++++
 test/gtest/conv_2d.hpp                       |  66 ------
 test/gtest/conv_igemm_mlir_xdlops.cpp        |   2 +-
 test/gtest/immed_conv2d_codecov.cpp          | 206 +++++++++++++++++++
 test/gtest/immed_conv3d_codecov.cpp          | 206 +++++++++++++++++++
 test/gtest/pooling2d_asymmetric.cpp          | 136 ++++++++++++
 test/gtest/pooling2d_codecov.cpp             | 136 ++++++++++++
 test/gtest/pooling2d_wide.cpp                | 136 ++++++++++++
 test/immed_conv2d.cpp                        |  62 ------
 test/{pooling2d.cpp => pooling2d.hpp}        |   2 -
 16 files changed, 1028 insertions(+), 165 deletions(-)
 create mode 100644 .github/workflows/update_develop_nightly.yml
 rename test/{immed_conv3d.cpp => conv3d.hpp} (91%)
 create mode 100644 test/gtest/conv3d_codecov.cpp
 delete mode 100644 test/gtest/conv_2d.hpp
 create mode 100644 test/gtest/immed_conv2d_codecov.cpp
 create mode 100644 test/gtest/immed_conv3d_codecov.cpp
 create mode 100644 test/gtest/pooling2d_asymmetric.cpp
 create mode 100644 test/gtest/pooling2d_codecov.cpp
 create mode 100644 test/gtest/pooling2d_wide.cpp
 delete mode 100644 test/immed_conv2d.cpp
 rename test/{pooling2d.cpp => pooling2d.hpp} (98%)

diff --git a/.github/workflows/update_develop_nightly.yml b/.github/workflows/update_develop_nightly.yml
new file mode 100644
index 0000000000..a8b5aabdb9
--- /dev/null
+++ b/.github/workflows/update_develop_nightly.yml
@@ -0,0 +1,20 @@
+name: Sync branch
+
+on:
+  pull_request:
+  workflow_dispatch:
+  push:
+    branches:
+      - develop
+jobs:
+  sync-branch:
+    name: Update nightly branch
+    runs-on: ubuntu-latest
+    steps:
+      - name: Checkout repository
+        uses: actions/checkout@main
+      - uses: connor-baer/action-sync-branch@main
+        with:
+          branch: develop_nightly
+          token: ${{ secrets.GITHUB_TOKEN }}
+          force: false
\ No newline at end of file
diff --git a/Jenkinsfile b/Jenkinsfile
index a103509f23..7e07d3deba 100644
--- a/Jenkinsfile
+++ b/Jenkinsfile
@@ -436,15 +436,15 @@ pipeline {
             description: "")
         booleanParam(
             name: "BUILD_SMOKE_FP32",
-            defaultValue: true,
+            defaultValue: env.BRANCH_NAME == env.NIGHTLY_BRANCH ? true : false,
             description: "")
         booleanParam(
             name: "BUILD_SMOKE_AUX1",
-            defaultValue: true,
+            defaultValue: env.BRANCH_NAME == env.NIGHTLY_BRANCH ? true : false,
             description: "")
         booleanParam(
             name: "BUILD_SMOKE_FP16_BF16_INT8",
-            defaultValue: true,
+            defaultValue: env.BRANCH_NAME == env.NIGHTLY_BRANCH ? true : false,
             description: "")
         booleanParam(
             name: "BUILD_FULL_TESTS",
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index d35ed0c43f..ffaeb93228 100644
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -673,20 +673,6 @@ function(add_custom_test NAME)
         set_tests_properties(${NAME} PROPERTIES DISABLED On)
     endif()
 endfunction()
-if(${CODECOV_TEST})
-    add_custom_test(test_conv3d_codecov
-        COMMAND $<TARGET_FILE:test_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 2 4 4 4 4 --weights 2 4 1 1 1 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
-    )
-    add_custom_test(test_immed_conv2d_codecov
-        COMMAND $<TARGET_FILE:test_immed_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --input  2 2 14 14 --weights 8 2 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
-    )
-    add_custom_test(test_immed_conv3d_codecov
-        COMMAND $<TARGET_FILE:test_immed_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 1 4 4 4 4 --weights 2 4 3 3 3 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
-    )
-    add_custom_test(test_pooling2d_codecov
-        COMMAND $<TARGET_FILE:test_pooling2d> ${MIOPEN_TEST_FLOAT_ARG} --input 1, 192, 28, 28 --lens 2 2 --strides 2 2 --pads 0 0 ${MIOPEN_TEST_FLAGS_ARGS}
-    )
-endif()
 
 if(${MIOPEN_TEST_WITH_MIOPENDRIVER})
     add_custom_test(test_miopendriver_regression_issue_1576 SKIP_UNLESS_ALL GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED
@@ -776,14 +762,6 @@ set(TEST_CONV_VERBOSE_F ${MIOPEN_TEST_FLOAT_ARG} --verbose --disable-backward-da
 set(TEST_CONV_VERBOSE_B ${MIOPEN_TEST_FLOAT_ARG} --verbose --disable-forward --disable-backward-weights)
 set(TEST_CONV_VERBOSE_W ${MIOPEN_TEST_FLOAT_ARG} --verbose --disable-forward --disable-backward-data)
 
-add_custom_test(test_pooling2d_asymmetric SKIP_UNLESS_ALL HALF_ENABLED GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED
-    COMMAND $<TARGET_FILE:test_pooling2d> ${MIOPEN_TEST_FLOAT_ARG} --all --dataset 1 --limit 0 ${MIOPEN_TEST_FLAGS_ARGS}
-)
-
-add_custom_test(test_pooling2d_wide SKIP_UNLESS_ALL HALF_ENABLED GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED
-    COMMAND $<TARGET_FILE:test_pooling2d> ${MIOPEN_TEST_FLOAT_ARG} --all --dataset 2 --limit 0 ${MIOPEN_TEST_FLAGS_ARGS}
-)
-
 set(IMPLICITGEMM_MLIR_ENV_F_XDLOPS ${IMPLICITGEMM_MLIR_ENV_BASE} MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmFwdXdlops)
 set(IMPLICITGEMM_MLIR_ENV_B_XDLOPS ${IMPLICITGEMM_MLIR_ENV_BASE} MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmBwdXdlops)
 set(IMPLICITGEMM_MLIR_ENV_W_XDLOPS ${IMPLICITGEMM_MLIR_ENV_BASE} MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmWrWXdlops)
diff --git a/test/immed_conv3d.cpp b/test/conv3d.hpp
similarity index 91%
rename from test/immed_conv3d.cpp
rename to test/conv3d.hpp
index a968cd6100..2b49c0f10e 100644
--- a/test/immed_conv3d.cpp
+++ b/test/conv3d.hpp
@@ -26,9 +26,9 @@
 #include "conv_common.hpp"
 
 template <class T>
-struct conv3d_driver : conv_driver<T, ConvApi::Immediate>
+struct conv3d_driver : conv_driver<T>
 {
-    conv3d_driver() : conv_driver<T, ConvApi::Immediate>()
+    conv3d_driver() : conv_driver<T>()
     {
         this->add(this->input_dims, "input");
         this->add(this->weight_tensor_dims, "weights");
@@ -37,16 +37,16 @@ struct conv3d_driver : conv_driver<T, ConvApi::Immediate>
                   this->generate_data_limited(this->get_batch_sizes(), 1, {8}));
         this->add(this->input_channels,
                   "input_channels",
-                  this->generate_data_limited(this->get_input_channels(), 1, {2}));
+                  this->generate_data_limited(this->get_input_channels(), 1, {32}));
         this->add(this->output_channels,
                   "output_channels",
-                  this->generate_data_limited(this->get_output_channels(), 1, {16}));
+                  this->generate_data_limited(this->get_output_channels(), 1, {32}));
         this->add(this->spatial_dim_elements,
                   "spatial_dim_elements",
                   this->generate_data_limited(this->get_3d_spatial_dims(), 1, {16, 16, 16}));
         this->add(this->filter_dims,
                   "filter_dims",
-                  this->generate_data_limited(this->get_3d_filter_dims(), 2, {5, 5, 5}));
+                  this->generate_data_limited(this->get_3d_filter_dims(), 2, {3, 3, 3}));
         this->add(this->pads_strides_dilations,
                   "pads_strides_dilations",
                   this->generate_data_limited(this->get_3d_pads_strides_dilations(), 2));
@@ -58,5 +58,3 @@ struct conv3d_driver : conv_driver<T, ConvApi::Immediate>
         this->add(this->out_layout, "out_layout", this->generate_data({"NCDHW"}));
     }
 };
-
-int main(int argc, const char* argv[]) { test_drive<conv3d_driver>(argc, argv); }
diff --git a/test/driver.hpp b/test/driver.hpp
index 825a0d4511..97896d6c5d 100644
--- a/test/driver.hpp
+++ b/test/driver.hpp
@@ -1242,6 +1242,11 @@ void test_drive_impl_1(std::string program_name, std::vector<std::string> as)
     Driver d{};
     d.program_name = program_name;
 
+    std::cout << program_name << " ";
+    for(const auto& str : as)
+        std::cout << str << " ";
+    std::cout << std::endl;
+
     std::set<std::string> keywords{
         "--help", "-h", "--half", "--float", "--double", "--int8", "--bfloat16"};
     d.parse(keyword_set{keywords});
@@ -1381,7 +1386,7 @@ template <template <class...> class Driver>
 void test_drive(int argc, const char* argv[])
 {
     std::vector<std::string> as(argv + 1, argv + argc);
-    as.emplace_back("--float");
+    // as.emplace_back("--float");
     for(auto&& arg : as)
     {
         if(arg == "--half")
diff --git a/test/gtest/CMakeLists.txt b/test/gtest/CMakeLists.txt
index 2e80fb5eac..1e3fb206f3 100644
--- a/test/gtest/CMakeLists.txt
+++ b/test/gtest/CMakeLists.txt
@@ -30,7 +30,7 @@ function(add_gtest TEST_NAME)
     endif()
     if(NOT WIN32) # TODO: cannot run on Windows due to missing DLL dependencies
       # Enable CMake to discover the test binary
-      gtest_discover_tests(test_${TEST_NAME} PROPERTIES ENVIRONMENT "MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR};MIOPEN_TEST_FLOAT_ARG=${MIOPEN_TEST_FLOAT_ARG};MIOPEN_TEST_ALL=${MIOPEN_TEST_ALL};MIOPEN_TEST_MLIR=${MIOPEN_TEST_MLIR};MIOPEN_TEST_COMPOSABLEKERNEL=${MIOPEN_TEST_COMPOSABLEKERNEL}")
+      gtest_discover_tests(test_${TEST_NAME} PROPERTIES ENVIRONMENT "MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR};MIOPEN_TEST_FLOAT_ARG=${MIOPEN_TEST_FLOAT_ARG};MIOPEN_TEST_ALL=${MIOPEN_TEST_ALL};MIOPEN_TEST_MLIR=${MIOPEN_TEST_MLIR};MIOPEN_TEST_COMPOSABLEKERNEL=${MIOPEN_TEST_COMPOSABLEKERNEL};CODECOV_TEST=${CODECOV_TEST}")
     endif()
 
     if(WIN32)
diff --git a/test/gtest/conv3d_codecov.cpp b/test/gtest/conv3d_codecov.cpp
new file mode 100644
index 0000000000..f85b072722
--- /dev/null
+++ b/test/gtest/conv3d_codecov.cpp
@@ -0,0 +1,172 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2023 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+#include "conv3d.hpp"
+#include "get_handle.hpp"
+#include <miopen/env.hpp>
+#include <gtest/gtest.h>
+
+MIOPEN_DECLARE_ENV_VAR_BOOL(CODECOV_TEST)
+MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLAGS_ARGS)
+
+static bool SkipTest(void) { return !miopen::IsEnabled(ENV(CODECOV_TEST)); }
+
+void GetArgs(const std::string& param, std::vector<std::string>& tokens)
+{
+    std::stringstream ss(param);
+    std::istream_iterator<std::string> begin(ss);
+    std::istream_iterator<std::string> end;
+    while(begin != end)
+        tokens.push_back(*begin++);
+}
+
+class Conv3dFloat : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+class Conv3dHalf : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+class Conv3dBFloat16 : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+class Conv3dInt8 : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+void Run3dDriver(miopenDataType_t prec)
+{
+
+    std::vector<std::string> params;
+    switch(prec)
+    {
+    case miopenHalf: params = Conv3dHalf::GetParam(); break;
+    case miopenBFloat16: params = Conv3dBFloat16::GetParam(); break;
+    case miopenFloat: params = Conv3dFloat::GetParam(); break;
+    case miopenInt8: params = Conv3dInt8::GetParam(); break;
+    case miopenFloat8:
+    case miopenBFloat8:
+    case miopenInt32:
+    case miopenDouble:
+        FAIL() << "miopenInt32, miopenDouble, miopenFloat8, miopenBFloat8 "
+                  "data type not supported by "
+                  "conv3d_codecov test";
+
+    default: params = Conv3dFloat::GetParam();
+    }
+
+    for(const auto& test_value : params)
+    {
+        std::vector<std::string> tokens;
+        GetArgs(test_value, tokens);
+        std::vector<const char*> ptrs;
+
+        std::transform(tokens.begin(), tokens.end(), std::back_inserter(ptrs), [](const auto& str) {
+            return str.data();
+        });
+
+        testing::internal::CaptureStderr();
+        test_drive<conv3d_driver>(ptrs.size(), ptrs.data());
+        auto capture = testing::internal::GetCapturedStderr();
+        std::cout << capture;
+    }
+};
+
+bool IsTestSupportedForDevice(const miopen::Handle& handle) { return true; }
+
+TEST_P(Conv3dFloat, FloatTest)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run3dDriver(miopenFloat);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+TEST_P(Conv3dHalf, HalfTest)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run3dDriver(miopenHalf);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+TEST_P(Conv3dBFloat16, BFloat16Test)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run3dDriver(miopenBFloat16);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+TEST_P(Conv3dInt8, Int8Test)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run3dDriver(miopenInt8);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+std::vector<std::string> GetTestCases(const std::string& precision)
+{
+    const auto& flag_arg = miopen::GetStringEnv(ENV(MIOPEN_TEST_FLAGS_ARGS));
+
+    const std::vector<std::string> test_cases = {
+        // clang-format off
+    {"test_conv3d " + precision + " --input 2 4 4 4 4 --weights 2 4 1 1 1 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 "+flag_arg}
+        // clang-format on
+    };
+
+    return test_cases;
+}
+
+INSTANTIATE_TEST_SUITE_P(Conv3D, Conv3dFloat, testing::Values(GetTestCases("--float")));
+
+INSTANTIATE_TEST_SUITE_P(Conv3D, Conv3dHalf, testing::Values(GetTestCases("--half")));
+
+INSTANTIATE_TEST_SUITE_P(Conv3D, Conv3dBFloat16, testing::Values(GetTestCases("--bfloat16")));
+
+INSTANTIATE_TEST_SUITE_P(Conv3D, Conv3dInt8, testing::Values(GetTestCases("--int8")));
diff --git a/test/gtest/conv_2d.hpp b/test/gtest/conv_2d.hpp
deleted file mode 100644
index df73453ef2..0000000000
--- a/test/gtest/conv_2d.hpp
+++ /dev/null
@@ -1,66 +0,0 @@
-/*******************************************************************************
- *
- * MIT License
- *
- * Copyright (c) 2023 Advanced Micro Devices, Inc.
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- *
- *******************************************************************************/
-#include "conv_common.hpp"
-
-template <class T>
-struct conv2d_driver : conv_driver<T>
-{
-    conv2d_driver() : conv_driver<T>()
-    {
-        this->add(this->input_dims, "input");
-        this->add(this->weight_tensor_dims, "weights");
-        this->add(this->batch_size,
-                  "batch_size",
-                  this->generate_data_limited(this->get_batch_sizes(), 1));
-        this->add(this->input_channels,
-                  "input_channels",
-                  this->generate_data_limited(this->get_input_channels(), 1, {32}));
-        this->add(this->output_channels,
-                  "output_channels",
-                  this->generate_data_limited(this->get_output_channels(), 1, {64}));
-        this->add(this->spatial_dim_elements,
-                  "spatial_dim_elements",
-                  this->generate_data_limited(this->get_2d_spatial_dims(), 1, {28, 28}));
-        this->add(this->filter_dims,
-                  "filter_dims",
-                  this->generate_data_limited(this->get_2d_filter_dims(), 2, {3, 3}));
-        this->add(this->pads_strides_dilations,
-                  "pads_strides_dilations",
-                  this->generate_data_limited(this->get_2d_pads_strides_dilations(), 2));
-        this->add(this->trans_output_pads,
-                  "trans_output_pads",
-                  this->generate_data(this->get_2d_trans_output_pads()));
-        this->add(this->in_layout, "in_layout", this->generate_data({"NCHW"}));
-        this->add(this->fil_layout, "fil_layout", this->generate_data({"NCHW"}));
-        this->add(this->out_layout, "out_layout", this->generate_data({"NCHW"}));
-        this->add(this->deterministic, "deterministic", this->generate_data({false}));
-        this->add(this->tensor_vect, "tensor_vect", this->generate_data({0}));
-        this->add(this->vector_length, "vector_length", this->generate_data({1}));
-        // Only valid for int8 input and weights
-        this->add(this->output_type, "output_type", this->generate_data({"int32"}));
-        this->add(this->int8_vectorize, "int8_vectorize", this->generate_data({false}));
-    }
-};
diff --git a/test/gtest/conv_igemm_mlir_xdlops.cpp b/test/gtest/conv_igemm_mlir_xdlops.cpp
index bcc9495204..3d339d9c07 100644
--- a/test/gtest/conv_igemm_mlir_xdlops.cpp
+++ b/test/gtest/conv_igemm_mlir_xdlops.cpp
@@ -4,7 +4,7 @@
 #include <gtest/gtest.h>
 #include <miopen/miopen.h>
 #include <miopen/env.hpp>
-#include "conv_2d.hpp"
+#include "conv2d.hpp"
 #include "get_handle.hpp"
 
 MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLOAT_ARG)
diff --git a/test/gtest/immed_conv2d_codecov.cpp b/test/gtest/immed_conv2d_codecov.cpp
new file mode 100644
index 0000000000..ba2a9de09c
--- /dev/null
+++ b/test/gtest/immed_conv2d_codecov.cpp
@@ -0,0 +1,206 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2023 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+#include "conv_common.hpp"
+#include "get_handle.hpp"
+#include <miopen/env.hpp>
+#include <gtest/gtest.h>
+
+MIOPEN_DECLARE_ENV_VAR_BOOL(CODECOV_TEST)
+MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLAGS_ARGS)
+
+template <class T>
+struct conv2d_driver : conv_driver<T, ConvApi::Immediate>
+{
+    conv2d_driver() : conv_driver<T, ConvApi::Immediate>()
+    {
+        this->add(this->input_dims, "input");
+        this->add(this->weight_tensor_dims, "weights");
+        this->add(this->batch_size,
+                  "batch_size",
+                  this->generate_data_limited(this->get_batch_sizes(), 1, {16}));
+        this->add(this->input_channels,
+                  "input_channels",
+                  this->generate_data_limited(this->get_input_channels(), 1, {32}));
+        this->add(this->output_channels,
+                  "output_channels",
+                  this->generate_data_limited(this->get_output_channels(), 1, {32}));
+        this->add(this->spatial_dim_elements,
+                  "spatial_dim_elements",
+                  this->generate_data_limited(this->get_2d_spatial_dims(), 1, {56, 56}));
+        this->add(this->filter_dims,
+                  "filter_dims",
+                  this->generate_data_limited(this->get_2d_filter_dims(), 2, {3, 3}));
+        this->add(this->pads_strides_dilations,
+                  "pads_strides_dilations",
+                  this->generate_data_limited(this->get_2d_pads_strides_dilations(), 2));
+        this->add(this->trans_output_pads,
+                  "trans_output_pads",
+                  this->generate_data_limited(this->get_2d_trans_output_pads(), 1));
+        this->add(this->in_layout, "in_layout", this->generate_data({"NCHW"}));
+        this->add(this->fil_layout, "fil_layout", this->generate_data({"NCHW"}));
+        this->add(this->out_layout, "out_layout", this->generate_data({"NCHW"}));
+    }
+};
+
+static bool SkipTest(void) { return !miopen::IsEnabled(ENV(CODECOV_TEST)); }
+
+void GetArgs(const std::string& param, std::vector<std::string>& tokens)
+{
+    std::stringstream ss(param);
+    std::istream_iterator<std::string> begin(ss);
+    std::istream_iterator<std::string> end;
+    while(begin != end)
+        tokens.push_back(*begin++);
+}
+
+class Conv2dFloat : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+class Conv2dHalf : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+class Conv2dBFloat16 : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+class Conv2dInt8 : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+void Run2dDriver(miopenDataType_t prec)
+{
+
+    std::vector<std::string> params;
+    switch(prec)
+    {
+    case miopenHalf: params = Conv2dHalf::GetParam(); break;
+    case miopenBFloat16: params = Conv2dBFloat16::GetParam(); break;
+    case miopenFloat: params = Conv2dFloat::GetParam(); break;
+    case miopenInt8: params = Conv2dInt8::GetParam(); break;
+    case miopenFloat8:
+    case miopenBFloat8:
+    case miopenInt32:
+    case miopenDouble:
+        FAIL() << "miopenInt32, miopenDouble, miopenFloat8, miopenBFloat8 "
+                  "data type not supported by "
+                  "immed_conv2d_codecov test";
+
+    default: params = Conv2dFloat::GetParam();
+    }
+
+    for(const auto& test_value : params)
+    {
+        std::vector<std::string> tokens;
+        GetArgs(test_value, tokens);
+        std::vector<const char*> ptrs;
+
+        std::transform(tokens.begin(), tokens.end(), std::back_inserter(ptrs), [](const auto& str) {
+            return str.data();
+        });
+
+        testing::internal::CaptureStderr();
+        test_drive<conv2d_driver>(ptrs.size(), ptrs.data());
+        auto capture = testing::internal::GetCapturedStderr();
+        std::cout << capture;
+    }
+};
+
+bool IsTestSupportedForDevice(const miopen::Handle& handle) { return true; }
+
+TEST_P(Conv2dFloat, FloatTest)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run2dDriver(miopenFloat);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+TEST_P(Conv2dHalf, HalfTest)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run2dDriver(miopenHalf);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+TEST_P(Conv2dBFloat16, BFloat16Test)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run2dDriver(miopenBFloat16);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+TEST_P(Conv2dInt8, Int8Test)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run2dDriver(miopenInt8);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+std::vector<std::string> GetTestCases(const std::string& precision)
+{
+    const auto& flag_arg = miopen::GetStringEnv(ENV(MIOPEN_TEST_FLAGS_ARGS));
+
+    const std::vector<std::string> test_cases = {
+        // clang-format off
+    {"test_immed_conv2d " + precision + " --input  2 2 14 14 --weights 8 2 3 3 --pads_strides_dilations 0 0 1 1 1 1 "+flag_arg}
+        // clang-format on
+    };
+
+    return test_cases;
+}
+
+INSTANTIATE_TEST_SUITE_P(ImmedConv2D, Conv2dFloat, testing::Values(GetTestCases("--float")));
+
+INSTANTIATE_TEST_SUITE_P(ImmedConv2D, Conv2dHalf, testing::Values(GetTestCases("--half")));
+
+INSTANTIATE_TEST_SUITE_P(ImmedConv2D, Conv2dBFloat16, testing::Values(GetTestCases("--bfloat16")));
+
+INSTANTIATE_TEST_SUITE_P(ImmedConv2D, Conv2dInt8, testing::Values(GetTestCases("--int8")));
diff --git a/test/gtest/immed_conv3d_codecov.cpp b/test/gtest/immed_conv3d_codecov.cpp
new file mode 100644
index 0000000000..3f22479d91
--- /dev/null
+++ b/test/gtest/immed_conv3d_codecov.cpp
@@ -0,0 +1,206 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2023 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+#include "conv_common.hpp"
+#include "get_handle.hpp"
+#include <miopen/env.hpp>
+#include <gtest/gtest.h>
+
+MIOPEN_DECLARE_ENV_VAR_BOOL(CODECOV_TEST)
+MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLAGS_ARGS)
+
+template <class T>
+struct conv3d_driver : conv_driver<T, ConvApi::Immediate>
+{
+    conv3d_driver() : conv_driver<T, ConvApi::Immediate>()
+    {
+        this->add(this->input_dims, "input");
+        this->add(this->weight_tensor_dims, "weights");
+        this->add(this->batch_size,
+                  "batch_size",
+                  this->generate_data_limited(this->get_batch_sizes(), 1, {8}));
+        this->add(this->input_channels,
+                  "input_channels",
+                  this->generate_data_limited(this->get_input_channels(), 1, {2}));
+        this->add(this->output_channels,
+                  "output_channels",
+                  this->generate_data_limited(this->get_output_channels(), 1, {16}));
+        this->add(this->spatial_dim_elements,
+                  "spatial_dim_elements",
+                  this->generate_data_limited(this->get_3d_spatial_dims(), 1, {16, 16, 16}));
+        this->add(this->filter_dims,
+                  "filter_dims",
+                  this->generate_data_limited(this->get_3d_filter_dims(), 2, {5, 5, 5}));
+        this->add(this->pads_strides_dilations,
+                  "pads_strides_dilations",
+                  this->generate_data_limited(this->get_3d_pads_strides_dilations(), 2));
+        this->add(this->trans_output_pads,
+                  "trans_output_pads",
+                  this->generate_data_limited(this->get_3d_trans_output_pads(), 1));
+        this->add(this->in_layout, "in_layout", this->generate_data({"NCDHW"}));
+        this->add(this->fil_layout, "fil_layout", this->generate_data({"NCDHW"}));
+        this->add(this->out_layout, "out_layout", this->generate_data({"NCDHW"}));
+    }
+};
+
+static bool SkipTest(void) { return !miopen::IsEnabled(ENV(CODECOV_TEST)); }
+
+void GetArgs(const std::string& param, std::vector<std::string>& tokens)
+{
+    std::stringstream ss(param);
+    std::istream_iterator<std::string> begin(ss);
+    std::istream_iterator<std::string> end;
+    while(begin != end)
+        tokens.push_back(*begin++);
+}
+
+class Conv3dFloat : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+class Conv3dHalf : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+class Conv3dBFloat16 : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+class Conv3dInt8 : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+void Run3dDriver(miopenDataType_t prec)
+{
+
+    std::vector<std::string> params;
+    switch(prec)
+    {
+    case miopenHalf: params = Conv3dHalf::GetParam(); break;
+    case miopenBFloat16: params = Conv3dBFloat16::GetParam(); break;
+    case miopenFloat: params = Conv3dFloat::GetParam(); break;
+    case miopenInt8: params = Conv3dInt8::GetParam(); break;
+    case miopenFloat8:
+    case miopenBFloat8:
+    case miopenInt32:
+    case miopenDouble:
+        FAIL() << "miopenInt32, miopenDouble, miopenFloat8, miopenBFloat8 "
+                  "data type not supported by "
+                  "immed_conv3d_codecov test";
+
+    default: params = Conv3dFloat::GetParam();
+    }
+
+    for(const auto& test_value : params)
+    {
+        std::vector<std::string> tokens;
+        GetArgs(test_value, tokens);
+        std::vector<const char*> ptrs;
+
+        std::transform(tokens.begin(), tokens.end(), std::back_inserter(ptrs), [](const auto& str) {
+            return str.data();
+        });
+
+        testing::internal::CaptureStderr();
+        test_drive<conv3d_driver>(ptrs.size(), ptrs.data());
+        auto capture = testing::internal::GetCapturedStderr();
+        std::cout << capture;
+    }
+};
+
+bool IsTestSupportedForDevice(const miopen::Handle& handle) { return true; }
+
+TEST_P(Conv3dFloat, FloatTest)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run3dDriver(miopenFloat);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+TEST_P(Conv3dHalf, HalfTest)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run3dDriver(miopenHalf);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+TEST_P(Conv3dBFloat16, BFloat16Test)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run3dDriver(miopenBFloat16);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+TEST_P(Conv3dInt8, Int8Test)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run3dDriver(miopenInt8);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+std::vector<std::string> GetTestCases(const std::string& precision)
+{
+    const auto& flag_arg = miopen::GetStringEnv(ENV(MIOPEN_TEST_FLAGS_ARGS));
+
+    const std::vector<std::string> test_cases = {
+        // clang-format off
+    {"test_immed_conv3d " + precision + " --input 1 4 4 4 4 --weights 2 4 3 3 3 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 "+flag_arg}
+        // clang-format on
+    };
+
+    return test_cases;
+}
+
+INSTANTIATE_TEST_SUITE_P(ImmedConv3D, Conv3dFloat, testing::Values(GetTestCases("--float")));
+
+INSTANTIATE_TEST_SUITE_P(ImmedConv3D, Conv3dHalf, testing::Values(GetTestCases("--half")));
+
+INSTANTIATE_TEST_SUITE_P(ImmedConv3D, Conv3dBFloat16, testing::Values(GetTestCases("--bfloat16")));
+
+INSTANTIATE_TEST_SUITE_P(ImmedConv3D, Conv3dInt8, testing::Values(GetTestCases("--int8")));
diff --git a/test/gtest/pooling2d_asymmetric.cpp b/test/gtest/pooling2d_asymmetric.cpp
new file mode 100644
index 0000000000..894210b813
--- /dev/null
+++ b/test/gtest/pooling2d_asymmetric.cpp
@@ -0,0 +1,136 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2019 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+
+#include "pooling2d.hpp"
+#include "get_handle.hpp"
+#include <miopen/env.hpp>
+#include <gtest/gtest.h>
+
+MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL)
+MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLAGS_ARGS)
+
+static bool SkipTest(void) { return miopen::IsDisabled(ENV(MIOPEN_TEST_ALL)); }
+
+void GetArgs(const std::string& param, std::vector<std::string>& tokens)
+{
+    std::stringstream ss(param);
+    std::istream_iterator<std::string> begin(ss);
+    std::istream_iterator<std::string> end;
+    while(begin != end)
+        tokens.push_back(*begin++);
+}
+
+class Pooling2dFloat : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+class Pooling2dHalf : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+void Run2dDriver(miopenDataType_t prec)
+{
+
+    std::vector<std::string> params;
+    switch(prec)
+    {
+    case miopenFloat: params = Pooling2dFloat::GetParam(); break;
+    case miopenHalf: params = Pooling2dHalf::GetParam(); break;
+    case miopenBFloat16:
+    case miopenInt8:
+    case miopenFloat8:
+    case miopenBFloat8:
+    case miopenInt32:
+    case miopenDouble:
+        FAIL()
+            << "miopenBFloat16, miopenInt8, miopenInt32, miopenDouble, miopenFloat8, miopenBFloat8 "
+               "data type not supported by "
+               "immed_conv2d_codecov test";
+
+    default: params = Pooling2dFloat::GetParam();
+    }
+
+    for(const auto& test_value : params)
+    {
+        std::vector<std::string> tokens;
+        GetArgs(test_value, tokens);
+        std::vector<const char*> ptrs;
+
+        std::transform(tokens.begin(), tokens.end(), std::back_inserter(ptrs), [](const auto& str) {
+            return str.data();
+        });
+
+        testing::internal::CaptureStderr();
+        test_drive<pooling2d_driver>(ptrs.size(), ptrs.data());
+        auto capture = testing::internal::GetCapturedStderr();
+        std::cout << capture;
+    }
+};
+
+bool IsTestSupportedForDevice(const miopen::Handle& handle) { return true; }
+
+TEST_P(Pooling2dFloat, FloatTest)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run2dDriver(miopenFloat);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+TEST_P(Pooling2dHalf, HalfTest)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run2dDriver(miopenHalf);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+std::vector<std::string> GetTestCases(const std::string& precision)
+{
+    const auto& flag_arg = miopen::GetStringEnv(ENV(MIOPEN_TEST_FLAGS_ARGS));
+
+    const std::vector<std::string> test_cases = {
+        // clang-format off
+    {"test_pooling2d " + precision + " --all --dataset 1 --limit 0 "+flag_arg}
+        // clang-format on
+    };
+
+    return test_cases;
+}
+
+INSTANTIATE_TEST_SUITE_P(Pooling2D, Pooling2dFloat, testing::Values(GetTestCases("--float")));
+
+INSTANTIATE_TEST_SUITE_P(Pooling2D, Pooling2dHalf, testing::Values(GetTestCases("--half")));
diff --git a/test/gtest/pooling2d_codecov.cpp b/test/gtest/pooling2d_codecov.cpp
new file mode 100644
index 0000000000..52daa50f3f
--- /dev/null
+++ b/test/gtest/pooling2d_codecov.cpp
@@ -0,0 +1,136 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2019 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+
+#include "pooling2d.hpp"
+#include "get_handle.hpp"
+#include <miopen/env.hpp>
+#include <gtest/gtest.h>
+
+MIOPEN_DECLARE_ENV_VAR_BOOL(CODECOV_TEST)
+MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLAGS_ARGS)
+
+static bool SkipTest(void) { return !miopen::IsEnabled(ENV(CODECOV_TEST)); }
+
+void GetArgs(const std::string& param, std::vector<std::string>& tokens)
+{
+    std::stringstream ss(param);
+    std::istream_iterator<std::string> begin(ss);
+    std::istream_iterator<std::string> end;
+    while(begin != end)
+        tokens.push_back(*begin++);
+}
+
+class Pooling2dFloat : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+class Pooling2dHalf : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+void Run2dDriver(miopenDataType_t prec)
+{
+
+    std::vector<std::string> params;
+    switch(prec)
+    {
+    case miopenHalf: params = Pooling2dHalf::GetParam(); break;
+    case miopenFloat: params = Pooling2dFloat::GetParam(); break;
+    case miopenBFloat16:
+    case miopenInt8:
+    case miopenFloat8:
+    case miopenBFloat8:
+    case miopenInt32:
+    case miopenDouble:
+        FAIL()
+            << "miopenBFloat16, miopenInt8, miopenInt32, miopenDouble, miopenFloat8, miopenBFloat8 "
+               "data type not supported by "
+               "immed_conv2d_codecov test";
+
+    default: params = Pooling2dFloat::GetParam();
+    }
+
+    for(const auto& test_value : params)
+    {
+        std::vector<std::string> tokens;
+        GetArgs(test_value, tokens);
+        std::vector<const char*> ptrs;
+
+        std::transform(tokens.begin(), tokens.end(), std::back_inserter(ptrs), [](const auto& str) {
+            return str.data();
+        });
+
+        testing::internal::CaptureStderr();
+        test_drive<pooling2d_driver>(ptrs.size(), ptrs.data());
+        auto capture = testing::internal::GetCapturedStderr();
+        std::cout << capture;
+    }
+};
+
+bool IsTestSupportedForDevice(const miopen::Handle& handle) { return true; }
+
+TEST_P(Pooling2dFloat, FloatTest)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run2dDriver(miopenFloat);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+TEST_P(Pooling2dHalf, HalfTest)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run2dDriver(miopenHalf);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+std::vector<std::string> GetTestCases(const std::string& precision)
+{
+    const auto& flag_arg = miopen::GetStringEnv(ENV(MIOPEN_TEST_FLAGS_ARGS));
+
+    const std::vector<std::string> test_cases = {
+        // clang-format off
+    {"test_pooling2d " + precision + "--input 1, 192, 28, 28 --lens 2 2 --strides 2 2 --pads 0 0 "+flag_arg}
+        // clang-format on
+    };
+
+    return test_cases;
+}
+
+INSTANTIATE_TEST_SUITE_P(Pooling2D, Pooling2dFloat, testing::Values(GetTestCases("--float")));
+
+INSTANTIATE_TEST_SUITE_P(Pooling2D, Pooling2dHalf, testing::Values(GetTestCases("--half")));
diff --git a/test/gtest/pooling2d_wide.cpp b/test/gtest/pooling2d_wide.cpp
new file mode 100644
index 0000000000..159da6f608
--- /dev/null
+++ b/test/gtest/pooling2d_wide.cpp
@@ -0,0 +1,136 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2019 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+
+#include "pooling2d.hpp"
+#include "get_handle.hpp"
+#include <miopen/env.hpp>
+#include <gtest/gtest.h>
+
+MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL)
+MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLAGS_ARGS)
+
+static bool SkipTest(void) { return miopen::IsDisabled(ENV(MIOPEN_TEST_ALL)); }
+
+void GetArgs(const std::string& param, std::vector<std::string>& tokens)
+{
+    std::stringstream ss(param);
+    std::istream_iterator<std::string> begin(ss);
+    std::istream_iterator<std::string> end;
+    while(begin != end)
+        tokens.push_back(*begin++);
+}
+
+class Pooling2dFloat : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+class Pooling2dHalf : public testing::TestWithParam<std::vector<std::string>>
+{
+};
+
+void Run2dDriver(miopenDataType_t prec)
+{
+
+    std::vector<std::string> params;
+    switch(prec)
+    {
+    case miopenFloat: params = Pooling2dFloat::GetParam(); break;
+    case miopenHalf: params = Pooling2dHalf::GetParam(); break;
+    case miopenBFloat16:
+    case miopenInt8:
+    case miopenFloat8:
+    case miopenBFloat8:
+    case miopenInt32:
+    case miopenDouble:
+        FAIL()
+            << "miopenBFloat16, miopenInt8, miopenInt32, miopenDouble, miopenFloat8, miopenBFloat8 "
+               "data type not supported by "
+               "immed_conv2d_codecov test";
+
+    default: params = Pooling2dFloat::GetParam();
+    }
+
+    for(const auto& test_value : params)
+    {
+        std::vector<std::string> tokens;
+        GetArgs(test_value, tokens);
+        std::vector<const char*> ptrs;
+
+        std::transform(tokens.begin(), tokens.end(), std::back_inserter(ptrs), [](const auto& str) {
+            return str.data();
+        });
+
+        testing::internal::CaptureStderr();
+        test_drive<pooling2d_driver>(ptrs.size(), ptrs.data());
+        auto capture = testing::internal::GetCapturedStderr();
+        std::cout << capture;
+    }
+};
+
+bool IsTestSupportedForDevice(const miopen::Handle& handle) { return true; }
+
+TEST_P(Pooling2dFloat, FloatTest)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run2dDriver(miopenFloat);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+TEST_P(Pooling2dHalf, HalfTest)
+{
+    const auto& handle = get_handle();
+    if(IsTestSupportedForDevice(handle) && !SkipTest())
+    {
+        Run2dDriver(miopenHalf);
+    }
+    else
+    {
+        GTEST_SKIP();
+    }
+};
+
+std::vector<std::string> GetTestCases(const std::string precision)
+{
+    const auto& flag_arg = miopen::GetStringEnv(ENV(MIOPEN_TEST_FLAGS_ARGS));
+
+    const std::vector<std::string> test_cases = {
+        // clang-format off
+    {"test_pooling2d " + precision + " --all --dataset 2 --limit 0 "+flag_arg}
+        // clang-format on
+    };
+
+    return test_cases;
+}
+
+INSTANTIATE_TEST_SUITE_P(Pooling2D, Pooling2dFloat, testing::Values(GetTestCases("--float")));
+
+INSTANTIATE_TEST_SUITE_P(Pooling2D, Pooling2dHalf, testing::Values(GetTestCases("--half")));
diff --git a/test/immed_conv2d.cpp b/test/immed_conv2d.cpp
deleted file mode 100644
index 6649c69d65..0000000000
--- a/test/immed_conv2d.cpp
+++ /dev/null
@@ -1,62 +0,0 @@
-/*******************************************************************************
- *
- * MIT License
- *
- * Copyright (c) 2019 Advanced Micro Devices, Inc.
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- *
- *******************************************************************************/
-#include "conv_common.hpp"
-
-template <class T>
-struct conv2d_driver : conv_driver<T, ConvApi::Immediate>
-{
-    conv2d_driver() : conv_driver<T, ConvApi::Immediate>()
-    {
-        this->add(this->input_dims, "input");
-        this->add(this->weight_tensor_dims, "weights");
-        this->add(this->batch_size,
-                  "batch_size",
-                  this->generate_data_limited(this->get_batch_sizes(), 1, {16}));
-        this->add(this->input_channels,
-                  "input_channels",
-                  this->generate_data_limited(this->get_input_channels(), 1, {32}));
-        this->add(this->output_channels,
-                  "output_channels",
-                  this->generate_data_limited(this->get_output_channels(), 1, {32}));
-        this->add(this->spatial_dim_elements,
-                  "spatial_dim_elements",
-                  this->generate_data_limited(this->get_2d_spatial_dims(), 1, {56, 56}));
-        this->add(this->filter_dims,
-                  "filter_dims",
-                  this->generate_data_limited(this->get_2d_filter_dims(), 2, {3, 3}));
-        this->add(this->pads_strides_dilations,
-                  "pads_strides_dilations",
-                  this->generate_data_limited(this->get_2d_pads_strides_dilations(), 2));
-        this->add(this->trans_output_pads,
-                  "trans_output_pads",
-                  this->generate_data_limited(this->get_2d_trans_output_pads(), 1));
-        this->add(this->in_layout, "in_layout", this->generate_data({"NCHW"}));
-        this->add(this->fil_layout, "fil_layout", this->generate_data({"NCHW"}));
-        this->add(this->out_layout, "out_layout", this->generate_data({"NCHW"}));
-    }
-};
-
-int main(int argc, const char* argv[]) { test_drive<conv2d_driver>(argc, argv); }
diff --git a/test/pooling2d.cpp b/test/pooling2d.hpp
similarity index 98%
rename from test/pooling2d.cpp
rename to test/pooling2d.hpp
index c0df88c481..128e81cce2 100644
--- a/test/pooling2d.cpp
+++ b/test/pooling2d.hpp
@@ -105,5 +105,3 @@ struct pooling2d_driver : pooling_driver<T>
         this->add(this->wsidx, "wsidx", this->generate_data({0, 1}));
     }
 };
-
-int main(int argc, const char* argv[]) { test_drive<pooling2d_driver>(argc, argv); }