diff --git a/include/onnxruntime/core/framework/op_kernel.h b/include/onnxruntime/core/framework/op_kernel.h
index aff045e214..31a93d1518 100644
--- a/include/onnxruntime/core/framework/op_kernel.h
+++ b/include/onnxruntime/core/framework/op_kernel.h
@@ -359,4 +359,12 @@ using BuildKernelCreateInfoFn = KernelCreateInfo (*)();
         static_cast<KernelCreatePtrFn>([](const OpKernelInfo& info) -> OpKernel* { return new __VA_ARGS__(info); }));        \
   }
 
+// Use within macro definitions to create a custom vector of constraints.
+// Example: #define REG_KERNEL(OP, VERSION, KERNEL_CLASS, Type, ...)
+//  .TypeConstraint("T", BuildKernelDefConstraints<Type, __VA_ARGS_>())
+template <typename T, typename... Types>
+inline std::vector<MLDataType> BuildKernelDefConstraints() {
+  return {DataTypeImpl::GetTensorType<T>(), DataTypeImpl::GetTensorType<Types>()...};
+}
+
 }  // namespace onnxruntime
diff --git a/onnxruntime/core/optimizer/conv_activation_fusion.cc b/onnxruntime/core/optimizer/conv_activation_fusion.cc
index 831b44576d..cd0457f32a 100644
--- a/onnxruntime/core/optimizer/conv_activation_fusion.cc
+++ b/onnxruntime/core/optimizer/conv_activation_fusion.cc
@@ -108,7 +108,7 @@ Status ConvActivationFusion::ApplyImpl(Graph& graph, bool& modified, int graph_l
         !graph_utils::IsSupportedOptypeVersionAndDomain(next_node, "Tanh", {6})) {
       if (graph_utils::IsSupportedOptypeVersionAndDomain(next_node, "LeakyRelu", {6})) {
         activation_params.push_back(graph_utils::GetNodeAttribute(next_node, "alpha")->f());
-      } else if (graph_utils::IsSupportedOptypeVersionAndDomain(next_node, "Clip", {6, 11})) {
+      } else if (graph_utils::IsSupportedOptypeVersionAndDomain(next_node, "Clip", {6, 11, 12})) {
         float min, max;
         if (GetClipConstantMinMax(graph, next_node, min, max)) {
           activation_params.push_back(min);
diff --git a/onnxruntime/core/optimizer/nchwc_transformer.cc b/onnxruntime/core/optimizer/nchwc_transformer.cc
index 148c48da1b..6da13740c1 100644
--- a/onnxruntime/core/optimizer/nchwc_transformer.cc
+++ b/onnxruntime/core/optimizer/nchwc_transformer.cc
@@ -925,7 +925,7 @@ void NchwcTransformerImpl::Transform(Node& node) {
   if (graph_utils::IsSupportedOptypeVersionAndDomain(node, "Conv", {1, 11}) ||
       graph_utils::IsSupportedOptypeVersionAndDomain(node, "FusedConv", {1}, kMSDomain)) {
     TransformConv(node);
-  } else if (graph_utils::IsSupportedOptypeVersionAndDomain(node, "MaxPool", {1, 8, 10, 11}) ||
+  } else if (graph_utils::IsSupportedOptypeVersionAndDomain(node, "MaxPool", {1, 8, 10, 11, 12}) ||
              graph_utils::IsSupportedOptypeVersionAndDomain(node, "AveragePool", {1, 7, 10, 11}) ||
              graph_utils::IsSupportedOptypeVersionAndDomain(node, "GlobalMaxPool", {1}) ||
              graph_utils::IsSupportedOptypeVersionAndDomain(node, "GlobalAveragePool", {1})) {
diff --git a/onnxruntime/core/optimizer/relu_clip_fusion.cc b/onnxruntime/core/optimizer/relu_clip_fusion.cc
index b2143610d5..e0b86f9a7f 100644
--- a/onnxruntime/core/optimizer/relu_clip_fusion.cc
+++ b/onnxruntime/core/optimizer/relu_clip_fusion.cc
@@ -122,7 +122,7 @@ bool FuseReluClip::SatisfyCondition(const Graph& graph, const Node& node, const
   // as Clip will apply the minimum. If the Clip 'min' value is < 0 we need
   // to update it to 0 to apply what the Relu would have done. We do that in Apply.
   const auto& next_node = *node.OutputNodesBegin();
-  if (!graph_utils::IsSupportedOptypeVersionAndDomain(next_node, "Clip", {6, 11}) ||
+  if (!graph_utils::IsSupportedOptypeVersionAndDomain(next_node, "Clip", {6, 11, 12}) ||
       next_node.GetExecutionProviderType() != node.GetExecutionProviderType()) {
     return false;
   }
diff --git a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc
index c44afe8d7f..47c87eb2ec 100644
--- a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc
+++ b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc
@@ -78,10 +78,9 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain,
 class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, 7, float, Sum);
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, float, Sum);
 class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, 7, float, Min);
-class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, float, Min);
+class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, 11, Min);
 class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, 7, float, Max);
-class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, float, Max);
-class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, double, Max);
+class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, 11, Max);
 class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 1, Not);
 class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 7, And);
 class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 7, Or);
@@ -322,7 +321,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain,
 class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 10, ReverseSequence);
 
 // opset 11
-class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, Clip);
+class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 11, Clip);
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, float, CumSum);
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, double, CumSum);
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, int32_t, CumSum);
@@ -441,6 +440,10 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain,
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float, ArgMin);
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t, ArgMin);
 
+class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, Clip);
+
+class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, Min);
+class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, Max);
 class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, MaxPool);
 
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float, ReduceMax);
@@ -543,11 +546,10 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, float, Sum)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, 7,
                                                                             float, Min)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, float, Min)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, 11, Min)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, 7,
                                                                             float, Max)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, float, Max)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, double, Max)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, 11, Max)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 1, Not)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 7, And)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 7, Or)>,
@@ -1004,7 +1006,8 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 10, ReverseSequence)>,
 
       //opset 11
-      BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, Clip)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 11, Clip)>,
+
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, float,
                                                                   CumSum)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, double,
@@ -1142,6 +1145,12 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t,
                                                                   ArgMin)>,
 
+      BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12,Clip)>,
+
+      BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, Min)>,
+
+      BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, Max)>,
+
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, MaxPool)>,
 
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float,
diff --git a/onnxruntime/core/providers/cpu/math/clip.cc b/onnxruntime/core/providers/cpu/math/clip.cc
index dc99582ddc..75ae6d4cdd 100644
--- a/onnxruntime/core/providers/cpu/math/clip.cc
+++ b/onnxruntime/core/providers/cpu/math/clip.cc
@@ -2,6 +2,7 @@
 // Licensed under the MIT License.
 
 #include "core/providers/cpu/math/clip.h"
+#include "core/framework/data_types_internal.h"
 
 namespace onnxruntime {
 
@@ -12,10 +13,57 @@ ONNX_CPU_OPERATOR_VERSIONED_KERNEL(
     KernelDefBuilder().MayInplace(0, 0).TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),
     Clip_6<float>);
 
