[HotFix] revert #2580 and re-enable smoke tests

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",

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)

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")

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)

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}));
+    }
+};

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)