refactor(miopen): remove dead in-file code across driver, batch norm, tensor ops, and CMake

projects/miopen/driver/bn_driver.hpp

@@ -62,8 +62,6 @@
 #define RMSTOL_FP32 1e-4
 #define RMSTOL_FP16 2e-3
 
-#define MIO_DRIVER_BN_REFERENCE_COMPUTE_3D_AS_2D 1 // Resolves issue #1974
-
 // #define BN_RUNFOR_PROFILER
 
 template <typename TInput,

projects/miopen/driver/conv_driver.hpp

@@ -7,7 +7,6 @@
 #include "InputFlags.hpp"
 #include "conv_common.hpp"
 #include "driver.hpp"
-#include "mloConvHost.hpp"
 #include "random.hpp"
 #include "rocrand_wrapper.hpp"
 #include "tensor_driver.hpp"

projects/miopen/driver/driver.hpp

@@ -54,8 +54,6 @@ using bfloat8_fnuz = miopen_f8::hip_f8<miopen_f8::hip_f8_type::bf8>;
 #include <hip/hip_runtime_api.h>
 #include <functional>
 
-#define UNPACK_VEC4(v) (v[0]), (v[1]), (v[2]), (v[3])
-
 // Use values which are distinctively greater then miopenStatus_t,
 // so that these can be ORed with any miopen status code
 // without loss of information.