-ONNX_CPU_OPERATOR_KERNEL(
+ONNX_CPU_OPERATOR_VERSIONED_KERNEL(
     Clip,
     11,
+    11,
     KernelDefBuilder().MayInplace(0, 0).TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),
-    Clip<float>);
+    Clip);
+
+#define REG_KERNEL_NONTEMPL(OP_TYPE, VERSION, KERNEL_CLASS, ...)          \
+  ONNX_CPU_OPERATOR_KERNEL(                                                     \
+      OP_TYPE,                                                                  \
+      VERSION,                                                                  \
+      KernelDefBuilder()                                                        \
+          .MayInplace(0, 0)                                                     \
+          .TypeConstraint("T", BuildKernelDefConstraints<__VA_ARGS__>()), \
+      KERNEL_CLASS);
+
+REG_KERNEL_NONTEMPL(Clip, 12, Clip, float, double, int8_t, uint8_t, int64_t, uint64_t);
+
+template <typename T>
+struct Clip::ComputeImpl {
+  void operator()(const Tensor* X, const Tensor* min, const Tensor* max, Tensor* Y) const {
+    auto min_val = std::numeric_limits<T>::lowest();
+    auto max_val = std::numeric_limits<T>::max();
+    if (min) {
+      ORT_ENFORCE(min->Shape().NumDimensions() == 0, "min should be a scalar.");
+      min_val = *(min->template Data<T>());
+    }
+    if (max) {
+      ORT_ENFORCE(max->Shape().NumDimensions() == 0, "max should be a scalar.");
+      max_val = *(max->template Data<T>());
+    }
+
+    EigenVectorMap<T>(Y->template MutableData<T>(), Y->Shape().Size()) =
+        ConstEigenVectorMap<T>(X->template Data<T>(), X->Shape().Size())
+            .cwiseMax(min_val)
+            .cwiseMin(max_val);
+  }
+};
+
+Status Clip::Compute(OpKernelContext* ctx) const {
+  const auto* X = ctx->Input<Tensor>(0);
+  const auto* min = ctx->Input<Tensor>(1);
+  const auto* max = ctx->Input<Tensor>(2);
+  Tensor* Y = ctx->Output(0, X->Shape());
+
+  utils::MLTypeCallDispatcher<ComputeImpl, float, double, int8_t, uint8_t, int64_t, uint64_t>
+      t_disp(X->GetElementType());
+
+  t_disp.Invoke(X, min, max, Y);
+
+  return Status::OK();
+}
 
 }  // namespace onnxruntime
