From db3c07608130e4853bb6d9db66fcf57f95a864e9 Mon Sep 17 00:00:00 2001
From: Jeff Daily <jeff.daily@amd.com>
Date: Mon, 8 Jan 2024 03:06:45 -0800
Subject: [PATCH] [ROCm] do not use failed miopen fusion compile (#19012)

The FusedConv operator for the ROCm EP could fail to compile the fused
operation, in which case it should not attempt to use the failed fusion
plan. In addition, the hash for the miopenConvolutionDescriptor_t for
newer ROCm versions was failing to use all components of the descriptor.
---
 onnxruntime/contrib_ops/rocm/fused_conv.cc | 12 ++++++++++--
 1 file changed, 10 insertions(+), 2 deletions(-)

diff --git a/onnxruntime/contrib_ops/rocm/fused_conv.cc b/onnxruntime/contrib_ops/rocm/fused_conv.cc
index d597e0d57f..63804f79a3 100644
--- a/onnxruntime/contrib_ops/rocm/fused_conv.cc
+++ b/onnxruntime/contrib_ops/rocm/fused_conv.cc
@@ -76,7 +76,12 @@ struct FNVHash {
   void HashConvolutionDescriptor(miopenConvolutionDescriptor_t cdesc) {
     int spatial_dim = 1;
 #if ROCM_VERSION >= 50500
-    miopenGetConvolutionSpatialDim(cdesc, &spatial_dim);
+    MIOPEN_CALL(miopenGetConvolutionSpatialDim(cdesc, &spatial_dim));
+    std::vector<int> pads{spatial_dim};
+    std::vector<int> strides{spatial_dim};
+    std::vector<int> dilations{spatial_dim};
+    miopenConvolutionMode_t mode;
+    MIOPEN_CALL(miopenGetConvolutionNdDescriptor(cdesc, spatial_dim, &spatial_dim, pads.data(), strides.data(), dilations.data(), &mode));
 #else
     // Previous versions of MIOpen doesn't provide API to probe the dimension of a
     // miopenConvolutionDescriptor_t, so we have to guess.
@@ -100,11 +105,12 @@ struct FNVHash {
     pads.resize(spatial_dim);
     strides.resize(spatial_dim);
     dilations.resize(spatial_dim);
+#endif
     (*this) << spatial_dim;
     (*this) << pads;
     (*this) << strides;
     (*this) << dilations;
-#endif
+    (*this) << mode;
   }
 
  private:
@@ -313,6 +319,8 @@ class FusedConv : public onnxruntime::rocm::Conv<T, false> {
       auto ret = miopenCompileFusionPlan(handle, fusion->plan);
       if (miopenStatusSuccess == ret) {
         fusion->compiled_on.insert(handle);
+      } else {
+        return ret;
       }
       return miopenStatusSuccess;
     }