Merge from trunk revision 530203d6e3244c25eda4124f0fa5756ca9a5683e.

13 files changed, 1416 insertions, 1206 deletions

diff --git a/gcc/config/aarch64/aarch64-simd-builtins.def b/gcc/config/aarch64/aarch64-simd-builtins.def
index 48e481c..b787cb9 100644
--- a/gcc/config/aarch64/aarch64-simd-builtins.def
+++ b/gcc/config/aarch64/aarch64-simd-builtins.def
@@ -240,6 +240,18 @@
   BUILTIN_VQW (TERNOP, smlsl_hi, 0, NONE)
   BUILTIN_VQW (TERNOPU, umlsl_hi, 0, NONE)
 
+  /* Implemented by aarch64_<su>mlsl_hi_n<mode>.  */
+  BUILTIN_VQ_HSI (TERNOP, smlsl_hi_n, 0, NONE)
+  BUILTIN_VQ_HSI (TERNOPU, umlsl_hi_n, 0, NONE)
+
+  /* Implemented by aarch64_<su>mlal_hi<mode>.  */
+  BUILTIN_VQW (TERNOP, smlal_hi, 0, NONE)
+  BUILTIN_VQW (TERNOPU, umlal_hi, 0, NONE)
+
+  /* Implemented by aarch64_<su>mlal_hi_n<mode>.  */
+  BUILTIN_VQ_HSI (TERNOP, smlal_hi_n, 0, NONE)
+  BUILTIN_VQ_HSI (TERNOPU, umlal_hi_n, 0, NONE)
+
   BUILTIN_VSQN_HSDI (UNOPUS, sqmovun, 0, NONE)
 
   /* Implemented by aarch64_sqxtun2<mode>.  */
@@ -302,6 +314,16 @@
   BUILTIN_VD_HSI (QUADOPU_LANE, vec_umlsl_lane_, 0, NONE)
   BUILTIN_VD_HSI (QUADOPU_LANE, vec_umlsl_laneq_, 0, NONE)
 
+  BUILTIN_VQ_HSI (QUADOP_LANE, smlal_hi_lane, 0, NONE)
+  BUILTIN_VQ_HSI (QUADOP_LANE, smlal_hi_laneq, 0, NONE)
+  BUILTIN_VQ_HSI (QUADOPU_LANE, umlal_hi_lane, 0, NONE)
+  BUILTIN_VQ_HSI (QUADOPU_LANE, umlal_hi_laneq, 0, NONE)
+
+  BUILTIN_VQ_HSI (QUADOP_LANE, smlsl_hi_lane, 0, NONE)
+  BUILTIN_VQ_HSI (QUADOP_LANE, smlsl_hi_laneq, 0, NONE)
+  BUILTIN_VQ_HSI (QUADOPU_LANE, umlsl_hi_lane, 0, NONE)
+  BUILTIN_VQ_HSI (QUADOPU_LANE, umlsl_hi_laneq, 0, NONE)
+
   BUILTIN_VSD_HSI (BINOP, sqdmull, 0, NONE)
   BUILTIN_VSD_HSI (TERNOP_LANE, sqdmull_lane, 0, NONE)
   BUILTIN_VSD_HSI (TERNOP_LANE, sqdmull_laneq, 0, NONE)
diff --git a/gcc/config/aarch64/aarch64-simd.md b/gcc/config/aarch64/aarch64-simd.md
index 767d673..393bab1 100644
--- a/gcc/config/aarch64/aarch64-simd.md
+++ b/gcc/config/aarch64/aarch64-simd.md
@@ -1976,7 +1976,7 @@
   [(set_attr "type" "neon_mla_<Vetype>_long")]
 )
 