diff --git a/onnxruntime/core/providers/cpu/math/clip.h b/onnxruntime/core/providers/cpu/math/clip.h
index b4ef64398d..b365aa33fd 100644
--- a/onnxruntime/core/providers/cpu/math/clip.h
+++ b/onnxruntime/core/providers/cpu/math/clip.h
@@ -9,59 +9,53 @@
 
 namespace onnxruntime {
 
-template <typename T>
-class Clip_6 final : public OpKernel {
+namespace clip_internal {
+
+template<typename T>
+class Clip_6Base {
  public:
-  Clip_6(const OpKernelInfo& info) : OpKernel(info) {
-    ORT_ENFORCE(info.GetAttr<T>("max", &max_).IsOK());
-    ORT_ENFORCE(info.GetAttr<T>("min", &min_).IsOK());
+  explicit Clip_6Base(const OpKernelInfo& info) {
+    auto min_val = std::numeric_limits<T>::lowest();
+    auto max_val = std::numeric_limits<T>::max();
+    info.GetAttrOrDefault("min", &min_, min_val);
+    info.GetAttrOrDefault("max", &max_, max_val);
+    ORT_ENFORCE(min_ <= max_);
   }
-
-  Status Compute(OpKernelContext* ctx) const override {
-    const auto* X = ctx->Input<Tensor>(0);
-    Tensor* Y = ctx->Output(0, X->Shape());
-    EigenVectorMap<T>(Y->template MutableData<T>(), Y->Shape().Size()) =
-        ConstEigenVectorMap<T>(X->template Data<T>(), X->Shape().Size())
-            .cwiseMax(min_)
-            .cwiseMin(max_);
-    return Status::OK();
-  }
-
- private:
+ protected:
   T max_;
   T min_;
 };
+} // namespace clip_internal
 
 template <typename T>
-class Clip final : public OpKernel {
+class Clip_6 final : public clip_internal::Clip_6Base<T>, public OpKernel {
  public:
-  Clip(const OpKernelInfo& info) : OpKernel(info) {
+  explicit Clip_6(const OpKernelInfo& info) : clip_internal::Clip_6Base<T>(info), OpKernel(info) {
   }
 
   Status Compute(OpKernelContext* ctx) const override {
     const auto* X = ctx->Input<Tensor>(0);
-    const auto* min = ctx->Input<Tensor>(1);
-    const auto* max = ctx->Input<Tensor>(2);
     Tensor* Y = ctx->Output(0, X->Shape());
-
-    auto min_val = -std::numeric_limits<T>::infinity();
-    auto max_val = std::numeric_limits<T>::infinity();
-    if (min) {
-      ORT_ENFORCE(min->Shape().NumDimensions() == 0, "min should be a scalar.");
-      min_val = *(min->template Data<T>());
-    }
-    if (max) {
-      ORT_ENFORCE(max->Shape().NumDimensions() == 0, "max should be a scalar.");
-      max_val = *(max->template Data<T>());
-    }
-
     EigenVectorMap<T>(Y->template MutableData<T>(), Y->Shape().Size()) =
         ConstEigenVectorMap<T>(X->template Data<T>(), X->Shape().Size())
-            .cwiseMax(min_val)
-            .cwiseMin(max_val);
-
+            .cwiseMax(this->min_)
+            .cwiseMin(this->max_);
     return Status::OK();
   }
 };
 
+// Since version 11. Min and Max are inputs
+// version 12 adds type support
+class Clip final : public OpKernel {
+ public:
+  explicit Clip(const OpKernelInfo& info) : OpKernel(info) {
+  }
+
+  Status Compute(OpKernelContext* ctx) const override;
+
+ private:
+  template<typename T>
+  struct ComputeImpl;
+};
+
 }  // namespace onnxruntime
diff --git a/onnxruntime/core/providers/cpu/math/element_wise_ops.cc b/onnxruntime/core/providers/cpu/math/element_wise_ops.cc
index 97316c48f5..aaaf6e70b4 100644
--- a/onnxruntime/core/providers/cpu/math/element_wise_ops.cc
+++ b/onnxruntime/core/providers/cpu/math/element_wise_ops.cc
@@ -1,6 +1,7 @@
 // Copyright (c) Microsoft Corporation. All rights reserved.
 // Licensed under the MIT License.
 
+#include "core/framework/data_types_internal.h"
 #include "core/providers/cpu/math/element_wise_ops.h"
 #include <unsupported/Eigen/SpecialFunctions>
 #include "core/util/math.h"
@@ -46,6 +47,27 @@ namespace onnxruntime {
           .TypeConstraint("T1", DataTypeImpl::GetTensorType<bool>()),                                           \
       KERNEL_CLASS<TYPE>);
 
+// var args are type constraints
+#define REG_ELEMENTWISE_KERNEL_NONT(OP_TYPE, VERSION, KERNEL_CLASS, ...)   \
+  ONNX_CPU_OPERATOR_KERNEL(                                                      \
+      OP_TYPE,                                                                   \
+      VERSION,                                                                   \
+      KernelDefBuilder()                                                         \
+          .TypeConstraint("T", BuildKernelDefConstraints<__VA_ARGS__>())   \
+          .TypeConstraint("T1", BuildKernelDefConstraints<__VA_ARGS__>()), \
+      KERNEL_CLASS);
+
+// var args are type constraints
+#define REG_ELEMENTWISE_VERSIONED_KERNEL_NONT(OP_TYPE, VERSION_FROM, VERSION_TO, KERNEL_CLASS, ...) \
+  ONNX_CPU_OPERATOR_VERSIONED_KERNEL(                                                                     \
+      OP_TYPE,                                                                                            \
+      VERSION_FROM,                                                                                       \
+      VERSION_TO,                                                                                         \
+      KernelDefBuilder()                                                                                  \
+          .TypeConstraint("T", BuildKernelDefConstraints<__VA_ARGS__>())                            \
+          .TypeConstraint("T1", BuildKernelDefConstraints<__VA_ARGS__>()),                          \
+      KERNEL_CLASS);
+
 REG_ELEMENTWISE_TYPED_KERNEL(Add, 7, float, Add);
 REG_ELEMENTWISE_TYPED_KERNEL(Add, 7, double, Add);
 REG_ELEMENTWISE_TYPED_KERNEL(Add, 7, int32_t, Add);
@@ -103,12 +125,13 @@ REG_ELEMENTWISE_TYPED_KERNEL(Log, 6, float, Log);
 REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(Sum, 6, 7, float, Sum_6);
 REG_ELEMENTWISE_TYPED_KERNEL(Sum, 8, float, Sum_8);
 
-REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(Min, 6, 7, float, Min_6);
-REG_ELEMENTWISE_TYPED_KERNEL(Min, 8, float, Min_8);
-
 REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(Max, 6, 7, float, Max_6);
-REG_ELEMENTWISE_TYPED_KERNEL(Max, 8, float, Max_8);
-REG_ELEMENTWISE_TYPED_KERNEL(Max, 8, double, Max_8);
+REG_ELEMENTWISE_VERSIONED_KERNEL_NONT(Max, 8, 11, Max_8, float, double);
+REG_ELEMENTWISE_KERNEL_NONT(Max, 12, Max_8, float, double, MLFloat16, int32_t, uint32_t, int64_t, uint64_t);
+
+REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(Min, 6, 7, float, Min_6);
+REG_ELEMENTWISE_VERSIONED_KERNEL_NONT(Min, 8, 11, Min_8, float);
+REG_ELEMENTWISE_KERNEL_NONT(Min, 12, Min_8, float, double, MLFloat16, int32_t, uint32_t, int64_t, uint64_t);
 
 REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Less, 7, 9, float, Less);
 REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Less, 7, 9, double, Less);
@@ -343,13 +366,21 @@ Status Min_6<float>::Compute(OpKernelContext* ctx) const {
   return Status::OK();
 }
 
-template <>
-Status Min_8<float>::Compute(OpKernelContext* context) const {
-  return BroadcastVariadic<float, float>(
-      Node(), *context,
-      [](EigenVectorMap<float> output, float input0, ConstEigenVectorMap<float> input1) { output = input1.array().min(input0); },
-      [](EigenVectorMap<float> output, ConstEigenVectorMap<float> input0, float input1) { output = input0.array().min(input1); },
-      [](EigenVectorMap<float> output, ConstEigenVectorMap<float> input0, ConstEigenVectorMap<float> input1) { output = input0.array().min(input1.array()); });
+template <typename T>
+struct Min_8::ComputeImpl {
+  Status operator()(const Min_8* inst, OpKernelContext* context) const {
+    return BroadcastVariadic<T, T>(
+        inst->Node(), *context,
+        [](EigenVectorMap<T> output, T input0, ConstEigenVectorMap<T> input1) { output = input1.array().min(input0); },
+        [](EigenVectorMap<T> output, ConstEigenVectorMap<T> input0, T input1) { output = input0.array().min(input1); },
+        [](EigenVectorMap<T> output, ConstEigenVectorMap<T> input0, ConstEigenVectorMap<T> input1) { output = input0.array().min(input1.array()); });
+  }
+};
+
+Status Min_8::Compute(OpKernelContext* context) const {
+  utils::MLTypeCallDispatcherRet<Status, ComputeImpl, float, double, MLFloat16, int32_t, uint32_t, int64_t, uint64_t>
+      t_disp(context->Input<Tensor>(0)->GetElementType());
+  return t_disp.Invoke(this, context);
 }
 
 template <>
@@ -371,12 +402,20 @@ Status Max_6<float>::Compute(OpKernelContext* ctx) const {
 }
 
 template <typename T>
-Status Max_8<T>::Compute(OpKernelContext* context) const {
-  return BroadcastVariadic<T, T>(
-      Node(), *context,
-      [](EigenVectorMap<T> output, T input0, ConstEigenVectorMap<T> input1) { output = input1.array().max(input0); },
-      [](EigenVectorMap<T> output, ConstEigenVectorMap<T> input0, T input1) { output = input0.array().max(input1); },
-      [](EigenVectorMap<T> output, ConstEigenVectorMap<T> input0, ConstEigenVectorMap<T> input1) { output = input0.array().max(input1.array()); });
+struct Max_8::ComputeImpl {
+  Status operator()(const Max_8* inst, OpKernelContext* context) const {
+    return BroadcastVariadic<T, T>(
+        inst->Node(), *context,
+        [](EigenVectorMap<T> output, T input0, ConstEigenVectorMap<T> input1) { output = input1.array().max(input0); },
+        [](EigenVectorMap<T> output, ConstEigenVectorMap<T> input0, T input1) { output = input0.array().max(input1); },
+        [](EigenVectorMap<T> output, ConstEigenVectorMap<T> input0, ConstEigenVectorMap<T> input1) { output = input0.array().max(input1.array()); });
+  }
+};
+
+Status Max_8::Compute(OpKernelContext* context) const {
+  utils::MLTypeCallDispatcherRet<Status, ComputeImpl, float, double, MLFloat16, int32_t, uint32_t, int64_t, uint64_t>
+      t_disp(context->Input<Tensor>(0)->GetElementType());
+  return t_disp.Invoke(this, context);
 }
 
 Status Not::Compute(OpKernelContext* context) const {
diff --git a/onnxruntime/core/providers/cpu/math/element_wise_ops.h b/onnxruntime/core/providers/cpu/math/element_wise_ops.h
index e84123d1ed..79e53f5b9e 100644
--- a/onnxruntime/core/providers/cpu/math/element_wise_ops.h
+++ b/onnxruntime/core/providers/cpu/math/element_wise_ops.h
@@ -165,13 +165,19 @@ class Min_6 final : public OpKernel {
   Status Compute(OpKernelContext* context) const override;
 };
 
-template <typename T>
+// Max versions 8 - 12
+// Version 8 added broadcast
+// Version 12 added types support
 class Min_8 final : public OpKernel {
  public:
   Min_8(const OpKernelInfo& info) : OpKernel(info) {
   }
 
   Status Compute(OpKernelContext* context) const override;
+
+ private:
+  template <typename T>
+  struct ComputeImpl;
 };
 
 template <typename T>
@@ -183,13 +189,20 @@ class Max_6 final : public OpKernel {
   Status Compute(OpKernelContext* context) const override;
 };
 
-template <typename T>
+// Max versions 8 - 12
+// Version 8 added broadcast
+// Version 12 added types support
 class Max_8 final : public OpKernel {
  public:
   Max_8(const OpKernelInfo& info) : OpKernel(info) {
   }
 
   Status Compute(OpKernelContext* context) const override;
+
+private:
+
+ template<typename T>
+ struct ComputeImpl;
 };
 
 class Not final : public OpKernel {
diff --git a/onnxruntime/core/providers/cuda/cu_inc/common.cuh b/onnxruntime/core/providers/cuda/cu_inc/common.cuh
index 8fee721a55..901b2c52dc 100644
--- a/onnxruntime/core/providers/cuda/cu_inc/common.cuh
+++ b/onnxruntime/core/providers/cuda/cu_inc/common.cuh
@@ -129,7 +129,7 @@ template <>
 __device__ __inline__ double _Round(double a) { return rint(a); }
 
 template <>
-__device__ __inline__ half _Round(half a) { 
+__device__ __inline__ half _Round(half a) {
 #if __CUDA_ARCH__ < 530
   return half(rintf((float)a));
 #else
@@ -196,9 +196,33 @@ __device__ __inline__ half _Pow(half a, half b) { return half(powf((float)a, (fl
 template <typename T>
 __device__ __inline__ T _Min(T a, T b) { return a < b ? a : b; }
 
+template <>
+__device__ __inline__ int _Min(int a, int b) { return a < b ? a : b; }
+
+template <>
+__device__ __inline__ unsigned int _Min(unsigned int a, unsigned int b) { return a < b ? a : b; }
+
+template <>
+__device__ __inline__ long long _Min(long long a, long long b) { return a < b ? a : b; }
+
+template <>
+__device__ __inline__ unsigned long long _Min(unsigned long long a, unsigned long long b) { return a < b ? a : b; }
+
 template <typename T>
 __device__ __inline__ T _Max(T a, T b) { return a > b ? a : b; }
 
+template <>
+__device__ __inline__ int _Max(int a, int b) { return a > b ? a : b; }
+
+template <>
+__device__ __inline__ unsigned int _Max(unsigned int a, unsigned int b) { return a > b ? a : b; }
+
+template <>
+__device__ __inline__ long long _Max(long long a, long long b) { return a > b ? a : b; }
+
+template <>
+__device__ __inline__ unsigned long long _Max(unsigned long long a, unsigned long long b) { return a > b ? a : b; }
+
 template <typename T>
 __device__ __inline__ T _Abs(T a) { return a > (T)0 ? a : -a; }
 
@@ -226,12 +250,6 @@ __device__ __inline__ T _Gelu(T a) {
 #define CUDA_LONG int32_t
 #endif
 
-#define IDX2C(i, j, ld) (((j) * (ld)) + (i))  // 0 based indexing
-
-// ---------------------------------------------------------------------------
-// GridDim -- helper to choose the CUDA grid dimensions
-// ---------------------------------------------------------------------------
-
 template <class INT, class INT2>
 static INT CeilDiv(INT a, INT2 b)  // ceil(a/b)
 {
@@ -241,54 +259,13 @@ static INT CeilDiv(INT a, INT2 b)  // ceil(a/b)
 struct GridDim {
   enum : CUDA_LONG {
     maxThreadsPerBlock = 256,  // max threads per block
-    maxWarpsPerBlock = 32,     // max warps per block
     maxElementsPerThread = 4,  // max element processed per thread
   };
-
-  // use these for launching
-  //   GridDim grid(NN);
-  //   kernel<<<grid.m_blocksPerGrid, grid.m_threadsPerBlock, ...>>>(...)
-  int blocks_per_grid_, threads_per_block_;  // (these may in the future be extended to multi-dimensional ones)
-  CUDA_LONG N_;
-
-  GridDim(CUDA_LONG N)  // linear grid
-  {
-    N_ = N;
-    if (N == 0)  // CUDA will fail to launch with 0 blocks
-      N = 1;
-
-    // get device information
-    const auto& props = DeviceProp::GetDeviceProps();
-    CUDA_LONG numProcs = props.multiProcessorCount;
-    CUDA_LONG warpSize = props.warpSize;
-
-    // distribute warps evenly over processors
-    CUDA_LONG warpsPerProc = CeilDiv(N, numProcs * warpSize);
-
-    // if too many warps per block then reduce #warps
-    // This limits the number of threads to 512.
-    if (warpsPerProc > maxWarpsPerBlock) {
-      CUDA_LONG overBy = CeilDiv(warpsPerProc, maxWarpsPerBlock);  // we are over by this factor
-      warpsPerProc = CeilDiv(warpsPerProc, overBy);
-    }
-
-    // put it back together
-    threads_per_block_ = warpsPerProc * warpSize;  // =a multiple of 32 that is as close to 1024 as makes sense given NN
-    blocks_per_grid_ = CeilDiv(N, threads_per_block_);
-    if (blocks_per_grid_ == 1)
-      threads_per_block_ = N;  // don't launch more than necessary
-    assert(blocks_per_grid_ * threads_per_block_ >= N);
-  }
-
-  // compute our location on the grid
-  static __device__ CUDA_LONG GetLinearThreadId() {
-    return blockDim.x * blockIdx.x + threadIdx.x;
-  }
 };
 
-#define CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N) \
-  CUDA_LONG id = GridDim::GetLinearThreadId();     \
-  if (id >= N)                                     \
+#define CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N)      \
+  CUDA_LONG id = blockDim.x * blockIdx.x + threadIdx.x; \
+  if (id >= N)                                          \
     return;
 
 // CUDA_KERNEL_ASSERT is a macro that wraps an assert() call inside cuda kernels.
@@ -296,9 +273,9 @@ struct GridDim {
 // See http://docs.nvidia.com/cuda/cuda-c-programming-guide/#assertion
 #if defined(__APPLE__) || defined(__HIP_PLATFORM_HCC__)
 #define CUDA_KERNEL_ASSERT(...)
-#else // __APPLE__
+#else  // __APPLE__
 #define CUDA_KERNEL_ASSERT(...) assert(__VA_ARGS__)
-#endif // __APPLE__
+#endif  // __APPLE__
 
 }  // namespace cuda
 }  // namespace onnxruntime
diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc
index 55a7650580..514d6d50d6 100644
--- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc
+++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc
@@ -318,15 +318,29 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain,
 class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, float, Max);
 class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, double, Max);
 class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, MLFloat16, Max);
-class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, float, Max);
-class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, double, Max);
-class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, MLFloat16, Max);
+class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, float, Max);
+class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, double, Max);
+class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, MLFloat16, Max);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, Max);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int64_t, Max);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint32_t, Max);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint64_t, Max);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, Max);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, Max);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, Max);
 class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, float, Min);
 class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, double, Min);
 class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, MLFloat16, Min);
