testing changes

src/eliminate_pad.cpp

@@ -44,8 +44,8 @@ void eliminate_pad::update_op(T,
     std::array<size_t, 2> new_pads{static_cast<size_t>(pads[2]), static_cast<size_t>(pads[3])};
 
     T op = any_cast<T>(ins->get_operator());
-    if(op.padding_mode != op::padding_mode_t::default_)
-        return;
+    // if(op.padding_mode != op::padding_mode_t::default_)
+    //     return;
     op.padding = new_pads;
 
     std::vector<instruction_ref> new_inputs{ins->inputs()};

src/include/migraphx/op/convolution.hpp

@@ -44,8 +44,8 @@ struct convolution
         const shape& input   = inputs.at(0);
         const shape& weights = inputs.at(1);
         auto t               = input.type();
-        if(padding_mode == default_)
-        {
+        // if(padding_mode == default_)
+        // {
             return {t,
                     {
                         input.lens()[0],
@@ -63,32 +63,32 @@ struct convolution
                                     stride[1] +
                                 1)),
                     }};
-        }
-        else if(padding_mode == same)
-        {
-            return {t,
-                    {input.lens()[0],
-                     weights.lens()[0],
-                     static_cast<std::size_t>(
-                         std::ceil(static_cast<double>(input.lens()[2]) / stride[0])),
-                     static_cast<std::size_t>(
-                         std::ceil(static_cast<double>(input.lens()[3]) / stride[1]))}};
-        }
-        else if(padding_mode == valid)
-        {
-            return {
-                t,
-                {input.lens()[0],
-                 weights.lens()[0],
-                 static_cast<std::size_t>(std::ceil(
-                     static_cast<double>(input.lens()[2] - weights.lens()[2] + 1) / stride[0])),
-                 static_cast<std::size_t>(std::ceil(
-                     static_cast<double>(input.lens()[3] - weights.lens()[3] + 1) / stride[1]))}};
-        }
-        else
-        {
-            MIGRAPHX_THROW("Invalid padding mode");
-        }
+        // }
+        // else if(padding_mode == same)
+        // {
+        //     return {t,
+        //             {input.lens()[0],
+        //              weights.lens()[0],
+        //              static_cast<std::size_t>(
+        //                  std::ceil(static_cast<double>(input.lens()[2]) / stride[0])),
+        //              static_cast<std::size_t>(
+        //                  std::ceil(static_cast<double>(input.lens()[3]) / stride[1]))}};
+        // }
+        // else if(padding_mode == valid)
+        // {
+        //     return {
+        //         t,
+        //         {input.lens()[0],
+        //          weights.lens()[0],
+        //          static_cast<std::size_t>(std::ceil(
+        //              static_cast<double>(input.lens()[2] - weights.lens()[2] + 1) / stride[0])),
+        //          static_cast<std::size_t>(std::ceil(
+        //              static_cast<double>(input.lens()[3] - weights.lens()[3] + 1) / stride[1]))}};
+        // }
+        // else
+        // {
+        //     MIGRAPHX_THROW("Invalid padding mode");
+        // }
     }
 };
 

src/include/migraphx/op/pooling.hpp

@@ -48,8 +48,8 @@ struct pooling
         assert(lengths[0] <= (input.lens()[2] + 2 * padding[0]));
         assert(lengths[1] <= (input.lens()[3] + 2 * padding[1]));
 
-        if(padding_mode == default_)
-        {
+        // if(padding_mode == default_)
+        // {
             return {t,
                     {
                         input.lens()[0],
@@ -65,34 +65,34 @@ struct pooling
                                 input.lens()[3] + 2 * padding[1] - lengths[1], stride[1]) +
                                 1)),
                     }};