-(define_insn "*aarch64_<su>mlal_hi<mode>"
+(define_insn "aarch64_<su>mlal_hi<mode>_insn"
   [(set (match_operand:<VWIDE> 0 "register_operand" "=w")
         (plus:<VWIDE>
           (mult:<VWIDE>
@@ -1992,6 +1992,49 @@
   [(set_attr "type" "neon_mla_<Vetype>_long")]
 )
 
+(define_expand "aarch64_<su>mlal_hi<mode>"
+  [(match_operand:<VWIDE> 0 "register_operand")
+   (match_operand:<VWIDE> 1 "register_operand")
+   (ANY_EXTEND:<VWIDE>(match_operand:VQW 2 "register_operand"))
+   (match_operand:VQW 3 "register_operand")]
+  "TARGET_SIMD"
+{
+  rtx p = aarch64_simd_vect_par_cnst_half (<MODE>mode, <nunits>, true);
+  emit_insn (gen_aarch64_<su>mlal_hi<mode>_insn (operands[0], operands[1],
+						 operands[2], p, operands[3]));
+  DONE;
+}
+)
+
+(define_insn "aarch64_<su>mlal_hi_n<mode>_insn"
+  [(set (match_operand:<VWIDE> 0 "register_operand" "=w")
+        (plus:<VWIDE>
+          (mult:<VWIDE>
+              (ANY_EXTEND:<VWIDE> (vec_select:<VHALF>
+                 (match_operand:VQ_HSI 2 "register_operand" "w")
+                 (match_operand:VQ_HSI 3 "vect_par_cnst_hi_half" "")))
+              (ANY_EXTEND:<VWIDE> (vec_duplicate:<VCOND>
+	               (match_operand:<VEL> 4 "register_operand" "<h_con>"))))
+          (match_operand:<VWIDE> 1 "register_operand" "0")))]
+  "TARGET_SIMD"
+  "<su>mlal2\t%0.<Vwtype>, %2.<Vtype>, %4.<Vetype>[0]"
+  [(set_attr "type" "neon_mla_<Vetype>_long")]
+)
+
+(define_expand "aarch64_<su>mlal_hi_n<mode>"
+  [(match_operand:<VWIDE> 0 "register_operand")
+   (match_operand:<VWIDE> 1 "register_operand")
+   (ANY_EXTEND:<VWIDE>(match_operand:VQ_HSI 2 "register_operand"))
+   (match_operand:<VEL> 3 "register_operand")]
+  "TARGET_SIMD"
+{
+  rtx p = aarch64_simd_vect_par_cnst_half (<MODE>mode, <nunits>, true);
+  emit_insn (gen_aarch64_<su>mlal_hi_n<mode>_insn (operands[0],
+             operands[1], operands[2], p, operands[3]));
+  DONE;
+}
+)
+
 (define_insn "*aarch64_<su>mlsl_lo<mode>"
   [(set (match_operand:<VWIDE> 0 "register_operand" "=w")
         (minus:<VWIDE>
@@ -2038,6 +2081,35 @@
 }
 )
 
+(define_insn "aarch64_<su>mlsl_hi_n<mode>_insn"
+  [(set (match_operand:<VWIDE> 0 "register_operand" "=w")
+        (minus:<VWIDE>
+          (match_operand:<VWIDE> 1 "register_operand" "0")
+          (mult:<VWIDE>
+            (ANY_EXTEND:<VWIDE> (vec_select:<VHALF>
+              (match_operand:VQ_HSI 2 "register_operand" "w")
+              (match_operand:VQ_HSI 3 "vect_par_cnst_hi_half" "")))
+            (ANY_EXTEND:<VWIDE> (vec_duplicate:<VCOND>
+	            (match_operand:<VEL> 4 "register_operand" "<h_con>"))))))]
+  "TARGET_SIMD"
+  "<su>mlsl2\t%0.<Vwtype>, %2.<Vtype>, %4.<Vetype>[0]"
+  [(set_attr "type" "neon_mla_<Vetype>_long")]
+)
+
+(define_expand "aarch64_<su>mlsl_hi_n<mode>"
+  [(match_operand:<VWIDE> 0 "register_operand")
+   (match_operand:<VWIDE> 1 "register_operand")
+   (ANY_EXTEND:<VWIDE>(match_operand:VQ_HSI 2 "register_operand"))
+   (match_operand:<VEL> 3 "register_operand")]
+  "TARGET_SIMD"
+{
+  rtx p = aarch64_simd_vect_par_cnst_half (<MODE>mode, <nunits>, true);
+  emit_insn (gen_aarch64_<su>mlsl_hi_n<mode>_insn (operands[0],
+             operands[1], operands[2], p, operands[3]));
+  DONE;
+}
+)
+
 (define_insn "aarch64_<su>mlal<mode>"
   [(set (match_operand:<VWIDE> 0 "register_operand" "=w")
         (plus:<VWIDE>
@@ -2215,6 +2287,76 @@
   [(set_attr "type" "neon_mla_<Vetype>_scalar_long")]
 )
 
+(define_insn "aarch64_<su>mlal_hi_lane<mode>_insn"
+  [(set (match_operand:<VWIDE> 0 "register_operand" "=w")
+	(plus:<VWIDE>
+	  (mult:<VWIDE>
+	    (ANY_EXTEND:<VWIDE> (vec_select:<VHALF>
+	      (match_operand:VQ_HSI 2 "register_operand" "w")
+	      (match_operand:VQ_HSI 3 "vect_par_cnst_hi_half" "")))
+	    (ANY_EXTEND:<VWIDE> (vec_duplicate:<VHALF>
+	      (vec_select:<VEL>
+		(match_operand:<VCOND> 4 "register_operand" "<vwx>")
+		(parallel [(match_operand:SI 5 "immediate_operand" "i")])))))
+	  (match_operand:<VWIDE> 1 "register_operand" "0")))]
+  "TARGET_SIMD"
+  {
+    operands[5] = aarch64_endian_lane_rtx (<VCOND>mode, INTVAL (operands[5]));
+    return "<su>mlal2\\t%0.<Vwtype>, %2.<Vtype>, %4.<Vetype>[%5]";
+  }
+  [(set_attr "type" "neon_mla_<Vetype>_scalar_long")]
+)
+
+(define_expand "aarch64_<su>mlal_hi_lane<mode>"
+  [(match_operand:<VWIDE> 0 "register_operand")
+   (match_operand:<VWIDE> 1 "register_operand")
+   (ANY_EXTEND:<VWIDE>(match_operand:VQ_HSI 2 "register_operand"))
+   (match_operand:<VCOND> 3 "register_operand")
+   (match_operand:SI 4 "immediate_operand")]
+  "TARGET_SIMD"
+{
+  rtx p = aarch64_simd_vect_par_cnst_half (<MODE>mode, <nunits>, true);
+  emit_insn (gen_aarch64_<su>mlal_hi_lane<mode>_insn (operands[0],
+	     operands[1], operands[2], p, operands[3], operands[4]));
+  DONE;
+}
+)
+
+(define_insn "aarch64_<su>mlal_hi_laneq<mode>_insn"
+  [(set (match_operand:<VWIDE> 0 "register_operand" "=w")
+	(plus:<VWIDE>
+	  (mult:<VWIDE>
+	    (ANY_EXTEND:<VWIDE> (vec_select:<VHALF>
+	      (match_operand:VQ_HSI 2 "register_operand" "w")
+	      (match_operand:VQ_HSI 3 "vect_par_cnst_hi_half" "")))
+	    (ANY_EXTEND:<VWIDE> (vec_duplicate:<VHALF>
+	      (vec_select:<VEL>
+		(match_operand:<VCONQ> 4 "register_operand" "<vwx>")
+		(parallel [(match_operand:SI 5 "immediate_operand" "i")])))))
+	  (match_operand:<VWIDE> 1 "register_operand" "0")))]
+  "TARGET_SIMD"
+  {
+    operands[5] = aarch64_endian_lane_rtx (<VCONQ>mode, INTVAL (operands[5]));
+    return "<su>mlal2\\t%0.<Vwtype>, %2.<Vtype>, %4.<Vetype>[%5]";
+  }
+  [(set_attr "type" "neon_mla_<Vetype>_scalar_long")]
+)
+
+(define_expand "aarch64_<su>mlal_hi_laneq<mode>"
+  [(match_operand:<VWIDE> 0 "register_operand")
+   (match_operand:<VWIDE> 1 "register_operand")
+   (ANY_EXTEND:<VWIDE>(match_operand:VQ_HSI 2 "register_operand"))
+   (match_operand:<VCONQ> 3 "register_operand")
+   (match_operand:SI 4 "immediate_operand")]
+  "TARGET_SIMD"
+{
+  rtx p = aarch64_simd_vect_par_cnst_half (<MODE>mode, <nunits>, true);
+  emit_insn (gen_aarch64_<su>mlal_hi_laneq<mode>_insn (operands[0],
+	     operands[1], operands[2], p, operands[3], operands[4]));
+  DONE;
+}
+)
+
 (define_insn "aarch64_vec_<su>mlsl_lane<Qlane>"
   [(set (match_operand:<VWIDE> 0 "register_operand" "=w")
    (minus:<VWIDE>
@@ -2235,6 +2377,78 @@
   [(set_attr "type" "neon_mla_<Vetype>_scalar_long")]
 )
 
+(define_insn "aarch64_<su>mlsl_hi_lane<mode>_insn"
+  [(set (match_operand:<VWIDE> 0 "register_operand" "=w")
+	(minus:<VWIDE>
+	  (match_operand:<VWIDE> 1 "register_operand" "0")
+	  (mult:<VWIDE>
+	    (ANY_EXTEND:<VWIDE> (vec_select:<VHALF>
+	      (match_operand:VQ_HSI 2 "register_operand" "w")
+	      (match_operand:VQ_HSI 3 "vect_par_cnst_hi_half" "")))
+	    (ANY_EXTEND:<VWIDE> (vec_duplicate:<VHALF>
+	      (vec_select:<VEL>
+		(match_operand:<VCOND> 4 "register_operand" "<vwx>")
+		(parallel [(match_operand:SI 5 "immediate_operand" "i")]))))
+	  )))]
+  "TARGET_SIMD"
+  {
+    operands[5] = aarch64_endian_lane_rtx (<VCOND>mode, INTVAL (operands[5]));
+    return "<su>mlsl2\\t%0.<Vwtype>, %2.<Vtype>, %4.<Vetype>[%5]";
+  }
+  [(set_attr "type" "neon_mla_<Vetype>_scalar_long")]
+)
+
+(define_expand "aarch64_<su>mlsl_hi_lane<mode>"
+  [(match_operand:<VWIDE> 0 "register_operand")
+   (match_operand:<VWIDE> 1 "register_operand")
+   (ANY_EXTEND:<VWIDE>(match_operand:VQ_HSI 2 "register_operand"))
+   (match_operand:<VCOND> 3 "register_operand")
+   (match_operand:SI 4 "immediate_operand")]
+  "TARGET_SIMD"
+{
+  rtx p = aarch64_simd_vect_par_cnst_half (<MODE>mode, <nunits>, true);
+  emit_insn (gen_aarch64_<su>mlsl_hi_lane<mode>_insn (operands[0],
+	     operands[1], operands[2], p, operands[3], operands[4]));
+  DONE;
+}
+)
+
+(define_insn "aarch64_<su>mlsl_hi_laneq<mode>_insn"
+  [(set (match_operand:<VWIDE> 0 "register_operand" "=w")
+	(minus:<VWIDE>
+	  (match_operand:<VWIDE> 1 "register_operand" "0")
+	  (mult:<VWIDE>
+	    (ANY_EXTEND:<VWIDE> (vec_select:<VHALF>
+	      (match_operand:VQ_HSI 2 "register_operand" "w")
+	      (match_operand:VQ_HSI 3 "vect_par_cnst_hi_half" "")))
+	    (ANY_EXTEND:<VWIDE> (vec_duplicate:<VHALF>
+	      (vec_select:<VEL>
+		(match_operand:<VCONQ> 4 "register_operand" "<vwx>")
+		(parallel [(match_operand:SI 5 "immediate_operand" "i")]))))
+	  )))]
+  "TARGET_SIMD"
+  {
+    operands[5] = aarch64_endian_lane_rtx (<VCONQ>mode, INTVAL (operands[5]));
+    return "<su>mlsl2\\t%0.<Vwtype>, %2.<Vtype>, %4.<Vetype>[%5]";
+  }
+  [(set_attr "type" "neon_mla_<Vetype>_scalar_long")]
+)
+
+(define_expand "aarch64_<su>mlsl_hi_laneq<mode>"
+  [(match_operand:<VWIDE> 0 "register_operand")
+   (match_operand:<VWIDE> 1 "register_operand")
+   (ANY_EXTEND:<VWIDE>(match_operand:VQ_HSI 2 "register_operand"))
+   (match_operand:<VCONQ> 3 "register_operand")
+   (match_operand:SI 4 "immediate_operand")]
+  "TARGET_SIMD"
+{
+  rtx p = aarch64_simd_vect_par_cnst_half (<MODE>mode, <nunits>, true);
+  emit_insn (gen_aarch64_<su>mlsl_hi_laneq<mode>_insn (operands[0],
+	     operands[1], operands[2], p, operands[3], operands[4]));
+  DONE;
+}
+)
+
 ;; FP vector operations.
 ;; AArch64 AdvSIMD supports single-precision (32-bit) and 
 ;; double-precision (64-bit) floating-point data types and arithmetic as
diff --git a/gcc/config/aarch64/arm_neon.h b/gcc/config/aarch64/arm_neon.h
index 4b905d9..d50bd65 100644
--- a/gcc/config/aarch64/arm_neon.h
+++ b/gcc/config/aarch64/arm_neon.h
@@ -7152,236 +7152,138 @@ vmla_u32 (uint32x2_t __a, uint32x2_t __b, uint32x2_t __c)
                                                  (int32x2_t) __c);
 }
 
-#define vmlal_high_lane_s16(a, b, c, d)                                 \
-  __extension__                                                         \
-    ({                                                                  \
-       int16x4_t c_ = (c);                                              \
-       int16x8_t b_ = (b);                                              \
-       int32x4_t a_ = (a);                                              \
-       int32x4_t result;                                                \
-       __asm__ ("smlal2 %0.4s, %2.8h, %3.h[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "x"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline int32x4_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlal_high_lane_s16(int32x4_t __a, int16x8_t __b, int16x4_t __v,
+		    const int __lane)
+{
+  return __builtin_aarch64_smlal_hi_lanev8hi (__a, __b, __v, __lane);
+}
 
-#define vmlal_high_lane_s32(a, b, c, d)                                 \
-  __extension__                                                         \
-    ({                                                                  \
-       int32x2_t c_ = (c);                                              \
-       int32x4_t b_ = (b);                                              \
-       int64x2_t a_ = (a);                                              \
-       int64x2_t result;                                                \
-       __asm__ ("smlal2 %0.2d, %2.4s, %3.s[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "w"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline int64x2_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlal_high_lane_s32(int64x2_t __a, int32x4_t __b, int32x2_t __v,
+		    const int __lane)
+{
+  return __builtin_aarch64_smlal_hi_lanev4si (__a, __b, __v, __lane);
+}
 
-#define vmlal_high_lane_u16(a, b, c, d)                                 \
-  __extension__                                                         \
-    ({                                                                  \
-       uint16x4_t c_ = (c);                                             \
-       uint16x8_t b_ = (b);                                             \
-       uint32x4_t a_ = (a);                                             \
-       uint32x4_t result;                                               \
-       __asm__ ("umlal2 %0.4s, %2.8h, %3.h[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "x"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline uint32x4_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlal_high_lane_u16(uint32x4_t __a, uint16x8_t __b, uint16x4_t __v,
+		    const int __lane)
+{
+  return __builtin_aarch64_umlal_hi_lanev8hi_uuuus (__a, __b, __v, __lane);
+}
 
-#define vmlal_high_lane_u32(a, b, c, d)                                 \
-  __extension__                                                         \
-    ({                                                                  \
-       uint32x2_t c_ = (c);                                             \
-       uint32x4_t b_ = (b);                                             \
-       uint64x2_t a_ = (a);                                             \
-       uint64x2_t result;                                               \
-       __asm__ ("umlal2 %0.2d, %2.4s, %3.s[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "w"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline uint64x2_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlal_high_lane_u32(uint64x2_t __a, uint32x4_t __b, uint32x2_t __v,
+		    const int __lane)
+{
+  return __builtin_aarch64_umlal_hi_lanev4si_uuuus (__a, __b, __v, __lane);
+}
 
-#define vmlal_high_laneq_s16(a, b, c, d)                                \
-  __extension__                                                         \
-    ({                                                                  \
-       int16x8_t c_ = (c);                                              \
-       int16x8_t b_ = (b);                                              \
-       int32x4_t a_ = (a);                                              \
-       int32x4_t result;                                                \
-       __asm__ ("smlal2 %0.4s, %2.8h, %3.h[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "x"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline int32x4_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlal_high_laneq_s16(int32x4_t __a, int16x8_t __b, int16x8_t __v,
+		     const int __lane)
+{
+  return __builtin_aarch64_smlal_hi_laneqv8hi (__a, __b, __v, __lane);
+}
 
-#define vmlal_high_laneq_s32(a, b, c, d)                                \
-  __extension__                                                         \
-    ({                                                                  \
-       int32x4_t c_ = (c);                                              \
-       int32x4_t b_ = (b);                                              \
-       int64x2_t a_ = (a);                                              \
-       int64x2_t result;                                                \
-       __asm__ ("smlal2 %0.2d, %2.4s, %3.s[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "w"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline int64x2_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlal_high_laneq_s32(int64x2_t __a, int32x4_t __b, int32x4_t __v,
+		     const int __lane)
+{
+  return __builtin_aarch64_smlal_hi_laneqv4si (__a, __b, __v, __lane);
+}
 
-#define vmlal_high_laneq_u16(a, b, c, d)                                \
-  __extension__                                                         \
-    ({                                                                  \
-       uint16x8_t c_ = (c);                                             \
-       uint16x8_t b_ = (b);                                             \
-       uint32x4_t a_ = (a);                                             \
-       uint32x4_t result;                                               \
-       __asm__ ("umlal2 %0.4s, %2.8h, %3.h[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "x"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline uint32x4_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlal_high_laneq_u16(uint32x4_t __a, uint16x8_t __b, uint16x8_t __v,
+		     const int __lane)
+{
+  return __builtin_aarch64_umlal_hi_laneqv8hi_uuuus (__a, __b, __v, __lane);
+}
 
-#define vmlal_high_laneq_u32(a, b, c, d)                                \
-  __extension__                                                         \
-    ({                                                                  \
-       uint32x4_t c_ = (c);                                             \
-       uint32x4_t b_ = (b);                                             \
-       uint64x2_t a_ = (a);                                             \
-       uint64x2_t result;                                               \
-       __asm__ ("umlal2 %0.2d, %2.4s, %3.s[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "w"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline uint64x2_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlal_high_laneq_u32(uint64x2_t __a, uint32x4_t __b, uint32x4_t __v,
+		     const int __lane)
+{
+  return __builtin_aarch64_umlal_hi_laneqv4si_uuuus (__a, __b, __v, __lane);
+}
 
 __extension__ extern __inline int32x4_t
 __attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
 vmlal_high_n_s16 (int32x4_t __a, int16x8_t __b, int16_t __c)
 {
-  int32x4_t __result;
-  __asm__ ("smlal2 %0.4s,%2.8h,%3.h[0]"
-           : "=w"(__result)
-           : "0"(__a), "w"(__b), "x"(__c)
-           : /* No clobbers */);
-  return __result;
+  return __builtin_aarch64_smlal_hi_nv8hi (__a, __b, __c);
 }
 
 __extension__ extern __inline int64x2_t
 __attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
 vmlal_high_n_s32 (int64x2_t __a, int32x4_t __b, int32_t __c)
 {
-  int64x2_t __result;
-  __asm__ ("smlal2 %0.2d,%2.4s,%3.s[0]"
-           : "=w"(__result)
-           : "0"(__a), "w"(__b), "w"(__c)
-           : /* No clobbers */);
-  return __result;
+  return __builtin_aarch64_smlal_hi_nv4si (__a, __b, __c);
 }
 
 __extension__ extern __inline uint32x4_t
 __attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
 vmlal_high_n_u16 (uint32x4_t __a, uint16x8_t __b, uint16_t __c)
 {
-  uint32x4_t __result;
-  __asm__ ("umlal2 %0.4s,%2.8h,%3.h[0]"
-           : "=w"(__result)
-           : "0"(__a), "w"(__b), "x"(__c)
-           : /* No clobbers */);
-  return __result;
+  return __builtin_aarch64_umlal_hi_nv8hi_uuuu (__a, __b, __c);
 }
 
 __extension__ extern __inline uint64x2_t
 __attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
 vmlal_high_n_u32 (uint64x2_t __a, uint32x4_t __b, uint32_t __c)
 {
-  uint64x2_t __result;
-  __asm__ ("umlal2 %0.2d,%2.4s,%3.s[0]"
-           : "=w"(__result)
-           : "0"(__a), "w"(__b), "w"(__c)
-           : /* No clobbers */);
-  return __result;
+  return __builtin_aarch64_umlal_hi_nv4si_uuuu (__a, __b, __c);
 }
 
 __extension__ extern __inline int16x8_t
 __attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
 vmlal_high_s8 (int16x8_t __a, int8x16_t __b, int8x16_t __c)
 {
-  int16x8_t __result;
-  __asm__ ("smlal2 %0.8h,%2.16b,%3.16b"
-           : "=w"(__result)
-           : "0"(__a), "w"(__b), "w"(__c)
-           : /* No clobbers */);
-  return __result;
+  return __builtin_aarch64_smlal_hiv16qi (__a, __b, __c);
 }
 
 __extension__ extern __inline int32x4_t
 __attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
 vmlal_high_s16 (int32x4_t __a, int16x8_t __b, int16x8_t __c)
 {
-  int32x4_t __result;
-  __asm__ ("smlal2 %0.4s,%2.8h,%3.8h"
-           : "=w"(__result)
-           : "0"(__a), "w"(__b), "w"(__c)
-           : /* No clobbers */);
-  return __result;
+  return __builtin_aarch64_smlal_hiv8hi (__a, __b, __c);
 }
 
 __extension__ extern __inline int64x2_t
 __attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
 vmlal_high_s32 (int64x2_t __a, int32x4_t __b, int32x4_t __c)
 {
-  int64x2_t __result;
-  __asm__ ("smlal2 %0.2d,%2.4s,%3.4s"
-           : "=w"(__result)
-           : "0"(__a), "w"(__b), "w"(__c)
-           : /* No clobbers */);
-  return __result;
+  return __builtin_aarch64_smlal_hiv4si (__a, __b, __c);
 }
 
 __extension__ extern __inline uint16x8_t
 __attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
 vmlal_high_u8 (uint16x8_t __a, uint8x16_t __b, uint8x16_t __c)
 {
-  uint16x8_t __result;
-  __asm__ ("umlal2 %0.8h,%2.16b,%3.16b"
-           : "=w"(__result)
-           : "0"(__a), "w"(__b), "w"(__c)
-           : /* No clobbers */);
-  return __result;
+  return __builtin_aarch64_umlal_hiv16qi_uuuu (__a, __b, __c);
 }
 
 __extension__ extern __inline uint32x4_t
 __attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
 vmlal_high_u16 (uint32x4_t __a, uint16x8_t __b, uint16x8_t __c)
 {
-  uint32x4_t __result;
-  __asm__ ("umlal2 %0.4s,%2.8h,%3.8h"
-           : "=w"(__result)
-           : "0"(__a), "w"(__b), "w"(__c)
-           : /* No clobbers */);
-  return __result;
+  return __builtin_aarch64_umlal_hiv8hi_uuuu (__a, __b, __c);
 }
 
 __extension__ extern __inline uint64x2_t
 __attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
 vmlal_high_u32 (uint64x2_t __a, uint32x4_t __b, uint32x4_t __c)
 {
-  uint64x2_t __result;
-  __asm__ ("umlal2 %0.2d,%2.4s,%3.4s"
-           : "=w"(__result)
-           : "0"(__a), "w"(__b), "w"(__c)
-           : /* No clobbers */);
-  return __result;
+  return __builtin_aarch64_umlal_hiv4si_uuuu (__a, __b, __c);
 }
 
 __extension__ extern __inline int32x4_t
@@ -7696,164 +7598,96 @@ vmls_u32 (uint32x2_t __a, uint32x2_t __b, uint32x2_t __c)
                                                  (int32x2_t) __c);
 }
 
-#define vmlsl_high_lane_s16(a, b, c, d)                                 \
-  __extension__                                                         \
-    ({                                                                  \
-       int16x4_t c_ = (c);                                              \
-       int16x8_t b_ = (b);                                              \
-       int32x4_t a_ = (a);                                              \
-       int32x4_t result;                                                \
-       __asm__ ("smlsl2 %0.4s, %2.8h, %3.h[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "x"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline int32x4_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlsl_high_lane_s16(int32x4_t __a, int16x8_t __b, int16x4_t __v,
+		    const int __lane)
+{
+  return __builtin_aarch64_smlsl_hi_lanev8hi (__a, __b, __v, __lane);
+}
 
-#define vmlsl_high_lane_s32(a, b, c, d)                                 \
-  __extension__                                                         \
-    ({                                                                  \
-       int32x2_t c_ = (c);                                              \
-       int32x4_t b_ = (b);                                              \
-       int64x2_t a_ = (a);                                              \
-       int64x2_t result;                                                \
-       __asm__ ("smlsl2 %0.2d, %2.4s, %3.s[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "w"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline int64x2_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlsl_high_lane_s32(int64x2_t __a, int32x4_t __b, int32x2_t __v,
+		    const int __lane)
+{
+  return __builtin_aarch64_smlsl_hi_lanev4si (__a, __b, __v, __lane);
+}
 
-#define vmlsl_high_lane_u16(a, b, c, d)                                 \
-  __extension__                                                         \
-    ({                                                                  \
-       uint16x4_t c_ = (c);                                             \
-       uint16x8_t b_ = (b);                                             \
-       uint32x4_t a_ = (a);                                             \
-       uint32x4_t result;                                               \
-       __asm__ ("umlsl2 %0.4s, %2.8h, %3.h[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "x"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline uint32x4_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlsl_high_lane_u16(uint32x4_t __a, uint16x8_t __b, uint16x4_t __v,
+		    const int __lane)
+{
+  return __builtin_aarch64_umlsl_hi_lanev8hi_uuuus (__a, __b, __v, __lane);
+}
 
-#define vmlsl_high_lane_u32(a, b, c, d)                                 \
-  __extension__                                                         \
-    ({                                                                  \
-       uint32x2_t c_ = (c);                                             \
-       uint32x4_t b_ = (b);                                             \
-       uint64x2_t a_ = (a);                                             \
-       uint64x2_t result;                                               \
-       __asm__ ("umlsl2 %0.2d, %2.4s, %3.s[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "w"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline uint64x2_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlsl_high_lane_u32(uint64x2_t __a, uint32x4_t __b, uint32x2_t __v,
+		    const int __lane)
+{
+  return __builtin_aarch64_umlsl_hi_lanev4si_uuuus (__a, __b, __v, __lane);
+}
 
-#define vmlsl_high_laneq_s16(a, b, c, d)                                \
-  __extension__                                                         \
-    ({                                                                  \
-       int16x8_t c_ = (c);                                              \
-       int16x8_t b_ = (b);                                              \
-       int32x4_t a_ = (a);                                              \
-       int32x4_t result;                                                \
-       __asm__ ("smlsl2 %0.4s, %2.8h, %3.h[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "x"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline int32x4_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlsl_high_laneq_s16(int32x4_t __a, int16x8_t __b, int16x8_t __v,
+		     const int __lane)
+{
+  return __builtin_aarch64_smlsl_hi_laneqv8hi (__a, __b, __v, __lane);
+}
 
-#define vmlsl_high_laneq_s32(a, b, c, d)                                \
-  __extension__                                                         \
-    ({                                                                  \
-       int32x4_t c_ = (c);                                              \
-       int32x4_t b_ = (b);                                              \
-       int64x2_t a_ = (a);                                              \
-       int64x2_t result;                                                \
-       __asm__ ("smlsl2 %0.2d, %2.4s, %3.s[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "w"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline int64x2_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlsl_high_laneq_s32(int64x2_t __a, int32x4_t __b, int32x4_t __v,
+		     const int __lane)
+{
+  return __builtin_aarch64_smlsl_hi_laneqv4si (__a, __b, __v, __lane);
+}
 
-#define vmlsl_high_laneq_u16(a, b, c, d)                                \
-  __extension__                                                         \
-    ({                                                                  \
-       uint16x8_t c_ = (c);                                             \
-       uint16x8_t b_ = (b);                                             \
-       uint32x4_t a_ = (a);                                             \
-       uint32x4_t result;                                               \
-       __asm__ ("umlsl2 %0.4s, %2.8h, %3.h[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "x"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline uint32x4_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlsl_high_laneq_u16(uint32x4_t __a, uint16x8_t __b, uint16x8_t __v,
+		     const int __lane)
+{
+  return __builtin_aarch64_umlsl_hi_laneqv8hi_uuuus (__a, __b, __v, __lane);
+}
 
-#define vmlsl_high_laneq_u32(a, b, c, d)                                \
-  __extension__                                                         \
-    ({                                                                  \
-       uint32x4_t c_ = (c);                                             \
-       uint32x4_t b_ = (b);                                             \
-       uint64x2_t a_ = (a);                                             \
-       uint64x2_t result;                                               \
-       __asm__ ("umlsl2 %0.2d, %2.4s, %3.s[%4]"                         \
-                : "=w"(result)                                          \
-                : "0"(a_), "w"(b_), "w"(c_), "i"(d)                     \
-                : /* No clobbers */);                                   \
-       result;                                                          \
-     })
+__extension__ extern __inline uint64x2_t
+__attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
+vmlsl_high_laneq_u32(uint64x2_t __a, uint32x4_t __b, uint32x4_t __v,
+		     const int __lane)
+{
+  return __builtin_aarch64_umlsl_hi_laneqv4si_uuuus (__a, __b, __v, __lane);
+}
 
 __extension__ extern __inline int32x4_t
 __attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
 vmlsl_high_n_s16 (int32x4_t __a, int16x8_t __b, int16_t __c)
 {
-  int32x4_t __result;
-  __asm__ ("smlsl2 %0.4s, %2.8h, %3.h[0]"
-           : "=w"(__result)
-           : "0"(__a), "w"(__b), "x"(__c)
-           : /* No clobbers */);
-  return __result;
+  return __builtin_aarch64_smlsl_hi_nv8hi (__a, __b, __c);
 }
 
 __extension__ extern __inline int64x2_t
 __attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
 vmlsl_high_n_s32 (int64x2_t __a, int32x4_t __b, int32_t __c)
 {
-  int64x2_t __result;
-  __asm__ ("smlsl2 %0.2d, %2.4s, %3.s[0]"
-           : "=w"(__result)
-           : "0"(__a), "w"(__b), "w"(__c)
-           : /* No clobbers */);
-  return __result;
+  return __builtin_aarch64_smlsl_hi_nv4si (__a, __b, __c);
 }
 
 __extension__ extern __inline uint32x4_t
 __attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
 vmlsl_high_n_u16 (uint32x4_t __a, uint16x8_t __b, uint16_t __c)
 {
-  uint32x4_t __result;
-  __asm__ ("umlsl2 %0.4s, %2.8h, %3.h[0]"
-           : "=w"(__result)
-           : "0"(__a), "w"(__b), "x"(__c)
-           : /* No clobbers */);
-  return __result;
+  return __builtin_aarch64_umlsl_hi_nv8hi_uuuu (__a, __b, __c);
 }
 
 __extension__ extern __inline uint64x2_t
 __attribute__ ((__always_inline__, __gnu_inline__, __artificial__))
 vmlsl_high_n_u32 (uint64x2_t __a, uint32x4_t __b, uint32_t __c)
 {
-  uint64x2_t __result;
-  __asm__ ("umlsl2 %0.2d, %2.4s, %3.s[0]"
-           : "=w"(__result)
-           : "0"(__a), "w"(__b), "w"(__c)
-           : /* No clobbers */);
-  return __result;
+  return __builtin_aarch64_umlsl_hi_nv4si_uuuu (__a, __b, __c);
 }
 
 __extension__ extern __inline int16x8_t
diff --git a/gcc/config/gcn/gcn-opts.h b/gcc/config/gcn/gcn-opts.h
index ed9b451..ed67d01 100644
--- a/gcc/config/gcn/gcn-opts.h
+++ b/gcc/config/gcn/gcn-opts.h
@@ -22,7 +22,8 @@ enum processor_type
 {
   PROCESSOR_FIJI,    // gfx803
   PROCESSOR_VEGA10,  // gfx900
-  PROCESSOR_VEGA20   // gfx906
+  PROCESSOR_VEGA20,  // gfx906
+  PROCESSOR_GFX908   // as yet unnamed
 };
 
 /* Set in gcn_option_override.  */
diff --git a/gcc/config/gcn/gcn.c b/gcc/config/gcn/gcn.c
index 2351b24..e8bb0b6 100644
--- a/gcc/config/gcn/gcn.c
+++ b/gcc/config/gcn/gcn.c
@@ -2589,6 +2589,8 @@ gcn_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait,
 	return gcn_arch == PROCESSOR_VEGA10;
       if (strcmp (name, "gfx906") == 0)
 	return gcn_arch == PROCESSOR_VEGA20;
+      if (strcmp (name, "gfx908") == 0)
+	return gcn_arch == PROCESSOR_GFX908;
       return 0;
     default:
       gcc_unreachable ();
@@ -5030,6 +5032,7 @@ output_file_start (void)
     case PROCESSOR_FIJI: cpu = "gfx803"; break;
     case PROCESSOR_VEGA10: cpu = "gfx900"; break;
     case PROCESSOR_VEGA20: cpu = "gfx906"; break;
+    case PROCESSOR_GFX908: cpu = "gfx908+sram-ecc"; break;
     default: gcc_unreachable ();
     }
 
diff --git a/gcc/config/gcn/gcn.opt b/gcc/config/gcn/gcn.opt
index 7fd84f8..767d458 100644
--- a/gcc/config/gcn/gcn.opt
+++ b/gcc/config/gcn/gcn.opt
@@ -34,6 +34,9 @@ Enum(gpu_type) String(gfx900) Value(PROCESSOR_VEGA10)
 EnumValue
 Enum(gpu_type) String(gfx906) Value(PROCESSOR_VEGA20)
 
+EnumValue
+Enum(gpu_type) String(gfx908) Value(PROCESSOR_GFX908)
+
 march=
 Target RejectNegative Joined ToLower Enum(gpu_type) Var(gcn_arch) Init(PROCESSOR_FIJI)
 Specify the name of the target GPU.
diff --git a/gcc/config/gcn/mkoffload.c b/gcc/config/gcn/mkoffload.c
index eb1c717..dc9d518 100644
--- a/gcc/config/gcn/mkoffload.c
+++ b/gcc/config/gcn/mkoffload.c
@@ -51,6 +51,8 @@
 #define EF_AMDGPU_MACH_AMDGCN_GFX900 0x2c
 #undef  EF_AMDGPU_MACH_AMDGCN_GFX906
 #define EF_AMDGPU_MACH_AMDGCN_GFX906 0x2f
+#undef  EF_AMDGPU_MACH_AMDGCN_GFX908
+#define EF_AMDGPU_MACH_AMDGCN_GFX908 0x230  // Assume SRAM-ECC enabled.
 
 #ifndef R_AMDGPU_NONE
 #define R_AMDGPU_NONE		0
@@ -856,6 +858,8 @@ main (int argc, char **argv)
 	elf_arch = EF_AMDGPU_MACH_AMDGCN_GFX900;
       else if (strcmp (argv[i], "-march=gfx906") == 0)
 	elf_arch = EF_AMDGPU_MACH_AMDGCN_GFX906;
+      else if (strcmp (argv[i], "-march=gfx908") == 0)
+	elf_arch = EF_AMDGPU_MACH_AMDGCN_GFX908;
     }
 
   if (!(fopenacc ^ fopenmp))
diff --git a/gcc/config/gcn/t-gcn-hsa b/gcc/config/gcn/t-gcn-hsa
index bf47da7..ee4d9b3 100644
--- a/gcc/config/gcn/t-gcn-hsa
+++ b/gcc/config/gcn/t-gcn-hsa
@@ -42,8 +42,8 @@ ALL_HOST_OBJS += gcn-run.o
 gcn-run$(exeext): gcn-run.o
 	+$(LINKER) $(ALL_LINKERFLAGS) $(LDFLAGS) -o $@ $< -ldl
 
-MULTILIB_OPTIONS = march=gfx900/march=gfx906
-MULTILIB_DIRNAMES = gfx900 gfx906
+MULTILIB_OPTIONS = march=gfx900/march=gfx906/march=gfx908
+MULTILIB_DIRNAMES = gfx900 gfx906 gfx908
 
 gcn-tree.o: $(srcdir)/config/gcn/gcn-tree.c
 	$(COMPILE) $<
diff --git a/gcc/config/gcn/t-omp-device b/gcc/config/gcn/t-omp-device
index d9809d5..8461c43 100644
--- a/gcc/config/gcn/t-omp-device
+++ b/gcc/config/gcn/t-omp-device
@@ -1,4 +1,4 @@
 omp-device-properties-gcn: $(srcdir)/config/gcn/gcn.c
 	echo kind: gpu > $@
 	echo arch: gcn >> $@
-	echo isa: fiji gfx900 gfx906 >> $@
+	echo isa: fiji gfx900 gfx906 gfx908 >> $@
diff --git a/gcc/config/i386/mmx.md b/gcc/config/i386/mmx.md
index a6ddc71..9e5a4d1 100644
--- a/gcc/config/i386/mmx.md
+++ b/gcc/config/i386/mmx.md
@@ -1528,9 +1528,9 @@
    (set_attr "mode" "DI,TI,TI")])
 
 (define_expand "<insn><mode>3"
-  [(set (match_operand:MMXMODE248 0 "register_operand")
-        (any_lshift:MMXMODE248
-	  (match_operand:MMXMODE248 1 "register_operand")
+  [(set (match_operand:MMXMODE24 0 "register_operand")
+        (any_lshift:MMXMODE24
+	  (match_operand:MMXMODE24 1 "register_operand")
 	  (match_operand:DI 2 "nonmemory_operand")))]
   "TARGET_MMX_WITH_SSE")
 
diff --git a/gcc/config/rs6000/fusion.md b/gcc/config/rs6000/fusion.md
index 1ddbe7f..737a6da 100644
--- a/gcc/config/rs6000/fusion.md
+++ b/gcc/config/rs6000/fusion.md
@@ -357,2176 +357,2304 @@
 
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: and op and rtl and inv 0 comp 0
-;; inner: and op and rtl and inv 0 comp 0
+;; scalar and -> and
 (define_insn "*fuse_and_and"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "%r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "%r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    and %3,%1,%0\;and %3,%3,%2
-   and %0,%1,%0\;and %0,%0,%2
-   and %1,%1,%0\;and %1,%1,%2
+   and %3,%1,%0\;and %3,%3,%2
+   and %3,%1,%0\;and %3,%3,%2
    and %4,%1,%0\;and %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: and op and rtl and inv 0 comp 0
-;; inner: andc op andc rtl and inv 0 comp 1
+;; scalar andc -> and
 (define_insn "*fuse_andc_and"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    andc %3,%1,%0\;and %3,%3,%2
-   andc %0,%1,%0\;and %0,%0,%2
-   andc %1,%1,%0\;and %1,%1,%2
+   andc %3,%1,%0\;and %3,%3,%2
+   andc %3,%1,%0\;and %3,%3,%2
    andc %4,%1,%0\;and %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: and op and rtl and inv 0 comp 0
-;; inner: eqv op eqv rtl xor inv 1 comp 0
+;; scalar eqv -> and
 (define_insn "*fuse_eqv_and"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    eqv %3,%1,%0\;and %3,%3,%2
-   eqv %0,%1,%0\;and %0,%0,%2
-   eqv %1,%1,%0\;and %1,%1,%2
+   eqv %3,%1,%0\;and %3,%3,%2
+   eqv %3,%1,%0\;and %3,%3,%2
    eqv %4,%1,%0\;and %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: and op and rtl and inv 0 comp 0
-;; inner: nand op nand rtl ior inv 0 comp 3
+;; scalar nand -> and
 (define_insn "*fuse_nand_and"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nand %3,%1,%0\;and %3,%3,%2
-   nand %0,%1,%0\;and %0,%0,%2
-   nand %1,%1,%0\;and %1,%1,%2
+   nand %3,%1,%0\;and %3,%3,%2
+   nand %3,%1,%0\;and %3,%3,%2
    nand %4,%1,%0\;and %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: and op and rtl and inv 0 comp 0
-;; inner: nor op nor rtl and inv 0 comp 3
+;; scalar nor -> and
 (define_insn "*fuse_nor_and"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nor %3,%1,%0\;and %3,%3,%2
-   nor %0,%1,%0\;and %0,%0,%2
-   nor %1,%1,%0\;and %1,%1,%2
+   nor %3,%1,%0\;and %3,%3,%2
+   nor %3,%1,%0\;and %3,%3,%2
    nor %4,%1,%0\;and %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: and op and rtl and inv 0 comp 0
-;; inner: or op or rtl ior inv 0 comp 0
+;; scalar or -> and
 (define_insn "*fuse_or_and"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    or %3,%1,%0\;and %3,%3,%2
-   or %0,%1,%0\;and %0,%0,%2
-   or %1,%1,%0\;and %1,%1,%2
+   or %3,%1,%0\;and %3,%3,%2
+   or %3,%1,%0\;and %3,%3,%2
    or %4,%1,%0\;and %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: and op and rtl and inv 0 comp 0
-;; inner: orc op orc rtl ior inv 0 comp 1
+;; scalar orc -> and
 (define_insn "*fuse_orc_and"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    orc %3,%1,%0\;and %3,%3,%2
-   orc %0,%1,%0\;and %0,%0,%2
-   orc %1,%1,%0\;and %1,%1,%2
+   orc %3,%1,%0\;and %3,%3,%2
+   orc %3,%1,%0\;and %3,%3,%2
    orc %4,%1,%0\;and %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: and op and rtl and inv 0 comp 0
-;; inner: xor op xor rtl xor inv 0 comp 0
+;; scalar xor -> and
 (define_insn "*fuse_xor_and"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    xor %3,%1,%0\;and %3,%3,%2
-   xor %0,%1,%0\;and %0,%0,%2
-   xor %1,%1,%0\;and %1,%1,%2
+   xor %3,%1,%0\;and %3,%3,%2
+   xor %3,%1,%0\;and %3,%3,%2
    xor %4,%1,%0\;and %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: andc op andc rtl and inv 0 comp 1
-;; inner: and op and rtl and inv 0 comp 0
+;; scalar and -> andc
 (define_insn "*fuse_and_andc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    and %3,%1,%0\;andc %3,%3,%2
-   and %0,%1,%0\;andc %0,%0,%2
-   and %1,%1,%0\;andc %1,%1,%2
+   and %3,%1,%0\;andc %3,%3,%2
+   and %3,%1,%0\;andc %3,%3,%2
    and %4,%1,%0\;andc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: andc op andc rtl and inv 0 comp 1
-;; inner: andc op andc rtl and inv 0 comp 1
+;; scalar andc -> andc
 (define_insn "*fuse_andc_andc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    andc %3,%1,%0\;andc %3,%3,%2
-   andc %0,%1,%0\;andc %0,%0,%2
-   andc %1,%1,%0\;andc %1,%1,%2
+   andc %3,%1,%0\;andc %3,%3,%2
+   andc %3,%1,%0\;andc %3,%3,%2
    andc %4,%1,%0\;andc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: andc op andc rtl and inv 0 comp 1
-;; inner: eqv op eqv rtl xor inv 1 comp 0
+;; scalar eqv -> andc
 (define_insn "*fuse_eqv_andc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    eqv %3,%1,%0\;andc %3,%3,%2
-   eqv %0,%1,%0\;andc %0,%0,%2
-   eqv %1,%1,%0\;andc %1,%1,%2
+   eqv %3,%1,%0\;andc %3,%3,%2
+   eqv %3,%1,%0\;andc %3,%3,%2
    eqv %4,%1,%0\;andc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: andc op andc rtl and inv 0 comp 1
-;; inner: nand op nand rtl ior inv 0 comp 3
+;; scalar nand -> andc
 (define_insn "*fuse_nand_andc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nand %3,%1,%0\;andc %3,%3,%2
-   nand %0,%1,%0\;andc %0,%0,%2
-   nand %1,%1,%0\;andc %1,%1,%2
+   nand %3,%1,%0\;andc %3,%3,%2
+   nand %3,%1,%0\;andc %3,%3,%2
    nand %4,%1,%0\;andc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: andc op andc rtl and inv 0 comp 1
-;; inner: nor op nor rtl and inv 0 comp 3
+;; scalar nor -> andc
 (define_insn "*fuse_nor_andc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nor %3,%1,%0\;andc %3,%3,%2
-   nor %0,%1,%0\;andc %0,%0,%2
-   nor %1,%1,%0\;andc %1,%1,%2
+   nor %3,%1,%0\;andc %3,%3,%2
+   nor %3,%1,%0\;andc %3,%3,%2
    nor %4,%1,%0\;andc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: andc op andc rtl and inv 0 comp 1
-;; inner: or op or rtl ior inv 0 comp 0
+;; scalar or -> andc
 (define_insn "*fuse_or_andc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    or %3,%1,%0\;andc %3,%3,%2
-   or %0,%1,%0\;andc %0,%0,%2
-   or %1,%1,%0\;andc %1,%1,%2
+   or %3,%1,%0\;andc %3,%3,%2
+   or %3,%1,%0\;andc %3,%3,%2
    or %4,%1,%0\;andc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: andc op andc rtl and inv 0 comp 1
-;; inner: orc op orc rtl ior inv 0 comp 1
+;; scalar orc -> andc
 (define_insn "*fuse_orc_andc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    orc %3,%1,%0\;andc %3,%3,%2
-   orc %0,%1,%0\;andc %0,%0,%2
-   orc %1,%1,%0\;andc %1,%1,%2
+   orc %3,%1,%0\;andc %3,%3,%2
+   orc %3,%1,%0\;andc %3,%3,%2
    orc %4,%1,%0\;andc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: andc op andc rtl and inv 0 comp 1
-;; inner: xor op xor rtl xor inv 0 comp 0
+;; scalar xor -> andc
 (define_insn "*fuse_xor_andc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    xor %3,%1,%0\;andc %3,%3,%2
-   xor %0,%1,%0\;andc %0,%0,%2
-   xor %1,%1,%0\;andc %1,%1,%2
+   xor %3,%1,%0\;andc %3,%3,%2
+   xor %3,%1,%0\;andc %3,%3,%2
    xor %4,%1,%0\;andc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: eqv op eqv rtl xor inv 1 comp 0
-;; inner: and op and rtl and inv 0 comp 0
+;; scalar and -> eqv
 (define_insn "*fuse_and_eqv"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (not:GPR (xor:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (not:GPR (xor:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    and %3,%1,%0\;eqv %3,%3,%2
-   and %0,%1,%0\;eqv %0,%0,%2
-   and %1,%1,%0\;eqv %1,%1,%2
+   and %3,%1,%0\;eqv %3,%3,%2
+   and %3,%1,%0\;eqv %3,%3,%2
    and %4,%1,%0\;eqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: eqv op eqv rtl xor inv 1 comp 0
-;; inner: andc op andc rtl and inv 0 comp 1
+;; scalar andc -> eqv
 (define_insn "*fuse_andc_eqv"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (not:GPR (xor:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (not:GPR (xor:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    andc %3,%1,%0\;eqv %3,%3,%2
-   andc %0,%1,%0\;eqv %0,%0,%2
-   andc %1,%1,%0\;eqv %1,%1,%2
+   andc %3,%1,%0\;eqv %3,%3,%2
+   andc %3,%1,%0\;eqv %3,%3,%2
    andc %4,%1,%0\;eqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: eqv op eqv rtl xor inv 1 comp 0
-;; inner: eqv op eqv rtl xor inv 1 comp 0
+;; scalar eqv -> eqv
 (define_insn "*fuse_eqv_eqv"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (not:GPR (xor:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "%r,r,r,r"))) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (not:GPR (xor:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "%r,r,r,r")))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    eqv %3,%1,%0\;eqv %3,%3,%2
-   eqv %0,%1,%0\;eqv %0,%0,%2
-   eqv %1,%1,%0\;eqv %1,%1,%2
+   eqv %3,%1,%0\;eqv %3,%3,%2
+   eqv %3,%1,%0\;eqv %3,%3,%2
    eqv %4,%1,%0\;eqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: eqv op eqv rtl xor inv 1 comp 0
-;; inner: nand op nand rtl ior inv 0 comp 3
+;; scalar nand -> eqv
 (define_insn "*fuse_nand_eqv"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (not:GPR (xor:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (not:GPR (xor:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nand %3,%1,%0\;eqv %3,%3,%2
-   nand %0,%1,%0\;eqv %0,%0,%2
-   nand %1,%1,%0\;eqv %1,%1,%2
+   nand %3,%1,%0\;eqv %3,%3,%2
+   nand %3,%1,%0\;eqv %3,%3,%2
    nand %4,%1,%0\;eqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: eqv op eqv rtl xor inv 1 comp 0
-;; inner: nor op nor rtl and inv 0 comp 3
+;; scalar nor -> eqv
 (define_insn "*fuse_nor_eqv"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (not:GPR (xor:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (not:GPR (xor:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nor %3,%1,%0\;eqv %3,%3,%2
-   nor %0,%1,%0\;eqv %0,%0,%2
-   nor %1,%1,%0\;eqv %1,%1,%2
+   nor %3,%1,%0\;eqv %3,%3,%2
+   nor %3,%1,%0\;eqv %3,%3,%2
    nor %4,%1,%0\;eqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: eqv op eqv rtl xor inv 1 comp 0
-;; inner: or op or rtl ior inv 0 comp 0
+;; scalar or -> eqv
 (define_insn "*fuse_or_eqv"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (not:GPR (xor:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (not:GPR (xor:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    or %3,%1,%0\;eqv %3,%3,%2
-   or %0,%1,%0\;eqv %0,%0,%2
-   or %1,%1,%0\;eqv %1,%1,%2
+   or %3,%1,%0\;eqv %3,%3,%2
+   or %3,%1,%0\;eqv %3,%3,%2
    or %4,%1,%0\;eqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: eqv op eqv rtl xor inv 1 comp 0
-;; inner: orc op orc rtl ior inv 0 comp 1
+;; scalar orc -> eqv
 (define_insn "*fuse_orc_eqv"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (not:GPR (xor:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (not:GPR (xor:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    orc %3,%1,%0\;eqv %3,%3,%2
-   orc %0,%1,%0\;eqv %0,%0,%2
-   orc %1,%1,%0\;eqv %1,%1,%2
+   orc %3,%1,%0\;eqv %3,%3,%2
+   orc %3,%1,%0\;eqv %3,%3,%2
    orc %4,%1,%0\;eqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: eqv op eqv rtl xor inv 1 comp 0
-;; inner: xor op xor rtl xor inv 0 comp 0
+;; scalar xor -> eqv
 (define_insn "*fuse_xor_eqv"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (not:GPR (xor:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (not:GPR (xor:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    xor %3,%1,%0\;eqv %3,%3,%2
-   xor %0,%1,%0\;eqv %0,%0,%2
-   xor %1,%1,%0\;eqv %1,%1,%2
+   xor %3,%1,%0\;eqv %3,%3,%2
+   xor %3,%1,%0\;eqv %3,%3,%2
    xor %4,%1,%0\;eqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nand op nand rtl ior inv 0 comp 3
-;; inner: and op and rtl and inv 0 comp 0
+;; scalar and -> nand
 (define_insn "*fuse_and_nand"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (not:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (not:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    and %3,%1,%0\;nand %3,%3,%2
-   and %0,%1,%0\;nand %0,%0,%2
-   and %1,%1,%0\;nand %1,%1,%2
+   and %3,%1,%0\;nand %3,%3,%2
+   and %3,%1,%0\;nand %3,%3,%2
    and %4,%1,%0\;nand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nand op nand rtl ior inv 0 comp 3
-;; inner: andc op andc rtl and inv 0 comp 1
+;; scalar andc -> nand
 (define_insn "*fuse_andc_nand"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (not:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (not:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    andc %3,%1,%0\;nand %3,%3,%2
-   andc %0,%1,%0\;nand %0,%0,%2
-   andc %1,%1,%0\;nand %1,%1,%2
+   andc %3,%1,%0\;nand %3,%3,%2
+   andc %3,%1,%0\;nand %3,%3,%2
    andc %4,%1,%0\;nand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nand op nand rtl ior inv 0 comp 3
-;; inner: eqv op eqv rtl xor inv 1 comp 0
+;; scalar eqv -> nand
 (define_insn "*fuse_eqv_nand"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (not:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (not:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    eqv %3,%1,%0\;nand %3,%3,%2
-   eqv %0,%1,%0\;nand %0,%0,%2
-   eqv %1,%1,%0\;nand %1,%1,%2
+   eqv %3,%1,%0\;nand %3,%3,%2
+   eqv %3,%1,%0\;nand %3,%3,%2
    eqv %4,%1,%0\;nand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nand op nand rtl ior inv 0 comp 3
-;; inner: nand op nand rtl ior inv 0 comp 3
+;; scalar nand -> nand
 (define_insn "*fuse_nand_nand"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (not:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (not:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nand %3,%1,%0\;nand %3,%3,%2
-   nand %0,%1,%0\;nand %0,%0,%2
-   nand %1,%1,%0\;nand %1,%1,%2
+   nand %3,%1,%0\;nand %3,%3,%2
+   nand %3,%1,%0\;nand %3,%3,%2
    nand %4,%1,%0\;nand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nand op nand rtl ior inv 0 comp 3
-;; inner: nor op nor rtl and inv 0 comp 3
+;; scalar nor -> nand
 (define_insn "*fuse_nor_nand"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (not:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (not:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nor %3,%1,%0\;nand %3,%3,%2
-   nor %0,%1,%0\;nand %0,%0,%2
-   nor %1,%1,%0\;nand %1,%1,%2
+   nor %3,%1,%0\;nand %3,%3,%2
+   nor %3,%1,%0\;nand %3,%3,%2
    nor %4,%1,%0\;nand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nand op nand rtl ior inv 0 comp 3
-;; inner: or op or rtl ior inv 0 comp 0
+;; scalar or -> nand
 (define_insn "*fuse_or_nand"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (not:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (not:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    or %3,%1,%0\;nand %3,%3,%2
-   or %0,%1,%0\;nand %0,%0,%2
-   or %1,%1,%0\;nand %1,%1,%2
+   or %3,%1,%0\;nand %3,%3,%2
+   or %3,%1,%0\;nand %3,%3,%2
    or %4,%1,%0\;nand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nand op nand rtl ior inv 0 comp 3
-;; inner: orc op orc rtl ior inv 0 comp 1
+;; scalar orc -> nand
 (define_insn "*fuse_orc_nand"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (not:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (not:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    orc %3,%1,%0\;nand %3,%3,%2
-   orc %0,%1,%0\;nand %0,%0,%2
-   orc %1,%1,%0\;nand %1,%1,%2
+   orc %3,%1,%0\;nand %3,%3,%2
+   orc %3,%1,%0\;nand %3,%3,%2
    orc %4,%1,%0\;nand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nand op nand rtl ior inv 0 comp 3
-;; inner: xor op xor rtl xor inv 0 comp 0
+;; scalar xor -> nand
 (define_insn "*fuse_xor_nand"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    xor %3,%1,%0\;nand %3,%3,%2
-   xor %0,%1,%0\;nand %0,%0,%2
-   xor %1,%1,%0\;nand %1,%1,%2
+   xor %3,%1,%0\;nand %3,%3,%2
+   xor %3,%1,%0\;nand %3,%3,%2
    xor %4,%1,%0\;nand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nor op nor rtl and inv 0 comp 3
-;; inner: and op and rtl and inv 0 comp 0
+;; scalar and -> nor
 (define_insn "*fuse_and_nor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (not:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (not:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    and %3,%1,%0\;nor %3,%3,%2
-   and %0,%1,%0\;nor %0,%0,%2
-   and %1,%1,%0\;nor %1,%1,%2
+   and %3,%1,%0\;nor %3,%3,%2
+   and %3,%1,%0\;nor %3,%3,%2
    and %4,%1,%0\;nor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nor op nor rtl and inv 0 comp 3
-;; inner: andc op andc rtl and inv 0 comp 1
+;; scalar andc -> nor
 (define_insn "*fuse_andc_nor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (not:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (not:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    andc %3,%1,%0\;nor %3,%3,%2
-   andc %0,%1,%0\;nor %0,%0,%2
-   andc %1,%1,%0\;nor %1,%1,%2
+   andc %3,%1,%0\;nor %3,%3,%2
+   andc %3,%1,%0\;nor %3,%3,%2
    andc %4,%1,%0\;nor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nor op nor rtl and inv 0 comp 3
-;; inner: eqv op eqv rtl xor inv 1 comp 0
+;; scalar eqv -> nor
 (define_insn "*fuse_eqv_nor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (not:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (not:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    eqv %3,%1,%0\;nor %3,%3,%2
-   eqv %0,%1,%0\;nor %0,%0,%2
-   eqv %1,%1,%0\;nor %1,%1,%2
+   eqv %3,%1,%0\;nor %3,%3,%2
+   eqv %3,%1,%0\;nor %3,%3,%2
    eqv %4,%1,%0\;nor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nor op nor rtl and inv 0 comp 3
-;; inner: nand op nand rtl ior inv 0 comp 3
+;; scalar nand -> nor
 (define_insn "*fuse_nand_nor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (not:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (not:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nand %3,%1,%0\;nor %3,%3,%2
-   nand %0,%1,%0\;nor %0,%0,%2
-   nand %1,%1,%0\;nor %1,%1,%2
+   nand %3,%1,%0\;nor %3,%3,%2
+   nand %3,%1,%0\;nor %3,%3,%2
    nand %4,%1,%0\;nor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nor op nor rtl and inv 0 comp 3
-;; inner: nor op nor rtl and inv 0 comp 3
+;; scalar nor -> nor
 (define_insn "*fuse_nor_nor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (not:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (not:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nor %3,%1,%0\;nor %3,%3,%2
-   nor %0,%1,%0\;nor %0,%0,%2
-   nor %1,%1,%0\;nor %1,%1,%2
+   nor %3,%1,%0\;nor %3,%3,%2
+   nor %3,%1,%0\;nor %3,%3,%2
    nor %4,%1,%0\;nor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nor op nor rtl and inv 0 comp 3
-;; inner: or op or rtl ior inv 0 comp 0
+;; scalar or -> nor
 (define_insn "*fuse_or_nor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (not:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (not:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    or %3,%1,%0\;nor %3,%3,%2
-   or %0,%1,%0\;nor %0,%0,%2
-   or %1,%1,%0\;nor %1,%1,%2
+   or %3,%1,%0\;nor %3,%3,%2
+   or %3,%1,%0\;nor %3,%3,%2
    or %4,%1,%0\;nor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nor op nor rtl and inv 0 comp 3
-;; inner: orc op orc rtl ior inv 0 comp 1
+;; scalar orc -> nor
 (define_insn "*fuse_orc_nor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (not:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (not:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    orc %3,%1,%0\;nor %3,%3,%2
-   orc %0,%1,%0\;nor %0,%0,%2
-   orc %1,%1,%0\;nor %1,%1,%2
+   orc %3,%1,%0\;nor %3,%3,%2
+   orc %3,%1,%0\;nor %3,%3,%2
    orc %4,%1,%0\;nor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: nor op nor rtl and inv 0 comp 3
-;; inner: xor op xor rtl xor inv 0 comp 0
+;; scalar xor -> nor
 (define_insn "*fuse_xor_nor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (and:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (and:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    xor %3,%1,%0\;nor %3,%3,%2
-   xor %0,%1,%0\;nor %0,%0,%2
-   xor %1,%1,%0\;nor %1,%1,%2
+   xor %3,%1,%0\;nor %3,%3,%2
+   xor %3,%1,%0\;nor %3,%3,%2
    xor %4,%1,%0\;nor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: or op or rtl ior inv 0 comp 0
-;; inner: and op and rtl and inv 0 comp 0
+;; scalar and -> or
 (define_insn "*fuse_and_or"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    and %3,%1,%0\;or %3,%3,%2
-   and %0,%1,%0\;or %0,%0,%2
-   and %1,%1,%0\;or %1,%1,%2
+   and %3,%1,%0\;or %3,%3,%2
+   and %3,%1,%0\;or %3,%3,%2
    and %4,%1,%0\;or %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: or op or rtl ior inv 0 comp 0
-;; inner: andc op andc rtl and inv 0 comp 1
+;; scalar andc -> or
 (define_insn "*fuse_andc_or"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    andc %3,%1,%0\;or %3,%3,%2
-   andc %0,%1,%0\;or %0,%0,%2
-   andc %1,%1,%0\;or %1,%1,%2
+   andc %3,%1,%0\;or %3,%3,%2
+   andc %3,%1,%0\;or %3,%3,%2
    andc %4,%1,%0\;or %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: or op or rtl ior inv 0 comp 0
-;; inner: eqv op eqv rtl xor inv 1 comp 0
+;; scalar eqv -> or
 (define_insn "*fuse_eqv_or"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    eqv %3,%1,%0\;or %3,%3,%2
-   eqv %0,%1,%0\;or %0,%0,%2
-   eqv %1,%1,%0\;or %1,%1,%2
+   eqv %3,%1,%0\;or %3,%3,%2
+   eqv %3,%1,%0\;or %3,%3,%2
    eqv %4,%1,%0\;or %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: or op or rtl ior inv 0 comp 0
-;; inner: nand op nand rtl ior inv 0 comp 3
+;; scalar nand -> or
 (define_insn "*fuse_nand_or"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nand %3,%1,%0\;or %3,%3,%2
-   nand %0,%1,%0\;or %0,%0,%2
-   nand %1,%1,%0\;or %1,%1,%2
+   nand %3,%1,%0\;or %3,%3,%2
+   nand %3,%1,%0\;or %3,%3,%2
    nand %4,%1,%0\;or %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: or op or rtl ior inv 0 comp 0
-;; inner: nor op nor rtl and inv 0 comp 3
+;; scalar nor -> or
 (define_insn "*fuse_nor_or"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nor %3,%1,%0\;or %3,%3,%2
-   nor %0,%1,%0\;or %0,%0,%2
-   nor %1,%1,%0\;or %1,%1,%2
+   nor %3,%1,%0\;or %3,%3,%2
+   nor %3,%1,%0\;or %3,%3,%2
    nor %4,%1,%0\;or %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: or op or rtl ior inv 0 comp 0
-;; inner: or op or rtl ior inv 0 comp 0
+;; scalar or -> or
 (define_insn "*fuse_or_or"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "%r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "%r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    or %3,%1,%0\;or %3,%3,%2
-   or %0,%1,%0\;or %0,%0,%2
-   or %1,%1,%0\;or %1,%1,%2
+   or %3,%1,%0\;or %3,%3,%2
+   or %3,%1,%0\;or %3,%3,%2
    or %4,%1,%0\;or %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: or op or rtl ior inv 0 comp 0
-;; inner: orc op orc rtl ior inv 0 comp 1
+;; scalar orc -> or
 (define_insn "*fuse_orc_or"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    orc %3,%1,%0\;or %3,%3,%2
-   orc %0,%1,%0\;or %0,%0,%2
-   orc %1,%1,%0\;or %1,%1,%2
+   orc %3,%1,%0\;or %3,%3,%2
+   orc %3,%1,%0\;or %3,%3,%2
    orc %4,%1,%0\;or %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: or op or rtl ior inv 0 comp 0
-;; inner: xor op xor rtl xor inv 0 comp 0
+;; scalar xor -> or
 (define_insn "*fuse_xor_or"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    xor %3,%1,%0\;or %3,%3,%2
-   xor %0,%1,%0\;or %0,%0,%2
-   xor %1,%1,%0\;or %1,%1,%2
+   xor %3,%1,%0\;or %3,%3,%2
+   xor %3,%1,%0\;or %3,%3,%2
    xor %4,%1,%0\;or %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: orc op orc rtl ior inv 0 comp 1
-;; inner: and op and rtl and inv 0 comp 0
+;; scalar and -> orc
 (define_insn "*fuse_and_orc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    and %3,%1,%0\;orc %3,%3,%2
-   and %0,%1,%0\;orc %0,%0,%2
-   and %1,%1,%0\;orc %1,%1,%2
+   and %3,%1,%0\;orc %3,%3,%2
+   and %3,%1,%0\;orc %3,%3,%2
    and %4,%1,%0\;orc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: orc op orc rtl ior inv 0 comp 1
-;; inner: andc op andc rtl and inv 0 comp 1
+;; scalar andc -> orc
 (define_insn "*fuse_andc_orc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    andc %3,%1,%0\;orc %3,%3,%2
-   andc %0,%1,%0\;orc %0,%0,%2
-   andc %1,%1,%0\;orc %1,%1,%2
+   andc %3,%1,%0\;orc %3,%3,%2
+   andc %3,%1,%0\;orc %3,%3,%2
    andc %4,%1,%0\;orc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: orc op orc rtl ior inv 0 comp 1
-;; inner: eqv op eqv rtl xor inv 1 comp 0
+;; scalar eqv -> orc
 (define_insn "*fuse_eqv_orc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    eqv %3,%1,%0\;orc %3,%3,%2
-   eqv %0,%1,%0\;orc %0,%0,%2
-   eqv %1,%1,%0\;orc %1,%1,%2
+   eqv %3,%1,%0\;orc %3,%3,%2
+   eqv %3,%1,%0\;orc %3,%3,%2
    eqv %4,%1,%0\;orc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: orc op orc rtl ior inv 0 comp 1
-;; inner: nand op nand rtl ior inv 0 comp 3
+;; scalar nand -> orc
 (define_insn "*fuse_nand_orc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nand %3,%1,%0\;orc %3,%3,%2
-   nand %0,%1,%0\;orc %0,%0,%2
-   nand %1,%1,%0\;orc %1,%1,%2
+   nand %3,%1,%0\;orc %3,%3,%2
+   nand %3,%1,%0\;orc %3,%3,%2
    nand %4,%1,%0\;orc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: orc op orc rtl ior inv 0 comp 1
-;; inner: nor op nor rtl and inv 0 comp 3
+;; scalar nor -> orc
 (define_insn "*fuse_nor_orc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nor %3,%1,%0\;orc %3,%3,%2
-   nor %0,%1,%0\;orc %0,%0,%2
-   nor %1,%1,%0\;orc %1,%1,%2
+   nor %3,%1,%0\;orc %3,%3,%2
+   nor %3,%1,%0\;orc %3,%3,%2
    nor %4,%1,%0\;orc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: orc op orc rtl ior inv 0 comp 1
-;; inner: or op or rtl ior inv 0 comp 0
+;; scalar or -> orc
 (define_insn "*fuse_or_orc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    or %3,%1,%0\;orc %3,%3,%2
-   or %0,%1,%0\;orc %0,%0,%2
-   or %1,%1,%0\;orc %1,%1,%2
+   or %3,%1,%0\;orc %3,%3,%2
+   or %3,%1,%0\;orc %3,%3,%2
    or %4,%1,%0\;orc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: orc op orc rtl ior inv 0 comp 1
-;; inner: orc op orc rtl ior inv 0 comp 1
+;; scalar orc -> orc
 (define_insn "*fuse_orc_orc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    orc %3,%1,%0\;orc %3,%3,%2
-   orc %0,%1,%0\;orc %0,%0,%2
-   orc %1,%1,%0\;orc %1,%1,%2
+   orc %3,%1,%0\;orc %3,%3,%2
+   orc %3,%1,%0\;orc %3,%3,%2
    orc %4,%1,%0\;orc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: orc op orc rtl ior inv 0 comp 1
-;; inner: xor op xor rtl xor inv 0 comp 0
+;; scalar xor -> orc
 (define_insn "*fuse_xor_orc"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (ior:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (ior:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (not:GPR (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r"))))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    xor %3,%1,%0\;orc %3,%3,%2
-   xor %0,%1,%0\;orc %0,%0,%2
-   xor %1,%1,%0\;orc %1,%1,%2
+   xor %3,%1,%0\;orc %3,%3,%2
+   xor %3,%1,%0\;orc %3,%3,%2
    xor %4,%1,%0\;orc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: xor op xor rtl xor inv 0 comp 0
-;; inner: and op and rtl and inv 0 comp 0
+;; scalar and -> xor
 (define_insn "*fuse_and_xor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (xor:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (xor:GPR (and:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    and %3,%1,%0\;xor %3,%3,%2
-   and %0,%1,%0\;xor %0,%0,%2
-   and %1,%1,%0\;xor %1,%1,%2
+   and %3,%1,%0\;xor %3,%3,%2
+   and %3,%1,%0\;xor %3,%3,%2
    and %4,%1,%0\;xor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: xor op xor rtl xor inv 0 comp 0
-;; inner: andc op andc rtl and inv 0 comp 1
+;; scalar andc -> xor
 (define_insn "*fuse_andc_xor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (xor:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (xor:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    andc %3,%1,%0\;xor %3,%3,%2
-   andc %0,%1,%0\;xor %0,%0,%2
-   andc %1,%1,%0\;xor %1,%1,%2
+   andc %3,%1,%0\;xor %3,%3,%2
+   andc %3,%1,%0\;xor %3,%3,%2
    andc %4,%1,%0\;xor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: xor op xor rtl xor inv 0 comp 0
-;; inner: eqv op eqv rtl xor inv 1 comp 0
+;; scalar eqv -> xor
 (define_insn "*fuse_eqv_xor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (xor:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (xor:GPR (not:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    eqv %3,%1,%0\;xor %3,%3,%2
-   eqv %0,%1,%0\;xor %0,%0,%2
-   eqv %1,%1,%0\;xor %1,%1,%2
+   eqv %3,%1,%0\;xor %3,%3,%2
+   eqv %3,%1,%0\;xor %3,%3,%2
    eqv %4,%1,%0\;xor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: xor op xor rtl xor inv 0 comp 0
-;; inner: nand op nand rtl ior inv 0 comp 3
+;; scalar nand -> xor
 (define_insn "*fuse_nand_xor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (xor:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (xor:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nand %3,%1,%0\;xor %3,%3,%2
-   nand %0,%1,%0\;xor %0,%0,%2
-   nand %1,%1,%0\;xor %1,%1,%2
+   nand %3,%1,%0\;xor %3,%3,%2
+   nand %3,%1,%0\;xor %3,%3,%2
    nand %4,%1,%0\;xor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: xor op xor rtl xor inv 0 comp 0
-;; inner: nor op nor rtl and inv 0 comp 3
+;; scalar nor -> xor
 (define_insn "*fuse_nor_xor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (xor:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (xor:GPR (and:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (not:GPR (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    nor %3,%1,%0\;xor %3,%3,%2
-   nor %0,%1,%0\;xor %0,%0,%2
-   nor %1,%1,%0\;xor %1,%1,%2
+   nor %3,%1,%0\;xor %3,%3,%2
+   nor %3,%1,%0\;xor %3,%3,%2
    nor %4,%1,%0\;xor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: xor op xor rtl xor inv 0 comp 0
-;; inner: or op or rtl ior inv 0 comp 0
+;; scalar or -> xor
 (define_insn "*fuse_or_xor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (xor:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (xor:GPR (ior:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    or %3,%1,%0\;xor %3,%3,%2
-   or %0,%1,%0\;xor %0,%0,%2
-   or %1,%1,%0\;xor %1,%1,%2
+   or %3,%1,%0\;xor %3,%3,%2
+   or %3,%1,%0\;xor %3,%3,%2
    or %4,%1,%0\;xor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: xor op xor rtl xor inv 0 comp 0
-;; inner: orc op orc rtl ior inv 0 comp 1
+;; scalar orc -> xor
 (define_insn "*fuse_orc_xor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (xor:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (xor:GPR (ior:GPR (not:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r")) 
+                          (match_operand:GPR 1 "gpc_reg_operand" "r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    orc %3,%1,%0\;xor %3,%3,%2
-   orc %0,%1,%0\;xor %0,%0,%2
-   orc %1,%1,%0\;xor %1,%1,%2
+   orc %3,%1,%0\;xor %3,%3,%2
+   orc %3,%1,%0\;xor %3,%3,%2
    orc %4,%1,%0\;xor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: scalar outer: xor op xor rtl xor inv 0 comp 0
-;; inner: xor op xor rtl xor inv 0 comp 0
+;; scalar xor -> xor
 (define_insn "*fuse_xor_xor"
-  [(set (match_operand:GPR 3 "gpc_reg_operand" "=&r,0,1,r")
-        (xor:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") (match_operand:GPR 1 "gpc_reg_operand" "%r,r,r,r")) (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
-   (clobber (match_scratch:GPR 4 "=X,X,X,r"))]
+  [(set (match_operand:GPR 3 "gpc_reg_operand" "=0,1,&r,r")
+        (xor:GPR (xor:GPR (match_operand:GPR 0 "gpc_reg_operand" "r,r,r,r") 
+                          (match_operand:GPR 1 "gpc_reg_operand" "%r,r,r,r"))
+                 (match_operand:GPR 2 "gpc_reg_operand" "r,r,r,r")))
+   (clobber (match_scratch:GPR 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    xor %3,%1,%0\;xor %3,%3,%2
-   xor %0,%1,%0\;xor %0,%0,%2
-   xor %1,%1,%0\;xor %1,%1,%2
+   xor %3,%1,%0\;xor %3,%3,%2
+   xor %3,%1,%0\;xor %3,%3,%2
    xor %4,%1,%0\;xor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: and op vand rtl and inv 0 comp 0
-;; inner: and op vand rtl and inv 0 comp 0
+;; vector vand -> vand
 (define_insn "*fuse_vand_vand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "%v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "%v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vand %3,%1,%0\;vand %3,%3,%2
-   vand %0,%1,%0\;vand %0,%0,%2
-   vand %1,%1,%0\;vand %1,%1,%2
+   vand %3,%1,%0\;vand %3,%3,%2
+   vand %3,%1,%0\;vand %3,%3,%2
    vand %4,%1,%0\;vand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: and op vand rtl and inv 0 comp 0
-;; inner: andc op vandc rtl and inv 0 comp 1
+;; vector vandc -> vand
 (define_insn "*fuse_vandc_vand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vandc %3,%1,%0\;vand %3,%3,%2
-   vandc %0,%1,%0\;vand %0,%0,%2
-   vandc %1,%1,%0\;vand %1,%1,%2
+   vandc %3,%1,%0\;vand %3,%3,%2
+   vandc %3,%1,%0\;vand %3,%3,%2
    vandc %4,%1,%0\;vand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: and op vand rtl and inv 0 comp 0
-;; inner: eqv op veqv rtl xor inv 1 comp 0
+;; vector veqv -> vand
 (define_insn "*fuse_veqv_vand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    veqv %3,%1,%0\;vand %3,%3,%2
-   veqv %0,%1,%0\;vand %0,%0,%2
-   veqv %1,%1,%0\;vand %1,%1,%2
+   veqv %3,%1,%0\;vand %3,%3,%2
+   veqv %3,%1,%0\;vand %3,%3,%2
    veqv %4,%1,%0\;vand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: and op vand rtl and inv 0 comp 0
-;; inner: nand op vnand rtl ior inv 0 comp 3
+;; vector vnand -> vand
 (define_insn "*fuse_vnand_vand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnand %3,%1,%0\;vand %3,%3,%2
-   vnand %0,%1,%0\;vand %0,%0,%2
-   vnand %1,%1,%0\;vand %1,%1,%2
+   vnand %3,%1,%0\;vand %3,%3,%2
+   vnand %3,%1,%0\;vand %3,%3,%2
    vnand %4,%1,%0\;vand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: and op vand rtl and inv 0 comp 0
-;; inner: nor op vnor rtl and inv 0 comp 3
+;; vector vnor -> vand
 (define_insn "*fuse_vnor_vand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnor %3,%1,%0\;vand %3,%3,%2
-   vnor %0,%1,%0\;vand %0,%0,%2
-   vnor %1,%1,%0\;vand %1,%1,%2
+   vnor %3,%1,%0\;vand %3,%3,%2
+   vnor %3,%1,%0\;vand %3,%3,%2
    vnor %4,%1,%0\;vand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: and op vand rtl and inv 0 comp 0
-;; inner: or op vor rtl ior inv 0 comp 0
+;; vector vor -> vand
 (define_insn "*fuse_vor_vand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vor %3,%1,%0\;vand %3,%3,%2
-   vor %0,%1,%0\;vand %0,%0,%2
-   vor %1,%1,%0\;vand %1,%1,%2
+   vor %3,%1,%0\;vand %3,%3,%2
+   vor %3,%1,%0\;vand %3,%3,%2
    vor %4,%1,%0\;vand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: and op vand rtl and inv 0 comp 0
-;; inner: orc op vorc rtl ior inv 0 comp 1
+;; vector vorc -> vand
 (define_insn "*fuse_vorc_vand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vorc %3,%1,%0\;vand %3,%3,%2
-   vorc %0,%1,%0\;vand %0,%0,%2
-   vorc %1,%1,%0\;vand %1,%1,%2
+   vorc %3,%1,%0\;vand %3,%3,%2
+   vorc %3,%1,%0\;vand %3,%3,%2
    vorc %4,%1,%0\;vand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: and op vand rtl and inv 0 comp 0
-;; inner: xor op vxor rtl xor inv 0 comp 0
+;; vector vxor -> vand
 (define_insn "*fuse_vxor_vand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vxor %3,%1,%0\;vand %3,%3,%2
-   vxor %0,%1,%0\;vand %0,%0,%2
-   vxor %1,%1,%0\;vand %1,%1,%2
+   vxor %3,%1,%0\;vand %3,%3,%2
+   vxor %3,%1,%0\;vand %3,%3,%2
    vxor %4,%1,%0\;vand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: andc op vandc rtl and inv 0 comp 1
-;; inner: and op vand rtl and inv 0 comp 0
+;; vector vand -> vandc
 (define_insn "*fuse_vand_vandc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vand %3,%1,%0\;vandc %3,%3,%2
-   vand %0,%1,%0\;vandc %0,%0,%2
-   vand %1,%1,%0\;vandc %1,%1,%2
+   vand %3,%1,%0\;vandc %3,%3,%2
+   vand %3,%1,%0\;vandc %3,%3,%2
    vand %4,%1,%0\;vandc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: andc op vandc rtl and inv 0 comp 1
-;; inner: andc op vandc rtl and inv 0 comp 1
+;; vector vandc -> vandc
 (define_insn "*fuse_vandc_vandc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vandc %3,%1,%0\;vandc %3,%3,%2
-   vandc %0,%1,%0\;vandc %0,%0,%2
-   vandc %1,%1,%0\;vandc %1,%1,%2
+   vandc %3,%1,%0\;vandc %3,%3,%2
+   vandc %3,%1,%0\;vandc %3,%3,%2
    vandc %4,%1,%0\;vandc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: andc op vandc rtl and inv 0 comp 1
-;; inner: eqv op veqv rtl xor inv 1 comp 0
+;; vector veqv -> vandc
 (define_insn "*fuse_veqv_vandc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    veqv %3,%1,%0\;vandc %3,%3,%2
-   veqv %0,%1,%0\;vandc %0,%0,%2
-   veqv %1,%1,%0\;vandc %1,%1,%2
+   veqv %3,%1,%0\;vandc %3,%3,%2
+   veqv %3,%1,%0\;vandc %3,%3,%2
    veqv %4,%1,%0\;vandc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: andc op vandc rtl and inv 0 comp 1
-;; inner: nand op vnand rtl ior inv 0 comp 3
+;; vector vnand -> vandc
 (define_insn "*fuse_vnand_vandc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnand %3,%1,%0\;vandc %3,%3,%2
-   vnand %0,%1,%0\;vandc %0,%0,%2
-   vnand %1,%1,%0\;vandc %1,%1,%2
+   vnand %3,%1,%0\;vandc %3,%3,%2
+   vnand %3,%1,%0\;vandc %3,%3,%2
    vnand %4,%1,%0\;vandc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: andc op vandc rtl and inv 0 comp 1
-;; inner: nor op vnor rtl and inv 0 comp 3
+;; vector vnor -> vandc
 (define_insn "*fuse_vnor_vandc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnor %3,%1,%0\;vandc %3,%3,%2
-   vnor %0,%1,%0\;vandc %0,%0,%2
-   vnor %1,%1,%0\;vandc %1,%1,%2
+   vnor %3,%1,%0\;vandc %3,%3,%2
+   vnor %3,%1,%0\;vandc %3,%3,%2
    vnor %4,%1,%0\;vandc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: andc op vandc rtl and inv 0 comp 1
-;; inner: or op vor rtl ior inv 0 comp 0
+;; vector vor -> vandc
 (define_insn "*fuse_vor_vandc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vor %3,%1,%0\;vandc %3,%3,%2
-   vor %0,%1,%0\;vandc %0,%0,%2
-   vor %1,%1,%0\;vandc %1,%1,%2
+   vor %3,%1,%0\;vandc %3,%3,%2
+   vor %3,%1,%0\;vandc %3,%3,%2
    vor %4,%1,%0\;vandc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: andc op vandc rtl and inv 0 comp 1
-;; inner: orc op vorc rtl ior inv 0 comp 1
+;; vector vorc -> vandc
 (define_insn "*fuse_vorc_vandc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vorc %3,%1,%0\;vandc %3,%3,%2
-   vorc %0,%1,%0\;vandc %0,%0,%2
-   vorc %1,%1,%0\;vandc %1,%1,%2
+   vorc %3,%1,%0\;vandc %3,%3,%2
+   vorc %3,%1,%0\;vandc %3,%3,%2
    vorc %4,%1,%0\;vandc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: andc op vandc rtl and inv 0 comp 1
-;; inner: xor op vxor rtl xor inv 0 comp 0
+;; vector vxor -> vandc
 (define_insn "*fuse_vxor_vandc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vxor %3,%1,%0\;vandc %3,%3,%2
-   vxor %0,%1,%0\;vandc %0,%0,%2
-   vxor %1,%1,%0\;vandc %1,%1,%2
+   vxor %3,%1,%0\;vandc %3,%3,%2
+   vxor %3,%1,%0\;vandc %3,%3,%2
    vxor %4,%1,%0\;vandc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: eqv op veqv rtl xor inv 1 comp 0
-;; inner: and op vand rtl and inv 0 comp 0
+;; vector vand -> veqv
 (define_insn "*fuse_vand_veqv"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (not:VM (xor:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (not:VM (xor:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vand %3,%1,%0\;veqv %3,%3,%2
-   vand %0,%1,%0\;veqv %0,%0,%2
-   vand %1,%1,%0\;veqv %1,%1,%2
+   vand %3,%1,%0\;veqv %3,%3,%2
+   vand %3,%1,%0\;veqv %3,%3,%2
    vand %4,%1,%0\;veqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: eqv op veqv rtl xor inv 1 comp 0
-;; inner: andc op vandc rtl and inv 0 comp 1
+;; vector vandc -> veqv
 (define_insn "*fuse_vandc_veqv"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (not:VM (xor:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (not:VM (xor:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vandc %3,%1,%0\;veqv %3,%3,%2
-   vandc %0,%1,%0\;veqv %0,%0,%2
-   vandc %1,%1,%0\;veqv %1,%1,%2
+   vandc %3,%1,%0\;veqv %3,%3,%2
+   vandc %3,%1,%0\;veqv %3,%3,%2
    vandc %4,%1,%0\;veqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: eqv op veqv rtl xor inv 1 comp 0
-;; inner: eqv op veqv rtl xor inv 1 comp 0
+;; vector veqv -> veqv
 (define_insn "*fuse_veqv_veqv"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (not:VM (xor:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "%v,v,v,v"))) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (not:VM (xor:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "%v,v,v,v")))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    veqv %3,%1,%0\;veqv %3,%3,%2
-   veqv %0,%1,%0\;veqv %0,%0,%2
-   veqv %1,%1,%0\;veqv %1,%1,%2
+   veqv %3,%1,%0\;veqv %3,%3,%2
+   veqv %3,%1,%0\;veqv %3,%3,%2
    veqv %4,%1,%0\;veqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: eqv op veqv rtl xor inv 1 comp 0
-;; inner: nand op vnand rtl ior inv 0 comp 3
+;; vector vnand -> veqv
 (define_insn "*fuse_vnand_veqv"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (not:VM (xor:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (not:VM (xor:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnand %3,%1,%0\;veqv %3,%3,%2
-   vnand %0,%1,%0\;veqv %0,%0,%2
-   vnand %1,%1,%0\;veqv %1,%1,%2
+   vnand %3,%1,%0\;veqv %3,%3,%2
+   vnand %3,%1,%0\;veqv %3,%3,%2
    vnand %4,%1,%0\;veqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: eqv op veqv rtl xor inv 1 comp 0
-;; inner: nor op vnor rtl and inv 0 comp 3
+;; vector vnor -> veqv
 (define_insn "*fuse_vnor_veqv"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (not:VM (xor:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (not:VM (xor:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnor %3,%1,%0\;veqv %3,%3,%2
-   vnor %0,%1,%0\;veqv %0,%0,%2
-   vnor %1,%1,%0\;veqv %1,%1,%2
+   vnor %3,%1,%0\;veqv %3,%3,%2
+   vnor %3,%1,%0\;veqv %3,%3,%2
    vnor %4,%1,%0\;veqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: eqv op veqv rtl xor inv 1 comp 0
-;; inner: or op vor rtl ior inv 0 comp 0
+;; vector vor -> veqv
 (define_insn "*fuse_vor_veqv"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (not:VM (xor:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (not:VM (xor:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vor %3,%1,%0\;veqv %3,%3,%2
-   vor %0,%1,%0\;veqv %0,%0,%2
-   vor %1,%1,%0\;veqv %1,%1,%2
+   vor %3,%1,%0\;veqv %3,%3,%2
+   vor %3,%1,%0\;veqv %3,%3,%2
    vor %4,%1,%0\;veqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: eqv op veqv rtl xor inv 1 comp 0
-;; inner: orc op vorc rtl ior inv 0 comp 1
+;; vector vorc -> veqv
 (define_insn "*fuse_vorc_veqv"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (not:VM (xor:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (not:VM (xor:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vorc %3,%1,%0\;veqv %3,%3,%2
-   vorc %0,%1,%0\;veqv %0,%0,%2
-   vorc %1,%1,%0\;veqv %1,%1,%2
+   vorc %3,%1,%0\;veqv %3,%3,%2
+   vorc %3,%1,%0\;veqv %3,%3,%2
    vorc %4,%1,%0\;veqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: eqv op veqv rtl xor inv 1 comp 0
-;; inner: xor op vxor rtl xor inv 0 comp 0
+;; vector vxor -> veqv
 (define_insn "*fuse_vxor_veqv"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (not:VM (xor:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (not:VM (xor:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vxor %3,%1,%0\;veqv %3,%3,%2
-   vxor %0,%1,%0\;veqv %0,%0,%2
-   vxor %1,%1,%0\;veqv %1,%1,%2
+   vxor %3,%1,%0\;veqv %3,%3,%2
+   vxor %3,%1,%0\;veqv %3,%3,%2
    vxor %4,%1,%0\;veqv %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nand op vnand rtl ior inv 0 comp 3
-;; inner: and op vand rtl and inv 0 comp 0
+;; vector vand -> vnand
 (define_insn "*fuse_vand_vnand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (not:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (not:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vand %3,%1,%0\;vnand %3,%3,%2
-   vand %0,%1,%0\;vnand %0,%0,%2
-   vand %1,%1,%0\;vnand %1,%1,%2
+   vand %3,%1,%0\;vnand %3,%3,%2
+   vand %3,%1,%0\;vnand %3,%3,%2
    vand %4,%1,%0\;vnand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nand op vnand rtl ior inv 0 comp 3
-;; inner: andc op vandc rtl and inv 0 comp 1
+;; vector vandc -> vnand
 (define_insn "*fuse_vandc_vnand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (not:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (not:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vandc %3,%1,%0\;vnand %3,%3,%2
-   vandc %0,%1,%0\;vnand %0,%0,%2
-   vandc %1,%1,%0\;vnand %1,%1,%2
+   vandc %3,%1,%0\;vnand %3,%3,%2
+   vandc %3,%1,%0\;vnand %3,%3,%2
    vandc %4,%1,%0\;vnand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nand op vnand rtl ior inv 0 comp 3
-;; inner: eqv op veqv rtl xor inv 1 comp 0
+;; vector veqv -> vnand
 (define_insn "*fuse_veqv_vnand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (not:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (not:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    veqv %3,%1,%0\;vnand %3,%3,%2
-   veqv %0,%1,%0\;vnand %0,%0,%2
-   veqv %1,%1,%0\;vnand %1,%1,%2
+   veqv %3,%1,%0\;vnand %3,%3,%2
+   veqv %3,%1,%0\;vnand %3,%3,%2
    veqv %4,%1,%0\;vnand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nand op vnand rtl ior inv 0 comp 3
-;; inner: nand op vnand rtl ior inv 0 comp 3
+;; vector vnand -> vnand
 (define_insn "*fuse_vnand_vnand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (not:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (not:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnand %3,%1,%0\;vnand %3,%3,%2
-   vnand %0,%1,%0\;vnand %0,%0,%2
-   vnand %1,%1,%0\;vnand %1,%1,%2
+   vnand %3,%1,%0\;vnand %3,%3,%2
+   vnand %3,%1,%0\;vnand %3,%3,%2
    vnand %4,%1,%0\;vnand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nand op vnand rtl ior inv 0 comp 3
-;; inner: nor op vnor rtl and inv 0 comp 3
+;; vector vnor -> vnand
 (define_insn "*fuse_vnor_vnand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (not:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (not:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnor %3,%1,%0\;vnand %3,%3,%2
-   vnor %0,%1,%0\;vnand %0,%0,%2
-   vnor %1,%1,%0\;vnand %1,%1,%2
+   vnor %3,%1,%0\;vnand %3,%3,%2
+   vnor %3,%1,%0\;vnand %3,%3,%2
    vnor %4,%1,%0\;vnand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nand op vnand rtl ior inv 0 comp 3
-;; inner: or op vor rtl ior inv 0 comp 0
+;; vector vor -> vnand
 (define_insn "*fuse_vor_vnand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (not:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (not:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vor %3,%1,%0\;vnand %3,%3,%2
-   vor %0,%1,%0\;vnand %0,%0,%2
-   vor %1,%1,%0\;vnand %1,%1,%2
+   vor %3,%1,%0\;vnand %3,%3,%2
+   vor %3,%1,%0\;vnand %3,%3,%2
    vor %4,%1,%0\;vnand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nand op vnand rtl ior inv 0 comp 3
-;; inner: orc op vorc rtl ior inv 0 comp 1
+;; vector vorc -> vnand
 (define_insn "*fuse_vorc_vnand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (not:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (not:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vorc %3,%1,%0\;vnand %3,%3,%2
-   vorc %0,%1,%0\;vnand %0,%0,%2
-   vorc %1,%1,%0\;vnand %1,%1,%2
+   vorc %3,%1,%0\;vnand %3,%3,%2
+   vorc %3,%1,%0\;vnand %3,%3,%2
    vorc %4,%1,%0\;vnand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nand op vnand rtl ior inv 0 comp 3
-;; inner: xor op vxor rtl xor inv 0 comp 0
+;; vector vxor -> vnand
 (define_insn "*fuse_vxor_vnand"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vxor %3,%1,%0\;vnand %3,%3,%2
-   vxor %0,%1,%0\;vnand %0,%0,%2
-   vxor %1,%1,%0\;vnand %1,%1,%2
+   vxor %3,%1,%0\;vnand %3,%3,%2
+   vxor %3,%1,%0\;vnand %3,%3,%2
    vxor %4,%1,%0\;vnand %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nor op vnor rtl and inv 0 comp 3
-;; inner: and op vand rtl and inv 0 comp 0
+;; vector vand -> vnor
 (define_insn "*fuse_vand_vnor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (not:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (not:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vand %3,%1,%0\;vnor %3,%3,%2
-   vand %0,%1,%0\;vnor %0,%0,%2
-   vand %1,%1,%0\;vnor %1,%1,%2
+   vand %3,%1,%0\;vnor %3,%3,%2
+   vand %3,%1,%0\;vnor %3,%3,%2
    vand %4,%1,%0\;vnor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nor op vnor rtl and inv 0 comp 3
-;; inner: andc op vandc rtl and inv 0 comp 1
+;; vector vandc -> vnor
 (define_insn "*fuse_vandc_vnor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (not:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (not:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vandc %3,%1,%0\;vnor %3,%3,%2
-   vandc %0,%1,%0\;vnor %0,%0,%2
-   vandc %1,%1,%0\;vnor %1,%1,%2
+   vandc %3,%1,%0\;vnor %3,%3,%2
+   vandc %3,%1,%0\;vnor %3,%3,%2
    vandc %4,%1,%0\;vnor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nor op vnor rtl and inv 0 comp 3
-;; inner: eqv op veqv rtl xor inv 1 comp 0
+;; vector veqv -> vnor
 (define_insn "*fuse_veqv_vnor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (not:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (not:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    veqv %3,%1,%0\;vnor %3,%3,%2
-   veqv %0,%1,%0\;vnor %0,%0,%2
-   veqv %1,%1,%0\;vnor %1,%1,%2
+   veqv %3,%1,%0\;vnor %3,%3,%2
+   veqv %3,%1,%0\;vnor %3,%3,%2
    veqv %4,%1,%0\;vnor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nor op vnor rtl and inv 0 comp 3
-;; inner: nand op vnand rtl ior inv 0 comp 3
+;; vector vnand -> vnor
 (define_insn "*fuse_vnand_vnor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (not:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (not:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnand %3,%1,%0\;vnor %3,%3,%2
-   vnand %0,%1,%0\;vnor %0,%0,%2
-   vnand %1,%1,%0\;vnor %1,%1,%2
+   vnand %3,%1,%0\;vnor %3,%3,%2
+   vnand %3,%1,%0\;vnor %3,%3,%2
    vnand %4,%1,%0\;vnor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nor op vnor rtl and inv 0 comp 3
-;; inner: nor op vnor rtl and inv 0 comp 3
+;; vector vnor -> vnor
 (define_insn "*fuse_vnor_vnor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (not:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (not:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnor %3,%1,%0\;vnor %3,%3,%2
-   vnor %0,%1,%0\;vnor %0,%0,%2
-   vnor %1,%1,%0\;vnor %1,%1,%2
+   vnor %3,%1,%0\;vnor %3,%3,%2
+   vnor %3,%1,%0\;vnor %3,%3,%2
    vnor %4,%1,%0\;vnor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nor op vnor rtl and inv 0 comp 3
-;; inner: or op vor rtl ior inv 0 comp 0
+;; vector vor -> vnor
 (define_insn "*fuse_vor_vnor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (not:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (not:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vor %3,%1,%0\;vnor %3,%3,%2
-   vor %0,%1,%0\;vnor %0,%0,%2
-   vor %1,%1,%0\;vnor %1,%1,%2
+   vor %3,%1,%0\;vnor %3,%3,%2
+   vor %3,%1,%0\;vnor %3,%3,%2
    vor %4,%1,%0\;vnor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nor op vnor rtl and inv 0 comp 3
-;; inner: orc op vorc rtl ior inv 0 comp 1
+;; vector vorc -> vnor
 (define_insn "*fuse_vorc_vnor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (not:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (not:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vorc %3,%1,%0\;vnor %3,%3,%2
-   vorc %0,%1,%0\;vnor %0,%0,%2
-   vorc %1,%1,%0\;vnor %1,%1,%2
+   vorc %3,%1,%0\;vnor %3,%3,%2
+   vorc %3,%1,%0\;vnor %3,%3,%2
    vorc %4,%1,%0\;vnor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: nor op vnor rtl and inv 0 comp 3
-;; inner: xor op vxor rtl xor inv 0 comp 0
+;; vector vxor -> vnor
 (define_insn "*fuse_vxor_vnor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (and:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (and:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vxor %3,%1,%0\;vnor %3,%3,%2
-   vxor %0,%1,%0\;vnor %0,%0,%2
-   vxor %1,%1,%0\;vnor %1,%1,%2
+   vxor %3,%1,%0\;vnor %3,%3,%2
+   vxor %3,%1,%0\;vnor %3,%3,%2
    vxor %4,%1,%0\;vnor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: or op vor rtl ior inv 0 comp 0
-;; inner: and op vand rtl and inv 0 comp 0
+;; vector vand -> vor
 (define_insn "*fuse_vand_vor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vand %3,%1,%0\;vor %3,%3,%2
-   vand %0,%1,%0\;vor %0,%0,%2
-   vand %1,%1,%0\;vor %1,%1,%2
+   vand %3,%1,%0\;vor %3,%3,%2
+   vand %3,%1,%0\;vor %3,%3,%2
    vand %4,%1,%0\;vor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: or op vor rtl ior inv 0 comp 0
-;; inner: andc op vandc rtl and inv 0 comp 1
+;; vector vandc -> vor
 (define_insn "*fuse_vandc_vor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vandc %3,%1,%0\;vor %3,%3,%2
-   vandc %0,%1,%0\;vor %0,%0,%2
-   vandc %1,%1,%0\;vor %1,%1,%2
+   vandc %3,%1,%0\;vor %3,%3,%2
+   vandc %3,%1,%0\;vor %3,%3,%2
    vandc %4,%1,%0\;vor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: or op vor rtl ior inv 0 comp 0
-;; inner: eqv op veqv rtl xor inv 1 comp 0
+;; vector veqv -> vor
 (define_insn "*fuse_veqv_vor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    veqv %3,%1,%0\;vor %3,%3,%2
-   veqv %0,%1,%0\;vor %0,%0,%2
-   veqv %1,%1,%0\;vor %1,%1,%2
+   veqv %3,%1,%0\;vor %3,%3,%2
+   veqv %3,%1,%0\;vor %3,%3,%2
    veqv %4,%1,%0\;vor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: or op vor rtl ior inv 0 comp 0
-;; inner: nand op vnand rtl ior inv 0 comp 3
+;; vector vnand -> vor
 (define_insn "*fuse_vnand_vor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnand %3,%1,%0\;vor %3,%3,%2
-   vnand %0,%1,%0\;vor %0,%0,%2
-   vnand %1,%1,%0\;vor %1,%1,%2
+   vnand %3,%1,%0\;vor %3,%3,%2
+   vnand %3,%1,%0\;vor %3,%3,%2
    vnand %4,%1,%0\;vor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: or op vor rtl ior inv 0 comp 0
-;; inner: nor op vnor rtl and inv 0 comp 3
+;; vector vnor -> vor
 (define_insn "*fuse_vnor_vor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnor %3,%1,%0\;vor %3,%3,%2
-   vnor %0,%1,%0\;vor %0,%0,%2
-   vnor %1,%1,%0\;vor %1,%1,%2
+   vnor %3,%1,%0\;vor %3,%3,%2
+   vnor %3,%1,%0\;vor %3,%3,%2
    vnor %4,%1,%0\;vor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: or op vor rtl ior inv 0 comp 0
-;; inner: or op vor rtl ior inv 0 comp 0
+;; vector vor -> vor
 (define_insn "*fuse_vor_vor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "%v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "%v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vor %3,%1,%0\;vor %3,%3,%2
-   vor %0,%1,%0\;vor %0,%0,%2
-   vor %1,%1,%0\;vor %1,%1,%2
+   vor %3,%1,%0\;vor %3,%3,%2
+   vor %3,%1,%0\;vor %3,%3,%2
    vor %4,%1,%0\;vor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: or op vor rtl ior inv 0 comp 0
-;; inner: orc op vorc rtl ior inv 0 comp 1
+;; vector vorc -> vor
 (define_insn "*fuse_vorc_vor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vorc %3,%1,%0\;vor %3,%3,%2
-   vorc %0,%1,%0\;vor %0,%0,%2
-   vorc %1,%1,%0\;vor %1,%1,%2
+   vorc %3,%1,%0\;vor %3,%3,%2
+   vorc %3,%1,%0\;vor %3,%3,%2
    vorc %4,%1,%0\;vor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: or op vor rtl ior inv 0 comp 0
-;; inner: xor op vxor rtl xor inv 0 comp 0
+;; vector vxor -> vor
 (define_insn "*fuse_vxor_vor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vxor %3,%1,%0\;vor %3,%3,%2
-   vxor %0,%1,%0\;vor %0,%0,%2
-   vxor %1,%1,%0\;vor %1,%1,%2
+   vxor %3,%1,%0\;vor %3,%3,%2
+   vxor %3,%1,%0\;vor %3,%3,%2
    vxor %4,%1,%0\;vor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: orc op vorc rtl ior inv 0 comp 1
-;; inner: and op vand rtl and inv 0 comp 0
+;; vector vand -> vorc
 (define_insn "*fuse_vand_vorc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vand %3,%1,%0\;vorc %3,%3,%2
-   vand %0,%1,%0\;vorc %0,%0,%2
-   vand %1,%1,%0\;vorc %1,%1,%2
+   vand %3,%1,%0\;vorc %3,%3,%2
+   vand %3,%1,%0\;vorc %3,%3,%2
    vand %4,%1,%0\;vorc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: orc op vorc rtl ior inv 0 comp 1
-;; inner: andc op vandc rtl and inv 0 comp 1
+;; vector vandc -> vorc
 (define_insn "*fuse_vandc_vorc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vandc %3,%1,%0\;vorc %3,%3,%2
-   vandc %0,%1,%0\;vorc %0,%0,%2
-   vandc %1,%1,%0\;vorc %1,%1,%2
+   vandc %3,%1,%0\;vorc %3,%3,%2
+   vandc %3,%1,%0\;vorc %3,%3,%2
    vandc %4,%1,%0\;vorc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: orc op vorc rtl ior inv 0 comp 1
-;; inner: eqv op veqv rtl xor inv 1 comp 0
+;; vector veqv -> vorc
 (define_insn "*fuse_veqv_vorc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    veqv %3,%1,%0\;vorc %3,%3,%2
-   veqv %0,%1,%0\;vorc %0,%0,%2
-   veqv %1,%1,%0\;vorc %1,%1,%2
+   veqv %3,%1,%0\;vorc %3,%3,%2
+   veqv %3,%1,%0\;vorc %3,%3,%2
    veqv %4,%1,%0\;vorc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: orc op vorc rtl ior inv 0 comp 1
-;; inner: nand op vnand rtl ior inv 0 comp 3
+;; vector vnand -> vorc
 (define_insn "*fuse_vnand_vorc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnand %3,%1,%0\;vorc %3,%3,%2
-   vnand %0,%1,%0\;vorc %0,%0,%2
-   vnand %1,%1,%0\;vorc %1,%1,%2
+   vnand %3,%1,%0\;vorc %3,%3,%2
+   vnand %3,%1,%0\;vorc %3,%3,%2
    vnand %4,%1,%0\;vorc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: orc op vorc rtl ior inv 0 comp 1
-;; inner: nor op vnor rtl and inv 0 comp 3
+;; vector vnor -> vorc
 (define_insn "*fuse_vnor_vorc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnor %3,%1,%0\;vorc %3,%3,%2
-   vnor %0,%1,%0\;vorc %0,%0,%2
-   vnor %1,%1,%0\;vorc %1,%1,%2
+   vnor %3,%1,%0\;vorc %3,%3,%2
+   vnor %3,%1,%0\;vorc %3,%3,%2
    vnor %4,%1,%0\;vorc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: orc op vorc rtl ior inv 0 comp 1
-;; inner: or op vor rtl ior inv 0 comp 0
+;; vector vor -> vorc
 (define_insn "*fuse_vor_vorc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vor %3,%1,%0\;vorc %3,%3,%2
-   vor %0,%1,%0\;vorc %0,%0,%2
-   vor %1,%1,%0\;vorc %1,%1,%2
+   vor %3,%1,%0\;vorc %3,%3,%2
+   vor %3,%1,%0\;vorc %3,%3,%2
    vor %4,%1,%0\;vorc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: orc op vorc rtl ior inv 0 comp 1
-;; inner: orc op vorc rtl ior inv 0 comp 1
+;; vector vorc -> vorc
 (define_insn "*fuse_vorc_vorc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vorc %3,%1,%0\;vorc %3,%3,%2
-   vorc %0,%1,%0\;vorc %0,%0,%2
-   vorc %1,%1,%0\;vorc %1,%1,%2
+   vorc %3,%1,%0\;vorc %3,%3,%2
+   vorc %3,%1,%0\;vorc %3,%3,%2
    vorc %4,%1,%0\;vorc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: orc op vorc rtl ior inv 0 comp 1
-;; inner: xor op vxor rtl xor inv 0 comp 0
+;; vector vxor -> vorc
 (define_insn "*fuse_vxor_vorc"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (ior:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (ior:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (not:VM (match_operand:VM 2 "altivec_register_operand" "v,v,v,v"))))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vxor %3,%1,%0\;vorc %3,%3,%2
-   vxor %0,%1,%0\;vorc %0,%0,%2
-   vxor %1,%1,%0\;vorc %1,%1,%2
+   vxor %3,%1,%0\;vorc %3,%3,%2
+   vxor %3,%1,%0\;vorc %3,%3,%2
    vxor %4,%1,%0\;vorc %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: xor op vxor rtl xor inv 0 comp 0
-;; inner: and op vand rtl and inv 0 comp 0
+;; vector vand -> vxor
 (define_insn "*fuse_vand_vxor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (xor:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (xor:VM (and:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vand %3,%1,%0\;vxor %3,%3,%2
-   vand %0,%1,%0\;vxor %0,%0,%2
-   vand %1,%1,%0\;vxor %1,%1,%2
+   vand %3,%1,%0\;vxor %3,%3,%2
+   vand %3,%1,%0\;vxor %3,%3,%2
    vand %4,%1,%0\;vxor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: xor op vxor rtl xor inv 0 comp 0
-;; inner: andc op vandc rtl and inv 0 comp 1
+;; vector vandc -> vxor
 (define_insn "*fuse_vandc_vxor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (xor:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (xor:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vandc %3,%1,%0\;vxor %3,%3,%2
-   vandc %0,%1,%0\;vxor %0,%0,%2
-   vandc %1,%1,%0\;vxor %1,%1,%2
+   vandc %3,%1,%0\;vxor %3,%3,%2
+   vandc %3,%1,%0\;vxor %3,%3,%2
    vandc %4,%1,%0\;vxor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: xor op vxor rtl xor inv 0 comp 0
-;; inner: eqv op veqv rtl xor inv 1 comp 0
+;; vector veqv -> vxor
 (define_insn "*fuse_veqv_vxor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (xor:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (xor:VM (not:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    veqv %3,%1,%0\;vxor %3,%3,%2
-   veqv %0,%1,%0\;vxor %0,%0,%2
-   veqv %1,%1,%0\;vxor %1,%1,%2
+   veqv %3,%1,%0\;vxor %3,%3,%2
+   veqv %3,%1,%0\;vxor %3,%3,%2
    veqv %4,%1,%0\;vxor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: xor op vxor rtl xor inv 0 comp 0
-;; inner: nand op vnand rtl ior inv 0 comp 3
+;; vector vnand -> vxor
 (define_insn "*fuse_vnand_vxor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (xor:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (xor:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnand %3,%1,%0\;vxor %3,%3,%2
-   vnand %0,%1,%0\;vxor %0,%0,%2
-   vnand %1,%1,%0\;vxor %1,%1,%2
+   vnand %3,%1,%0\;vxor %3,%3,%2
+   vnand %3,%1,%0\;vxor %3,%3,%2
    vnand %4,%1,%0\;vxor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: xor op vxor rtl xor inv 0 comp 0
-;; inner: nor op vnor rtl and inv 0 comp 3
+;; vector vnor -> vxor
 (define_insn "*fuse_vnor_vxor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (xor:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (xor:VM (and:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (not:VM (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vnor %3,%1,%0\;vxor %3,%3,%2
-   vnor %0,%1,%0\;vxor %0,%0,%2
-   vnor %1,%1,%0\;vxor %1,%1,%2
+   vnor %3,%1,%0\;vxor %3,%3,%2
+   vnor %3,%1,%0\;vxor %3,%3,%2
    vnor %4,%1,%0\;vxor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: xor op vxor rtl xor inv 0 comp 0
-;; inner: or op vor rtl ior inv 0 comp 0
+;; vector vor -> vxor
 (define_insn "*fuse_vor_vxor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (xor:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (xor:VM (ior:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vor %3,%1,%0\;vxor %3,%3,%2
-   vor %0,%1,%0\;vxor %0,%0,%2
-   vor %1,%1,%0\;vxor %1,%1,%2
+   vor %3,%1,%0\;vxor %3,%3,%2
+   vor %3,%1,%0\;vxor %3,%3,%2
    vor %4,%1,%0\;vxor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: xor op vxor rtl xor inv 0 comp 0
-;; inner: orc op vorc rtl ior inv 0 comp 1
+;; vector vorc -> vxor
 (define_insn "*fuse_vorc_vxor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (xor:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 1 "altivec_register_operand" "v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (xor:VM (ior:VM (not:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v")) 
+                          (match_operand:VM 1 "altivec_register_operand" "v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vorc %3,%1,%0\;vxor %3,%3,%2
-   vorc %0,%1,%0\;vxor %0,%0,%2
-   vorc %1,%1,%0\;vxor %1,%1,%2
+   vorc %3,%1,%0\;vxor %3,%3,%2
+   vorc %3,%1,%0\;vxor %3,%3,%2
    vorc %4,%1,%0\;vxor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
    (set_attr "length" "8")])
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: vector outer: xor op vxor rtl xor inv 0 comp 0
-;; inner: xor op vxor rtl xor inv 0 comp 0
+;; vector vxor -> vxor
 (define_insn "*fuse_vxor_vxor"
-  [(set (match_operand:VM 3 "altivec_register_operand" "=&v,0,1,v")
-        (xor:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") (match_operand:VM 1 "altivec_register_operand" "%v,v,v,v")) (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
-   (clobber (match_scratch:VM 4 "=X,X,X,r"))]
+  [(set (match_operand:VM 3 "altivec_register_operand" "=0,1,&v,v")
+        (xor:VM (xor:VM (match_operand:VM 0 "altivec_register_operand" "v,v,v,v") 
+                          (match_operand:VM 1 "altivec_register_operand" "%v,v,v,v"))
+                 (match_operand:VM 2 "altivec_register_operand" "v,v,v,v")))
+   (clobber (match_scratch:VM 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    vxor %3,%1,%0\;vxor %3,%3,%2
-   vxor %0,%1,%0\;vxor %0,%0,%2
-   vxor %1,%1,%0\;vxor %1,%1,%2
+   vxor %3,%1,%0\;vxor %3,%3,%2
+   vxor %3,%1,%0\;vxor %3,%3,%2
    vxor %4,%1,%0\;vxor %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
diff --git a/gcc/config/rs6000/genfusion.pl b/gcc/config/rs6000/genfusion.pl
index 837af7a..e1c45f5 100755
--- a/gcc/config/rs6000/genfusion.pl
+++ b/gcc/config/rs6000/genfusion.pl
@@ -192,7 +192,8 @@ sub gen_2logical
 	  if ( ($inner_comp & 2) == 2 ) {
 	      $inner_arg1 = "(not:${mode} $inner_arg1)";
 	  }
-	  $inner_exp = "(${inner_rtl}:${mode} ${inner_arg0} ${inner_arg1})";
+	  $inner_exp = "(${inner_rtl}:${mode} ${inner_arg0} 
+                          ${inner_arg1})";
 	  if ( $inner_inv == 1 ) {
 	      $inner_exp = "(not:${mode} $inner_exp)";
 	  }
@@ -203,7 +204,8 @@ sub gen_2logical
 	  if ( ($outer_comp & 2) == 2 ) {
 	      $inner_exp = "(not:${mode} $inner_exp)";
 	  }
-	  $outer_exp = "(${outer_rtl}:${mode} ${inner_exp} ${outer_arg2})";
+	  $outer_exp = "(${outer_rtl}:${mode} ${inner_exp}
+                 ${outer_arg2})";
 	  if ( $outer_inv == 1 ) {
 	      $outer_exp = "(not:${mode} $outer_exp)";
 	  }
@@ -211,17 +213,16 @@ sub gen_2logical
 	  $insn =  <<"EOF";
 
 ;; logical-logical fusion pattern generated by gen_2logical
-;; kind: $kind outer: $outer op $outer_op rtl $outer_rtl inv $outer_inv comp $outer_comp
-;; inner: $inner op $inner_op rtl $inner_rtl inv $inner_inv comp $inner_comp
+;; $kind $inner_op -> $outer_op
 (define_insn "*fuse_${inner_op}_${outer_op}"
-  [(set (match_operand:${mode} 3 "${pred}" "=&${constraint},0,1,${constraint}")
+  [(set (match_operand:${mode} 3 "${pred}" "=0,1,&${constraint},${constraint}")
         ${outer_exp})
-   (clobber (match_scratch:${mode} 4 "=X,X,X,r"))]
+   (clobber (match_scratch:${mode} 4 "=X,X,X,&r"))]
   "(TARGET_P10_FUSION && TARGET_P10_FUSION_2LOGICAL)"
   "@
    ${inner_op} %3,%1,%0\\;${outer_op} %3,%3,%2
-   ${inner_op} %0,%1,%0\\;${outer_op} %0,%0,%2
-   ${inner_op} %1,%1,%0\\;${outer_op} %1,%1,%2
+   ${inner_op} %3,%1,%0\\;${outer_op} %3,%3,%2
+   ${inner_op} %3,%1,%0\\;${outer_op} %3,%3,%2
    ${inner_op} %4,%1,%0\\;${outer_op} %3,%4,%2"
   [(set_attr "type" "logical")
    (set_attr "cost" "6")
diff --git a/gcc/config/rs6000/t-rs6000 b/gcc/config/rs6000/t-rs6000
index e3a58bf..1541a65 100644
--- a/gcc/config/rs6000/t-rs6000
+++ b/gcc/config/rs6000/t-rs6000
@@ -47,8 +47,8 @@ rs6000-call.o: $(srcdir)/config/rs6000/rs6000-call.c
 	$(COMPILE) $<
 	$(POSTCOMPILE)
 
-$(srcdir)/config/rs6000/fusion.md: $(srcdir)/config/rs6000/genfusion.pl
-	$(srcdir)/config/rs6000/genfusion.pl > $(srcdir)/config/rs6000/fusion.md
+#$(srcdir)/config/rs6000/fusion.md: $(srcdir)/config/rs6000/genfusion.pl
+#	$(srcdir)/config/rs6000/genfusion.pl > $(srcdir)/config/rs6000/fusion.md
 
 $(srcdir)/config/rs6000/rs6000-tables.opt: $(srcdir)/config/rs6000/genopt.sh \
   $(srcdir)/config/rs6000/rs6000-cpus.def