-class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, float, Min);
-class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, double, Min);
-class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, MLFloat16, Min);
+class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, float, Min);
+class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, double, Min);
+class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, MLFloat16, Min);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, Min);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int64_t, Min);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint32_t, Min);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint64_t, Min);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, Min);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, Min);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, Min);
 class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, float, Greater);
 class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, double, Greater);
 class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, MLFloat16, Greater);
@@ -690,7 +704,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain,
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, Resize);
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, Resize);
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, uint8_t, Resize);
-class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Clip);
+class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, Clip);
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Pad);
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, Pad);
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, Pad);
@@ -707,6 +721,8 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain,
 class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, CumSum);
 
 // OpSet 12
+class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, Clip);
+
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, MaxPool);
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, MaxPool);
 class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, MaxPool);
@@ -812,15 +828,29 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) {
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, float, Max)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, double, Max)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, MLFloat16, Max)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, float, Max)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, double, Max)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, MLFloat16, Max)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, float, Max)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, double, Max)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, MLFloat16, Max)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, Max)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int64_t, Max)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint32_t, Max)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint64_t, Max)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, Max)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, Max)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, Max)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, float, Min)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, double, Min)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, MLFloat16, Min)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, float, Min)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, double, Min)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, MLFloat16, Min)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, float, Min)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, double, Min)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, MLFloat16, Min)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, Min)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int64_t, Min)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint32_t, Min)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint64_t, Min)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, Min)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, Min)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, Min)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, float, Greater)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, double, Greater)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, MLFloat16, Greater)>,
@@ -1187,7 +1217,7 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) {
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, Resize)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, Resize)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, uint8_t, Resize)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Clip)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, Clip)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Pad)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, Pad)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, Pad)>,
@@ -1200,6 +1230,8 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) {
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, CumSum)>,
 
       // OpSet 12
