Disable fusion when winograd is used except for 3x3 (#364)

* Disable fusion when winograd is used except for 3x3

* Formatting

src/targets/gpu/fuse_ops.cpp

@@ -13,6 +13,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
+MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_DISABLE_MIOPEN_FUSION)
+
 struct fusion
 {
     using op_t = miopenFusionOpDescriptor_t;
@@ -127,6 +129,8 @@ MIGRAPHX_PRED_MATCHER(bias_shape, instruction_ref ins)
 
 MIGRAPHX_PRED_MATCHER(fusable_conv, instruction_ref ins)
 {
+    if(enabled(MIGRAPHX_DISABLE_MIOPEN_FUSION{}))
+        return false;
     if(ins->name() != "gpu::convolution")
         return false;
     if(ins->get_shape().type() != shape::float_type)
@@ -139,6 +143,10 @@ MIGRAPHX_PRED_MATCHER(fusable_conv, instruction_ref ins)
     if(wei.lens()[1] > 512 and conv.algo != miopenConvolutionFwdAlgoWinograd)
         return false;
     auto op = conv.op;
+    // Dont fuse winograd for non-3x3s since there is no fused windograd for those configs
+    if(conv.algo == miopenConvolutionFwdAlgoWinograd and wei.lens()[2] != 3 and
+       wei.lens()[3] != 3 and op.stride == make_array<size_t>(1, 1))
+        return false;
     return contains({{0, 0}, {1, 1}, {2, 2}}, op.padding) and
            contains({{0, 0}, {1, 1}}, op.stride) and op.dilation == make_array<size_t>(1, 1);
 }