[CUDA] Fix clip operator (#25057)

### Description

The result of Clip op is not expected when min > max for CUDA. This
fixes the implementation to align with operator spec:
https://onnx.ai/onnx/operators/onnx__Clip.html

### Motivation and Context

ONNX backend test failure with onnx 1.18:

python onnx_backend_test_series.py
======================================================================
FAIL: test_clip_min_greater_than_max_cuda
(__main__.OnnxBackendNodeModelTest)
----------------------------------------------------------------------
DESIRED: array([1., 1., 1.], dtype=float32)
ACTUAL: array([2., 2., 1.], dtype=float32)

Author: tianleiwu

onnxruntime/core/providers/cuda/math/clip_impl.cu

@@ -8,10 +8,14 @@ namespace onnxruntime {
 namespace cuda {
 template <typename T>
 __global__ void _Clip(const T* input, T* output, const T* min, const T* max, T min_default, T max_default, size_t N) {
-  auto min_val = (min) ? *min : min_default; 
-  auto max_val = (max) ? *max : max_default; 
+  auto min_val = (min) ? *min : min_default;
+  auto max_val = (max) ? *max : max_default;
   CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N);
-  output[id] = (input[id] < min_val) ? min_val : ((input[id] > max_val) ? max_val : input[id]);
+
+  // output = Min(max, Max(input, min)). Note that min might be larger than max, so we need to compute in two steps.
+  auto value = input[id];
+  value = (value < min_val) ? min_val : value;
+  output[id] = (value > max_val) ? max_val : value;
 }
 
 template <typename T>