+      BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, Clip)>,
+
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, MaxPool)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, MaxPool)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, MaxPool)>,
@@ -1368,12 +1400,12 @@ CUDAExecutionProvider::GetCapability(const onnxruntime::GraphViewer& graph,
     for (auto registry : kernel_registries) {
       cuda_kernel_def = registry->TryFindKernel(node, Type());
 
-      // atleast one registry has a CUDA kernel for this node 
+      // atleast one registry has a CUDA kernel for this node
       if (cuda_kernel_def)
-          break;
+        break;
     }
 
-    // none of the provided registries has a CUDA kernel for this node  
+    // none of the provided registries has a CUDA kernel for this node
     if (cuda_kernel_def == nullptr) {
       // node is not in cuda exeuction provider if no kernel def found,
       // or if other execution provider already assigned to it
diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc
index 544558882e..dc57ab1f5c 100644
--- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc
+++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc
@@ -384,9 +384,11 @@ BINARY_LOGICALOP_REGISTER_UZILHFD(Greater, 9)
 BINARY_OP_REGISTER_VERSIONED_OIL(Equal, 7, 10)
 BINARY_OP_REGISTER_OIL(Equal, 11)
 BINARY_OP_REGISTER_VERSIONED_HFD(Greater, 7, 8)
-BINARY_OP_REGISTER_HFD(Max, 8)
+BINARY_OP_REGISTER_UZILHFD(Max, 12)
+BINARY_OP_REGISTER_VERSIONED_HFD(Max, 8, 11)
 BINARY_OP_REGISTER_VERSIONED_HFD(Max, 6, 7)
-BINARY_OP_REGISTER_HFD(Min, 8)
+BINARY_OP_REGISTER_UZILHFD(Min, 12)
+BINARY_OP_REGISTER_VERSIONED_HFD(Min, 8, 11)
 BINARY_OP_REGISTER_VERSIONED_HFD(Min, 6, 7)
 BINARY_LOGICALOP_REGISTER_UZILHFD(Less, 9)
 BINARY_OP_REGISTER_VERSIONED_HFD(Less, 7, 8)
diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu
index 832f41702f..ad4d428313 100644
--- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu
+++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu
@@ -92,8 +92,8 @@ SPECIALIZED_BINARY_ELEMENTWISE_IMPL(Xor, bool)
 SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(PRelu)
 SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Greater)
 SPECIALIZED_BINARY_ELEMENTWISE_IMPL_OIL(Equal)
-SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(Max)
-SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(Min)
+SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Max)
+SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Min)
 SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Less)
 
 }  // namespace cuda
diff --git a/onnxruntime/core/providers/cuda/math/clip.cc b/onnxruntime/core/providers/cuda/math/clip.cc
index 7bfa0f1fe4..55176624f7 100644
--- a/onnxruntime/core/providers/cuda/math/clip.cc
+++ b/onnxruntime/core/providers/cuda/math/clip.cc
@@ -8,61 +8,114 @@
 namespace onnxruntime {
 namespace cuda {
 
-#define REGISTER_KERNEL_TYPED(T)                                  \
-  ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX(                        \
-      Clip,                                                       \
-      kOnnxDomain,                                                \
-      6,                                                          \
-      10,                                                         \
-      T,                                                          \
-      kCudaExecutionProvider,                                     \
-      KernelDefBuilder()                                          \
-          .TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
-      Clip<T>);                                                   \
-  ONNX_OPERATOR_TYPED_KERNEL_EX(                                  \
-      Clip,                                                       \
-      kOnnxDomain,                                                \
-      11,                                                         \
-      T,                                                          \
-      kCudaExecutionProvider,                                     \
-      KernelDefBuilder()                                          \
-          .InputMemoryType<OrtMemTypeCPUInput>(1)                 \
-          .InputMemoryType<OrtMemTypeCPUInput>(2)                 \
-          .TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
-      Clip<T>);
+ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX(
+    Clip,
+    kOnnxDomain,
+    6,
+    10,
+    float,
+    kCudaExecutionProvider,
+    KernelDefBuilder()
+        .TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),
+    Clip_6<float>);
+
+ONNX_OPERATOR_VERSIONED_KERNEL_EX(
+    Clip,
+    kOnnxDomain,
+    11, 11,
+    kCudaExecutionProvider,
+    KernelDefBuilder()
+        .InputMemoryType<OrtMemTypeCPUInput>(1)
+        .InputMemoryType<OrtMemTypeCPUInput>(2)
+        .TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),
+    Clip);
+
+ONNX_OPERATOR_KERNEL_EX(
+    Clip,
+    kOnnxDomain,
+    12,
+    kCudaExecutionProvider,
+    KernelDefBuilder()
+        .InputMemoryType<OrtMemTypeCPUInput>(1)
+        .InputMemoryType<OrtMemTypeCPUInput>(2)
+        .TypeConstraint("T", BuildKernelDefConstraints<float, double, int8_t, uint8_t, int64_t, uint64_t>()),
+    Clip);
 
 template <typename T>
