rename the operator fp_conversion to convert

src/include/migraphx/op/fp_conversion.hpp → src/include/migraphx/op/convert.hpp

-#ifndef MIGRAPHX_GUARD_OPERATORS_FP_CONVERSION_HPP
-#define MIGRAPHX_GUARD_OPERATORS_FP_CONVERSION_HPP
+#ifndef MIGRAPHX_GUARD_OPERATORS_CONVERT_HPP
+#define MIGRAPHX_GUARD_OPERATORS_CONVERT_HPP
 
 #include <array>
 #include <migraphx/op/binary.hpp>
@@ -17,7 +17,7 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace op {
 
-struct fp_conversion
+struct convert
 {
     shape::type_t targe_type = shape::half_type;
 
@@ -27,7 +27,7 @@ struct fp_conversion
         return pack(f(self.targe_type, "target_type"));
     }
 
-    std::string name() const { return "fp_conversion"; }
+    std::string name() const { return "convert"; }
     shape compute_shape(std::vector<shape> inputs) const
     {
         check_shapes{inputs, *this}.has(1);

src/include/migraphx/operators.hpp

@@ -14,6 +14,7 @@
 #include <migraphx/op/common.hpp>
 #include <migraphx/op/concat.hpp>
 #include <migraphx/op/contiguous.hpp>
+#include <migraphx/op/convert.hpp>
 #include <migraphx/op/convolution.hpp>
 #include <migraphx/op/cosh.hpp>
 #include <migraphx/op/cos.hpp>
@@ -22,7 +23,6 @@
 #include <migraphx/op/elu.hpp>
 #include <migraphx/op/exp.hpp>
 #include <migraphx/op/flatten.hpp>
-#include <migraphx/op/fp_conversion.hpp>
 #include <migraphx/op/gather.hpp>
 #include <migraphx/op/gru.hpp>
 #include <migraphx/op/identity.hpp>

src/quantization.cpp

@@ -2,7 +2,7 @@
 #include <migraphx/program.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/iterator_for.hpp>
-#include <migraphx/op/fp_conversion.hpp>
+#include <migraphx/op/convert.hpp>
 #include <migraphx/stringutils.hpp>
 #include <migraphx/ranges.hpp>
 #include <utility>
@@ -35,11 +35,11 @@ instruction_ref insert_fp16(program& prog,
     {
         if(ins == std::prev(prog.end()))
         {
-            ins_fp16 = prog.add_instruction(op::fp_conversion{type}, ins);
+            ins_fp16 = prog.add_instruction(op::convert{type}, ins);
         }
         else
         {
-            ins_fp16 = prog.insert_instruction(std::next(ins), op::fp_conversion{}, ins);
+            ins_fp16 = prog.insert_instruction(std::next(ins), op::convert{}, ins);
         }
     }
     map_fp16[ins] = ins_fp16;
@@ -60,7 +60,7 @@ void quantize(program& prog, const std::vector<std::string>& ins_names)
 
         shape::type_t orig_type = ins->get_shape().type();
         // process all inputs, if input is a fp32 or fp64, convert it
-        // to a fp16 by adding a fp_conversion operator.
+        // to a fp16 by adding a convert operator.
         auto inputs = ins->inputs();
         std::vector<instruction_ref> converted_inputs;
         for(auto input : inputs)
@@ -68,10 +68,10 @@ void quantize(program& prog, const std::vector<std::string>& ins_names)
             auto s = input->get_shape();
             if(s.type() == shape::float_type || s.type() == shape::double_type)
             {
-                // if the input is a fp_conversion operator, uses its input
+                // if the input is a convert operator, uses its input
                 // as its current input
                 instruction_ref input_fp16{};
-                if(input->name() == "fp_conversion")
+                if(input->name() == "convert")
                 {
                     input_fp16 = input->inputs().front();
                 }
@@ -94,15 +94,15 @@ void quantize(program& prog, const std::vector<std::string>& ins_names)
             auto ins_shape = compute_shape(op, converted_inputs);
             if(ins_shape.type() != orig_type)
             {
-                // insert another fp_conversion instruction to convert it back
+                // insert another convert instruction to convert it back
                 if(ins == std::prev(prog.end()))
                 {
-                    prog.add_instruction(op::fp_conversion{orig_type}, ins);
+                    prog.add_instruction(op::convert{orig_type}, ins);
                 }
                 else
                 {
                     auto ins_orig_type =
-                        prog.insert_instruction(std::next(ins), op::fp_conversion{orig_type}, ins);
+                        prog.insert_instruction(std::next(ins), op::convert{orig_type}, ins);
                     prog.replace_instruction(ins, ins_orig_type);
                 }
             }

src/targets/cpu/lowering.cpp

@@ -727,10 +727,10 @@ struct cpu_logsoftmax
     }
 };
 
-struct cpu_fp_conversion
+struct cpu_convert
 {
-    op::fp_conversion op;
-    std::string name() const { return "cpu_fp_conversion"; }
+    op::convert op;
+    std::string name() const { return "cpu_convert"; }
     shape compute_shape(const std::vector<shape>& inputs) const { return op.compute_shape(inputs); }
 
     argument compute(context&, const shape& output_shape, std::vector<argument> args) const
@@ -853,7 +853,7 @@ struct cpu_apply
         apply_map["pad"]           = extend_op<cpu_pad, op::pad>();
         apply_map["concat"]        = extend_op<cpu_concat, op::concat>();
         apply_map["gather"]        = extend_op<cpu_gather, op::gather>();
-        apply_map["fp_conversion"] = extend_op<cpu_fp_conversion, op::fp_conversion>();
+        apply_map["convert"] = extend_op<cpu_convert, op::convert>();
         apply_map["logsoftmax"]    = extend_op<cpu_logsoftmax, op::logsoftmax>();
         apply_map["leaky_relu"]    = extend_op<cpu_unary<leaky_relu_op>, op::leaky_relu>();
         apply_map["elu"]           = extend_op<cpu_unary<elu_op>, op::elu>();

src/targets/gpu/CMakeLists.txt

@@ -27,7 +27,7 @@ add_library(migraphx_device
     device/add_relu.cpp
     device/contiguous.cpp
     device/logsoftmax.cpp
-    device/fp_conversion.cpp
+    device/convert.cpp
     device/mul.cpp
     device/concat.cpp
     device/pad.cpp
@@ -51,7 +51,7 @@ add_library(migraphx_gpu
     convolution.cpp
     softmax.cpp
     logsoftmax.cpp
-    fp_conversion.cpp
+    convert.cpp
     contiguous.cpp
     concat.cpp
     relu.cpp

src/targets/gpu/fp_conversion.cpp → src/targets/gpu/convert.cpp

-#include <migraphx/gpu/fp_conversion.hpp>
-#include <migraphx/gpu/device/fp_conversion.hpp>
+#include <migraphx/gpu/convert.hpp>
+#include <migraphx/gpu/device/convert.hpp>
 #include <migraphx/gpu/context.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-shape hip_fp_conversion::compute_shape(std::vector<shape> inputs) const
+shape hip_convert::compute_shape(std::vector<shape> inputs) const
 {
     inputs.pop_back();
     check_shapes{inputs}.packed();
     return op.compute_shape(inputs);
 }
 
-argument hip_fp_conversion::compute(context& ctx,
+argument hip_convert::compute(context& ctx,
                                     const shape& output_shape,
                                     const std::vector<argument>& args) const
 {
-    return device::fp_conversion(ctx.get_stream().get(), output_shape, args);
+    return device::convert(ctx.get_stream().get(), args[1], args[0]);
 }
 
 } // namespace gpu

src/targets/gpu/device/fp_conversion.cpp → src/targets/gpu/device/convert.cpp

-#include <migraphx/gpu/device/fp_conversion.hpp>
+#include <migraphx/gpu/device/convert.hpp>
 #include <migraphx/gpu/device/nary.hpp>
 
 namespace migraphx {
@@ -7,18 +7,18 @@ namespace gpu {
 namespace device {
 
 argument
-fp_conversion(hipStream_t stream, const shape& output_shape, const std::vector<argument>& args)
+convert(hipStream_t stream, const argument& result, const argument& arg)
 {
-    args.back().visit([&](auto output) {
-        args.front().visit([&](auto input) {
+    result.visit([&](auto output) {
+        arg.visit([&](auto input) {
             const auto* input_ptr = device_cast(input.data());
             auto* output_ptr      = device_cast(output.data());
             gs_launch(stream,
-                      output_shape.elements())([=](auto i) { output_ptr[i] = input_ptr[i]; });
+                      result.get_shape().elements())([=](auto i) { output_ptr[i] = input_ptr[i]; });
         });
     });
 
-    return args.back();
+    return result;
 }
 
 } // namespace device

src/targets/gpu/include/migraphx/gpu/fp_conversion.hpp → src/targets/gpu/include/migraphx/gpu/convert.hpp

-#ifndef MIGRAPHX_GUARD_RTGLIB_FP_CONVERSION_HPP
-#define MIGRAPHX_GUARD_RTGLIB_FP_CONVERSION_HPP
+#ifndef MIGRAPHX_GUARD_RTGLIB_CONVERT_HPP
+#define MIGRAPHX_GUARD_RTGLIB_CONVERT_HPP
 
 #include <migraphx/shape.hpp>
-#include <migraphx/op/fp_conversion.hpp>
+#include <migraphx/op/convert.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
@@ -10,10 +10,10 @@ namespace gpu {
 
 struct context;
 
-struct hip_fp_conversion
+struct hip_convert
 {
-    op::fp_conversion op;
-    std::string name() const { return "gpu::fp_conversion"; }
+    op::convert op;
+    std::string name() const { return "gpu::convert"; }
     shape compute_shape(std::vector<shape> inputs) const;
     argument
     compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;

src/targets/gpu/include/migraphx/gpu/device/fp_conversion.hpp → src/targets/gpu/include/migraphx/gpu/device/convert.hpp

 
-#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_FP_CONVERSION_HPP
-#define MIGRAPHX_GUARD_RTGLIB_DEVICE_FP_CONVERSION_HPP
+#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_CONVERT_HPP
+#define MIGRAPHX_GUARD_RTGLIB_DEVICE_CONVERT_HPP
 
 #include <migraphx/argument.hpp>
 #include <migraphx/config.hpp>
@@ -12,7 +12,7 @@ namespace gpu {
 namespace device {
 
 argument
-fp_conversion(hipStream_t stream, const shape& output_shape, const std::vector<argument>& args);
+convert(hipStream_t stream, const argument& result, const argument& arg);
 
 } // namespace device
 } // namespace gpu

src/targets/gpu/lowering.cpp

@@ -45,7 +45,7 @@
 #include <migraphx/gpu/pad.hpp>
 #include <migraphx/gpu/gather.hpp>
 #include <migraphx/gpu/lrn.hpp>
-#include <migraphx/gpu/fp_conversion.hpp>
+#include <migraphx/gpu/convert.hpp>
 #include <utility>
 #include <functional>
 #include <algorithm>
@@ -102,7 +102,7 @@ struct miopen_apply
         add_extend_op<hip_logsoftmax, op::logsoftmax>("logsoftmax");
         add_extend_op<hip_gather, op::gather>("gather");
         add_extend_op<hip_pad, op::pad>("pad");
-        add_extend_op<hip_fp_conversion, op::fp_conversion>("fp_conversion");
+        add_extend_op<hip_convert, op::convert>("convert");
 
         add_lrn_op();
         add_convolution_op();