-        }
-        else if(padding_mode == same)
-        {
-            return {t,
-                    {input.lens()[0],
-                     input.lens()[1],
-                     ceil_divide<std::size_t>(input.lens()[2], stride[0]),
-                     ceil_divide<std::size_t>(input.lens()[3], stride[1])}};
-        }
-        else if(padding_mode == valid)
-        {
-            return {
-                t,
-                {
-                    input.lens()[0],
-                    input.lens()[1],
-                    std::size_t(std::max<std::ptrdiff_t>(
-                        1,
-                        floor_divide<std::ptrdiff_t>(input.lens()[2] - lengths[0], stride[0]) + 1)),
-                    std::size_t(std::max<std::ptrdiff_t>(
-                        1,
-                        floor_divide<std::ptrdiff_t>(input.lens()[3] - lengths[1], stride[1]) + 1)),
-                }};
-        }
-        else
-        {
-            MIGRAPHX_THROW("Invalid padding mode");
-        }
+        // }
+        // else if(padding_mode == same)
+        // {
+        //     return {t,
+        //             {input.lens()[0],
+        //              input.lens()[1],
+        //              ceil_divide<std::size_t>(input.lens()[2], stride[0]),
+        //              ceil_divide<std::size_t>(input.lens()[3], stride[1])}};
+        // }
+        // else if(padding_mode == valid)
+        // {
+        //     return {
+        //         t,
+        //         {
+        //             input.lens()[0],
+        //             input.lens()[1],
+        //             std::size_t(std::max<std::ptrdiff_t>(
+        //                 1,
+        //                 floor_divide<std::ptrdiff_t>(input.lens()[2] - lengths[0], stride[0]) + 1)),
+        //             std::size_t(std::max<std::ptrdiff_t>(
+        //                 1,
+        //                 floor_divide<std::ptrdiff_t>(input.lens()[3] - lengths[1], stride[1]) + 1)),
+        //         }};
+        // }
+        // else
+        // {
+        //     MIGRAPHX_THROW("Invalid padding mode");
+        // }
     }
 };
 

src/targets/gpu/device/include/migraphx/gpu/device/types.hpp

@@ -64,9 +64,9 @@ host_type<T>* host_cast(T* x)
 }
 
 template <class T>
-device_type<T> device_cast(T x)
+device_type<T> device_cast(const T& x)
 {
-    return reinterpret_cast<device_type<T>>(x);
+    return reinterpret_cast<const device_type<T>&>(x);
 }
 
 template <class T>

src/targets/gpu/device/pad.cpp

@@ -4,6 +4,7 @@
 #include <migraphx/gpu/device/pad.hpp>
 #include <migraphx/gpu/device/tensor.hpp>
 #include <migraphx/gpu/device/launch.hpp>