-Status Clip<T>::ComputeInternal(OpKernelContext* ctx) const {
-  T min_val = min_;
-  T max_val = max_;
-  if (is_min_max_input_) {
-    const auto* min_input = ctx->Input<Tensor>(1);
-    const auto* max_input = ctx->Input<Tensor>(2);
-    if (min_input) {
-      ORT_ENFORCE(min_input->Shape().NumDimensions() == 0, "min should be a scalar.");
-      min_val = *(min_input->template Data<T>());
-    }
-    if (max_input) {
-      ORT_ENFORCE(max_input->Shape().NumDimensions() == 0, "max should be a scalar.");
-      max_val = *(max_input->template Data<T>());
-    }
-    ORT_ENFORCE(min_val <= max_val);
-  }
-
+Status Clip_6<T>::ComputeInternal(OpKernelContext* ctx) const {
   const Tensor& X = *ctx->Input<Tensor>(0);
   const TensorShape& input_shape{X.Shape()};
   Tensor* Y = ctx->Output(0, input_shape);
-  size_t count = input_shape.Size();
+  const size_t count = input_shape.Size();
   if (count > 0) {
     auto* y_data = Y->template MutableData<T>();
     const auto* x_data = X.template Data<T>();
-    ClipImpl<T>(x_data, y_data, min_val, max_val, count);
+    ClipImpl<T>(x_data, y_data, this->min_, this->max_, count);
   }
+  return Status::OK();
+}
+
+namespace clip_internal {
+template <typename T>
+struct LowMax {
+  constexpr static T low() {
+    return std::numeric_limits<T>::lowest();
+  }
+  constexpr static T max() {
+    return std::numeric_limits<T>::max();
+  }
+};
+
+template <>
+struct LowMax<MLFloat16> {
+  static MLFloat16 low() {
+    return MLFloat16(math::floatToHalf(std::numeric_limits<float>::lowest()));
+  }
+  static MLFloat16 max() {
+    return MLFloat16(math::floatToHalf(std::numeric_limits<float>::max()));
+  }
+};
+}  // namespace clip_internal
+
+template <typename T>
+struct Clip::ComputeImpl {
+  void operator()(const Tensor* X, const Tensor* min, const Tensor* max, Tensor* Y) const {
+    auto min_val = clip_internal::LowMax<T>::low();
+    auto max_val = clip_internal::LowMax<T>::max();
+
+    // 1-2 Input on CPU
+    if (min) {
+      ORT_ENFORCE(min->Shape().NumDimensions() == 0, "min should be a scalar.");
+      min_val = *(min->template Data<T>());
+    }
+
+    if (max) {
+      ORT_ENFORCE(max->Shape().NumDimensions() == 0, "max should be a scalar.");
+      max_val = *(max->template Data<T>());
+    }
+
+    const size_t count = X->Shape().Size();
+    if (count > 0) {
+      auto* y_data = Y->template MutableData<T>();
+      const auto* x_data = X->template Data<T>();
+      ClipImpl<T>(x_data, y_data, min_val, max_val, count);
+    }
+  }
+};
+
+Status Clip::ComputeInternal(OpKernelContext* ctx) const {
+  const auto* X = ctx->Input<Tensor>(0);
+  const auto* min = ctx->Input<Tensor>(1);
+  const auto* max = ctx->Input<Tensor>(2);
+  Tensor* Y = ctx->Output(0, X->Shape());
+
+  utils::MLTypeCallDispatcher<ComputeImpl, float, double, int8_t, uint8_t, int64_t, uint64_t>
+      t_disp(X->GetElementType());
+
+  t_disp.Invoke(X, min, max, Y);
 
   return Status::OK();
 }
 
-REGISTER_KERNEL_TYPED(float)
-
 }  // namespace cuda
 }  // namespace onnxruntime
diff --git a/onnxruntime/core/providers/cuda/math/clip.h b/onnxruntime/core/providers/cuda/math/clip.h
index d600dc1e9c..63ce6ada26 100644
--- a/onnxruntime/core/providers/cuda/math/clip.h
+++ b/onnxruntime/core/providers/cuda/math/clip.h
@@ -3,36 +3,32 @@
 
 #pragma once
 #include "core/providers/cuda/cuda_common.h"
+#include "core/providers/cpu/math/clip.h"
 
 namespace onnxruntime {
 namespace cuda {
 
 template <typename T>
+class Clip_6 final : public onnxruntime::clip_internal::Clip_6Base<T>,  public CudaKernel {
+ public:
+  explicit Clip_6(const OpKernelInfo& info) : onnxruntime::clip_internal::Clip_6Base<T>(info), CudaKernel{info} {
+  }
+
+  Status ComputeInternal(OpKernelContext* context) const override;
+};
+
+// Since version 11. Min and Max are inputs
+// version 12 adds type support
 class Clip final : public CudaKernel {
  public:
-  Clip(const OpKernelInfo& info) : CudaKernel{info}, is_min_max_input_(false) {
-    int start_version;
-    int end_version;
-    info.GetKernelDef().SinceVersion(&start_version, &end_version);
-
-    if (start_version < 11) {
-      auto min_val = -std::numeric_limits<T>::infinity();
-      auto max_val = std::numeric_limits<T>::infinity();
-      info.GetAttrOrDefault("min", &min_, min_val);
-      info.GetAttrOrDefault("max", &max_, max_val);
-      ORT_ENFORCE(min_ <= max_);
-    } else {
-      min_ = -std::numeric_limits<T>::infinity();
-      max_ = std::numeric_limits<T>::infinity();
-      is_min_max_input_ = true;
-    }
+  explicit Clip(const OpKernelInfo& info) : CudaKernel{info} {
   }
 
   Status ComputeInternal(OpKernelContext* context) const override;
 
  private:
-  T min_, max_;
-  bool is_min_max_input_;
+  template<typename T>
+  struct ComputeImpl;
 };
 
 }  // namespace cuda
diff --git a/onnxruntime/core/providers/cuda/math/clip_impl.cu b/onnxruntime/core/providers/cuda/math/clip_impl.cu
index a73c5ad812..3af9283bd8 100644
--- a/onnxruntime/core/providers/cuda/math/clip_impl.cu
+++ b/onnxruntime/core/providers/cuda/math/clip_impl.cu
@@ -27,6 +27,10 @@ void ClipImpl(const T* input_data, T* output_data, T min, T max, size_t count) {
 template void ClipImpl<float>(const float* input_data, float* output_data, float min, float max, size_t count);
 template void ClipImpl<double>(const double* input_data, double* output_data, double min, double max, size_t count);
 template void ClipImpl<MLFloat16>(const MLFloat16* input_data, MLFloat16* output_data, MLFloat16 min, MLFloat16 max, size_t count);
+template void ClipImpl<int8_t>(const int8_t* input_data, int8_t* output_data, int8_t min, int8_t max, size_t count);
+template void ClipImpl<uint8_t>(const uint8_t* input_data, uint8_t* output_data, uint8_t min, uint8_t max, size_t count);
+template void ClipImpl<int64_t>(const int64_t* input_data, int64_t* output_data, int64_t min, int64_t max, size_t count);
+template void ClipImpl<uint64_t>(const uint64_t* input_data, uint64_t* output_data, uint64_t min, uint64_t max, size_t count);
 
 }  // namespace cuda
 }  // namespace onnxruntime
diff --git a/onnxruntime/test/providers/cpu/math/clip_test.cc b/onnxruntime/test/providers/cpu/math/clip_test.cc
index b678bdf301..b87bc4683c 100644
--- a/onnxruntime/test/providers/cpu/math/clip_test.cc
+++ b/onnxruntime/test/providers/cpu/math/clip_test.cc
@@ -26,7 +26,7 @@ TEST(MathOpTest, Clip_6) {
 }
 
 TEST(MathOpTest, Clip_Default) {
-  OpTester test("Clip", 11);
+  OpTester test("Clip", -1);
 
   std::vector<int64_t> dims{3, 3};
   test.AddInput<float>("X", dims,
@@ -38,10 +38,79 @@ TEST(MathOpTest, Clip_Default) {
                          -1.3f, 3.5f, 64.0f,
                          -5.4f, 9.3f, 82.4f});
 
-  // nGraph does not support Clip opset 11 yet.
+  // nGraph does not support Clip opset 12 yet.
   test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kNGraphExecutionProvider});
 }
 