@@ -188,8 +186,7 @@ class GpumemTensor
 {
     std::unique_ptr<GPUMem> dev;
     tensor<Tgpu> host;
-    bool is_gpualloc         = false;
-    bool init_gpu_output_nan = false;
+    bool is_gpualloc = false;
 
 public:
     void SetGpuallocMode(bool v) { is_gpualloc = v; }

projects/miopen/driver/miopen_ConvBatchNormActivHost.hpp

@@ -219,161 +219,4 @@ void miopenActivationFwdHost(int neuron_type,
         c_res[i] = f(static_cast<Tref>(bot_ptr[i])); // f(data.at(i));
 }
 
-template <typename Tgpu /* the data type used in GPU computations (usually half) */,
-          typename Tref /* the data type used in CPU checkings (usually double) */>
-int miopenInferVerify(size_t size, const Tref* c_res, const Tgpu* top_ptr, Tref allowedEps)
-{
-    int match = 1;
-    for(size_t i = 0; i < size && match; i++)
-    {
-        Tref c_val     = c_res[i];
-        Tref g_val     = static_cast<Tref>(top_ptr[i]);
-        double err     = std::abs(c_val - g_val);
-        double err_rel = calculate_relative_error(c_val, g_val);
-        // if(err > 1e-6) printf("i: %d, cval: %f, gval: %f\n", i, c_val, g_val );
-        if((err > allowedEps && err_rel > allowedEps) || std::isnan(c_val) || std::isnan(g_val) ||
-           !std::isfinite(c_val) || !std::isfinite(g_val))
-        {
-            std::cout << "Difference in neuron layer: " << err << " too large at " << i
-                      << " c_v = " << c_val << " vs g_val = " << g_val
-                      << " tolerance = " << allowedEps << std::endl;
-            match = 0;
-        }
-    }
-
-    return (match);
-}
-
-template <typename Tgpu, typename Tref>
-int ConvForwardCPU(const std::vector<Tref>& in,
-                   std::vector<Tref>& outhost,
-                   const std::vector<Tgpu>& wei,
-                   const std::vector<Tgpu>& b,
-                   const int bias,
-                   miopenConvolutionDescriptor_t& convDesc,
-                   miopenTensorDescriptor_t& inputTensor,
-                   miopenTensorDescriptor_t& weightTensor,
-                   miopenTensorDescriptor_t& outputTensor)
-{
-
-    int in_n, in_c, in_h, in_w;
-    int in_nstride, in_cstride, in_hstride, in_wstride;
-    miopenDataType_t dt;
-    miopenGet4dTensorDescriptor(inputTensor,
-                                &dt,
-                                &in_n,
-                                &in_c,
-                                &in_h,
-                                &in_w,
-                                &in_nstride,
-                                &in_cstride,
-                                &in_hstride,
-                                &in_wstride);
-
-    int wei_n, wei_c, wei_h, wei_w;
-    int wei_nstride, wei_cstride, wei_hstride, wei_wstride;
-
-    miopenGet4dTensorDescriptor(weightTensor,
-                                &dt,
-                                &wei_n,
-                                &wei_c,
-                                &wei_h,
-                                &wei_w,
-                                &wei_nstride,
-                                &wei_cstride,
-                                &wei_hstride,
-                                &wei_wstride);
-
-    int out_n, out_c, out_h, out_w;
-    int out_nstride, out_cstride, out_hstride, out_wstride;
-    miopenGet4dTensorDescriptor(outputTensor,
-                                &dt,
-                                &out_n,
-                                &out_c,
-                                &out_h,
-                                &out_w,
-                                &out_nstride,
-                                &out_cstride,
-                                &out_hstride,
-                                &out_wstride);
-
-    int stride_h, stride_w, pad_h, pad_w, dilation_h, dilation_w;
-    miopenConvolutionMode_t mode;
-    miopenPaddingMode_t pmode = miopen::deref(convDesc).paddingMode;
-    miopenGetConvolutionDescriptor(
-        convDesc, &mode, &pad_h, &pad_w, &stride_h, &stride_w, &dilation_h, &dilation_w);
-
-    if(pmode == miopenPaddingSame)
-    {
-        pad_h = (in_h % stride_h == 0) ? (std::max((wei_h - stride_h), 0))
-                                       : (std::max((wei_h - (in_h % stride_h)), 0));
-        pad_w = (in_w % stride_w == 0) ? (std::max((wei_w - stride_w), 0))
-                                       : (std::max((wei_w - (in_w % stride_w)), 0));
-        pad_h /= 2;
-        pad_w /= 2;
-    }
-    else if(pmode == miopenPaddingValid)
-    {
-        pad_h = 0;
-        pad_w = 0;
-    }
-
-    if(out_h <= 0 || out_w <= 0)
-        throw std::runtime_error("Invalid Test Case: Check Output Dimension.");
-
-    miopenGet4dTensorDescriptor(weightTensor,
-                                &dt,
-                                &wei_n,
-                                &wei_c,
-                                &wei_h,
-                                &wei_w,
-                                &wei_nstride,
-                                &wei_cstride,
-                                &wei_hstride,
-                                &wei_wstride);
-
-    for(int o = 0; o < out_n; o++)
-    { // mini-batch size
-        for(int w = 0; w < out_c; w++)
-        { // out_channels (num filters)
-            for(int i = 0; i < out_h; i++)
-            { // output_height (from getforwardoutputdim())
-                int in_off_h = i * stride_h;
-                for(int j = 0; j < out_w; j++)
-                { // output_width (from getforwardoutputdim())
-                    Tgpu acc     = static_cast<Tgpu>(0.);
-                    int in_off_w = j * stride_w;
-                    for(int k = 0; k < in_c; k++)
-                    { // in_channels (RGB)
-                        for(int x = 0; x < wei_h; x++)
-                        {
-                            int in_x = in_off_h - pad_h + x * dilation_h;
-                            if(in_x >= 0 && in_x < in_h)
-                            {
-                                for(int y = 0; y < wei_w; y++)
-                                {
-                                    int in_y = in_off_w - pad_w + y * dilation_w;
-                                    if(in_y >= 0 && in_y < in_w)
-                                    {
-                                        acc += static_cast<Tgpu>(in.at(o * in_nstride +
-                                                                       k * in_cstride +
-                                                                       in_x * in_w + in_y)) *
-                                               static_cast<Tgpu>(wei.at(w * wei_nstride +
-                                                                        k * wei_cstride +
-                                                                        x * wei_hstride + y));
-                                    }
-                                }
-                            }
-                        }
-                    }
-                    acc = bias != 0 ? acc + static_cast<Tgpu>(b[w]) : acc;
-                    outhost[o * out_nstride + w * out_cstride + i * out_hstride + j] = acc;
-                }
-            }
-        }
-    }
-
-    return 0;
-}
-
 #endif