+#include <migraphx/float_equal.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
@@ -14,28 +15,30 @@ argument
 pad(hipStream_t stream, argument result, argument arg1, float value, std::vector<std::int64_t> pads)
 {
     std::size_t nelements = arg1.get_shape().elements();
-    // if(value == std::numeric_limits<float>::lowest())
-    // {
-    //     visit_all(result)([&](auto output) {
-    //             auto* outptr      = output.data();
-    //                 gs_launch(stream, nelements)([=](auto i) {
-    //                     outptr[i] = std::numeric_limits<typename
-    //                     decltype(output)::value_type>::lowest();
-    //                 });
-    //             });
-    // }
+    if(float_equal(value,std::numeric_limits<float>::lowest()))
+    {
+        visit_all(result)([&](auto output) {
+                auto* outptr      = device_cast(output.data());
+                auto val = device_cast(std::numeric_limits<typename
+                        decltype(output)::value_type>::lowest());
+                
+                    gs_launch(stream, nelements)([=](auto i) {
+                        outptr[i] = val;
+                    });
+                });
+    }
 
-    // else
-    // {
-    //     visit_all(result)([&](auto output) {
-    //             auto* outptr      = output.data();
-    //                 gs_launch(stream, nelements)([=](auto i) {
-    //                     outptr[i] = static_cast<typename decltype(output)::value_type>(value);
-    //                 });
-    //             });
-    // }
+    else
+    {
+        visit_all(result)([&](auto output) {
+                auto* outptr      = device_cast(output.data());
+                    gs_launch(stream, nelements)([=](auto i) {
+                        outptr[i] = value;
+                    });
+                });
+    }
 
-    nary(stream, result)([=] { return value; });
+    // nary(stream, result)([=] { return value; });
     visit_all(result, arg1)([&](auto output, auto input) {
         visit_tensor_size(result.get_shape().lens().size(), [&](auto ndim) {
             std::size_t offsets[ndim];

src/targets/gpu/include/migraphx/gpu/softmax.hpp

@@ -34,7 +34,7 @@ struct miopen_softmax
         return migraphx::reflect(self.op, f);
     }
 
-    std::string name() const { return "gpu::softmax"; }
+    std::string name() const { return "gpu::miopen_softmax"; }
     shape compute_shape(const std::vector<shape>& inputs) const;
     argument
     compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;

src/targets/gpu/lowering.cpp

@@ -100,6 +100,7 @@ struct miopen_apply
         add_extend_op<miopen_contiguous, op::contiguous>("contiguous");
         add_extend_op<hip_concat, op::concat>("concat");
         add_extend_op<hip_softmax, op::softmax>("softmax");
+        // add_extend_op<miopen_softmax, op::softmax>("softmax");
         add_extend_op<hip_logsoftmax, op::logsoftmax>("logsoftmax");
         add_extend_op<hip_gather, op::gather>("gather");
         add_extend_op<hip_pad, op::pad>("pad");

src/tf/CMakeLists.txt

@@ -31,7 +31,7 @@ rocm_install_targets(
 
 add_executable(read_tf read_tf.cpp)
 rocm_clang_tidy_check(read_tf)
-target_link_libraries(read_tf migraphx_tf)
+target_link_libraries(read_tf migraphx_tf migraphx_cpu)
 
 if(MIGRAPHX_ENABLE_GPU)
 add_executable(verify_tf verify_tf.cpp)

src/tf/tf.cpp

@@ -323,7 +323,7 @@ struct tf_parser
             const std::string& pad_mode = attributes.at("padding").s();
             if(pad_mode.find("SAME") != std::string::npos)
             {
-                // op.padding_mode                 = op::padding_mode_t::same;
+                op.padding_mode                 = op::padding_mode_t::same;
                 std::vector<size_t> weight_dims = weights->get_shape().lens();
                 size_t weight_h                 = weight_dims[2];
                 size_t weight_w                 = weight_dims[3];
@@ -424,7 +424,7 @@ struct tf_parser
 
             if(pad_mode.find("SAME") != std::string::npos)
             {
-                // op.padding_mode = op::padding_mode_t::same;
+                op.padding_mode = op::padding_mode_t::same;
                 std::vector<size_t> weight_dims = weights->get_shape().lens();
                 size_t weight_h                 = weight_dims[2];
                 size_t weight_w                 = weight_dims[3];
@@ -609,17 +609,13 @@ struct tf_parser
             const std::string& pad_mode = attributes.at("padding").s();
             if(pad_mode.find("SAME") != std::string::npos)
             {
-                // op.padding_mode = op::padding_mode_t::same;
+                op.padding_mode = op::padding_mode_t::same;
                 auto input_dims = l0->get_shape().lens();
                 size_t input_h  = input_dims[2];
                 size_t input_w  = input_dims[3];
                 std::vector<int64_t> pads(input_dims.size());
                 calculate_padding(0, pads, input_h, op.stride[0], 1, op.lengths[0]);
                 calculate_padding(1, pads, input_w, op.stride[1], 1, op.lengths[1]);
-                // for(auto pad : pads)
-                // {
-                //     std::cout << pad << std::endl;
-                // }
 
                 if(pads[0] != pads[2] || pads[1] != pads[3])
                 {