+TEST(MathOpTest, Clip_Default_int8) {
+  OpTester test("Clip", -1);
+
+  std::vector<int64_t> dims{3, 3};
+  test.AddInput<int8_t>("X", dims,
+                         {11, 4, 127,
+                          -1, 3, 64,
+                          -5, 9, 82});
+  test.AddOutput<int8_t>("Y", dims,
+                          {11, 4, 127,
+                           -1, 3, 64,
+                           -5, 9, 82});
+
+  // TensorRT, nGraph does not support Clip opset 12 yet.
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kNGraphExecutionProvider});
+}
+
+TEST(MathOpTest, Clip_Default_uint8) {
+  OpTester test("Clip", -1);
+
+  std::vector<int64_t> dims{3, 3};
+  test.AddInput<uint8_t>("X", dims,
+                        {11, 4, 255,
+                         1, 3, 64,
+                         5, 9, 82});
+  test.AddOutput<uint8_t>("Y", dims,
+                         {11, 4, 255,
+                          1, 3, 64,
+                          5, 9, 82});
+
+  // TensorRT, nGraph does not support Clip opset 12 yet.
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kNGraphExecutionProvider});
+}
+
+TEST(MathOpTest, Clip_Default_int64) {
+  OpTester test("Clip", -1);
+
+  std::vector<int64_t> dims{3, 3};
+  test.AddInput<int64_t>("X", dims,
+                        {11, 4, 432,
+                         -1, 3, 64,
+                         -5, 9, 82});
+  test.AddOutput<int64_t>("Y", dims,
+                         {11, 4, 432,
+                          -1, 3, 64,
+                          -5, 9, 82});
+
+  // TensorRT, nGraph does not support Clip opset 12 yet.
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kNGraphExecutionProvider});
+}
+
+
+TEST(MathOpTest, Clip_Default_uint64) {
+  OpTester test("Clip", -1);
+
+  std::vector<int64_t> dims{3, 3};
+  test.AddInput<uint64_t>("X", dims,
+                         {11, 4, 432,
+                          1, 3, 64,
+                          5, 9, 82});
+  test.AddOutput<uint64_t>("Y", dims,
+                          {11, 4, 432,
+                           1, 3, 64,
+                           5, 9, 82});
+
+  // TensorRT, nGraph does not support Clip opset 12 yet.
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kNGraphExecutionProvider});
+}
+
 TEST(MathOpTest, Clip) {
   OpTester test("Clip", 11);
 
@@ -57,14 +126,14 @@ TEST(MathOpTest, Clip) {
                          -5.0f, 0.0f, 5.0f,
                          -5.0f, 2.0f, 5.0f});
 
-  // nGraph and Tensorrt does not support Clip opset 11 yet.
+  // TensorRT, nGraph and Tensorrt does not support Clip opset 11 yet.
   test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kNGraphExecutionProvider, kTensorrtExecutionProvider});
 }
 
 TEST(MathOpTest, ClipDimWithZero) {
   std::vector<int64_t> dims{3, 0};  // dim with value of zero should be handled
 
-  OpTester test("Clip", 11);  // latest opset
+  OpTester test("Clip", -1);  // latest opset
   test.AddInput<float>("X", dims, {});
   test.AddInput<float>("min", {}, {-5});
   test.AddInput<float>("max", {}, {5});
diff --git a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc
index e842b63a50..d25e30c38d 100644
--- a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc
+++ b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc
@@ -700,7 +700,7 @@ TEST(MathOpTest, Sum_8_Test1) {
   //This test runs fine on CPU and GPU Plugins
   test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kOpenVINOExecutionProvider});
 #else
-  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});                    //TensorRT: Expected output shape [{3,3,3}] did not match run output shape [{3,1,1}] for sum
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});  //TensorRT: Expected output shape [{3,3,3}] did not match run output shape [{3,1,1}] for sum
 #endif
 }
 
@@ -783,6 +783,108 @@ TEST(MathOpTest, Min_8) {
   test.Run();
 }
 
+TEST(MathOpTest, Min_12_Float) {
+  OpTester test("Min", 12);
+  test.AddInput<float>("data_0", {1, 3},
+                       {1.0f, 2.0f, 3.0f});
+  test.AddInput<float>("data_2", {3, 3},
+                       {10.0f, 20.0f, 30.0f,
+                        40.0f, 50.0f, 60.0f,
+                        -70.0f, -80.0f, -90.0f});
+  test.AddInput<float>("data_1", {3, 1},
+                       {-1.0f, 20.0f, 300.0f});
+  test.AddOutput<float>("min", {3, 3},
+                        {-1.0f, -1.0f, -1.0f,
+                         1.0f, 2.0f, 3.0f,
+                         -70.0f, -80.0f, -90.0f});
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});  //TensorRT: Input batch size is inconsistent
+}
+
+TEST(MathOpTest, Min_12_Double) {
+  OpTester test("Min", 12);
+  test.AddInput<double>("data_0", {1, 3},
+                        {1.0f, 2.0f, 3.0f});
+  test.AddInput<double>("data_2", {3, 3},
+                        {10.0f, 20.0f, 30.0f,
+                         40.0f, 50.0f, 60.0f,
+                         -70.0f, -80.0f, -90.0f});
+  test.AddInput<double>("data_1", {3, 1},
+                        {-1.0f, 20.0f, 300.0f});
+  test.AddOutput<double>("min", {3, 3},
+                         {-1.0f, -1.0f, -1.0f,
+                          1.0f, 2.0f, 3.0f,
+                          -70.0f, -80.0f, -90.0f});
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});  //TensorRT: Input batch size is inconsistent
+}
+
+TEST(MathOpTest, Min_12_Int32) {
+  OpTester test("Min", 12);
+  test.AddInput<int32_t>("data_0", {1, 3},
+                         {1, 2, 3});
+  test.AddInput<int32_t>("data_2", {3, 3},
+                         {10, 20, 30,
+                          40, 50, 60,
+                          -70, -80, -90});
+  test.AddInput<int32_t>("data_1", {3, 1},
+                         {-1, 20, 300});
+  test.AddOutput<int32_t>("min", {3, 3},
+                          {-1, -1, -1,
+                           1, 2, 3,
+                           -70, -80, -90});
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});  //TensorRT: Input batch size is inconsistent
+}
+
+TEST(MathOpTest, Min_12_Int64) {
+  OpTester test("Min", 12);
+  test.AddInput<int64_t>("data_0", {1, 3},
+                         {1, 2, 3});
+  test.AddInput<int64_t>("data_2", {3, 3},
+                         {10, 20, 30,
+                          40, 50, 60,
+                          -70, -80, -90});
+  test.AddInput<int64_t>("data_1", {3, 1},
+                         {-1, 20, 300});
+  test.AddOutput<int64_t>("min", {3, 3},
+                          {-1, -1, -1,
+                           1, 2, 3,
+                           -70, -80, -90});
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});  //TensorRT: Input batch size is inconsistent
+}
+
+TEST(MathOpTest, Min_12_UInt32) {
+  OpTester test("Min", 12);
+  test.AddInput<uint32_t>("data_0", {1, 3},
+                          {1, 20, 30});
+  test.AddInput<uint32_t>("data_2", {3, 3},
+                          {10, 20, 30,
+                           40, 50, 60,
+                           70, 80, 90});
+  test.AddInput<uint32_t>("data_1", {3, 1},
+                          {1, 20, 30});
+  test.AddOutput<uint32_t>("min", {3, 3},
+                           {1, 1, 1,
+                            1, 20, 20,
+                            1, 20, 30});
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});  //TensorRT: Input batch size is inconsistent
+}
+
+TEST(MathOpTest, Min_12_UInt64) {
+  OpTester test("Min", 12);
+  test.AddInput<uint64_t>("data_0", {1, 3},
+                          {1, 20, 30});
+  test.AddInput<uint64_t>("data_2", {3, 3},
+                          {10, 20, 30,
+                           40, 50, 60,
+                           70, 80, 90});
+  test.AddInput<uint64_t>("data_1", {3, 1},
+                          {1, 20, 30});
+  test.AddOutput<uint64_t>("min", {3, 3},
+                           {1, 1, 1,
+                            1, 20, 20,
+                            1, 20, 30});
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});  //TensorRT: Input batch size is inconsistent
+}
+
 TEST(MathOpTest, Max_6) {
   OpTester test("Max", 6);
   std::vector<int64_t> dims{3, 3};
@@ -854,6 +956,108 @@ TEST(MathOpTest, Max_8_2inputbroadcast) {
   test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});  //TensorRT: Input batch size is inconsistent
 }
 
+TEST(MathOpTest, Max_12_Float) {
+  OpTester test("Max", 12);
+  test.AddInput<float>("data_0", {1, 3},
+                       {1.0f, 2.0f, 3.0f});
+  test.AddInput<float>("data_2", {3, 3},
+                       {10.0f, 20.0f, 30.0f,
+                        40.0f, 50.0f, 60.0f,
+                        70.0f, 80.0f, 90.0f});
+  test.AddInput<float>("data_1", {3, 1},
+                       {-1.0f, -2.0f, 300.0f});
+  test.AddOutput<float>("max", {3, 3},
+                        {10.0f, 20.0f, 30.0f,
+                         40.0f, 50.0f, 60.0f,
+                         300.0f, 300.0f, 300.0f});
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});  //TensorRT: Input batch size is inconsistent
+}
+
+TEST(MathOpTest, Max_12_Double) {
+  OpTester test("Max", 12);
+  test.AddInput<double>("data_0", {1, 3},
+                        {1.0, 2.0, 3.0});
+  test.AddInput<double>("data_2", {3, 3},
+                        {10.0, 20.0, 30.0,
+                         40.0, 50.0, 60.0,
+                         70.0, 80.0, 90.0});
+  test.AddInput<double>("data_1", {3, 1},
+                        {-1.0, -2.0, 300.0});
+  test.AddOutput<double>("max", {3, 3},
+                         {10.0, 20.0, 30.0,
+                          40.0, 50.0, 60.0,
+                          300.0, 300.0, 300.0});
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});  //TensorRT: Input batch size is inconsistent
+}
+
+TEST(MathOpTest, Max_12_Int32) {
+  OpTester test("Max", 12);
+  test.AddInput<int32_t>("data_0", {1, 3},
+                         {1, 2, 3});
+  test.AddInput<int32_t>("data_2", {3, 3},
+                         {10, 20, 30,
+                          40, 50, 60,
+                          70, 80, 90});
+  test.AddInput<int32_t>("data_1", {3, 1},
+                         {-1, -2, 300});
+  test.AddOutput<int32_t>("max", {3, 3},
+                          {10, 20, 30,
+                           40, 50, 60,
+                           300, 300, 300});
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});  //TensorRT: Input batch size is inconsistent
+}
+
+TEST(MathOpTest, Max_12_Int64) {
+  OpTester test("Max", 12);
+  test.AddInput<int64_t>("data_0", {1, 3},
+                         {1, 2, 3});
+  test.AddInput<int64_t>("data_2", {3, 3},
+                         {10, 20, 30,
+                          40, 50, 60,
+                          70, 80, 90});
+  test.AddInput<int64_t>("data_1", {3, 1},
+                         {-1, -2, 300});
+  test.AddOutput<int64_t>("max", {3, 3},
+                          {10, 20, 30,
+                           40, 50, 60,
+                           300, 300, 300});
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});  //TensorRT: Input batch size is inconsistent
+}
+
+TEST(MathOpTest, Max_12_UInt32) {
+  OpTester test("Max", 12);
+  test.AddInput<uint32_t>("data_0", {1, 3},
+                          {1, 2, 3});
+  test.AddInput<uint32_t>("data_2", {3, 3},
+                          {10, 20, 30,
+                           40, 50, 60,
+                           70, 80, 90});
+  test.AddInput<uint32_t>("data_1", {3, 1},
+                          {1, 2, 300});
+  test.AddOutput<uint32_t>("max", {3, 3},
+                           {10, 20, 30,
+                            40, 50, 60,
+                            300, 300, 300});
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});  //TensorRT: Input batch size is inconsistent
+}
+
+TEST(MathOpTest, Max_12_UInt64) {
+  OpTester test("Max", 12);
+  test.AddInput<uint64_t>("data_0", {1, 3},
+                          {1, 2, 3});
+  test.AddInput<uint64_t>("data_2", {3, 3},
+                          {10, 20, 30,
+                           40, 50, 60,
+                           70, 80, 90});
+  test.AddInput<uint64_t>("data_1", {3, 1},
+                          {1, 2, 300});
+  test.AddOutput<uint64_t>("max", {3, 3},
+                           {10, 20, 30,
+                            40, 50, 60,
+                            300, 300, 300});
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider});  //TensorRT: Input batch size is inconsistent
+}
+
 TEST(MathOpTest, Not) {
   OpTester test("Not");
   std::vector<int64_t> dims{2};