[v1,2/3] RISC-V: Add test cases for vwmulsu.vv reg overlap

diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/group_overlap.h b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/group_overlap.h
index 647622cae5a..9d6e5f2b804 100644
--- a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/group_overlap.h
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/group_overlap.h
@@ -391,6 +391,148 @@ 
     ST_F ((void *)out, vd14, VL); OUT += VL;                                \
     ST_F ((void *)out, vd15, VL); OUT += VL;                                \
 
+#define LOOP_DUAL_WIDEN_BINARY_BODY_SU_X4(NT, NUT, WT, LD_NF,               \
+					   LD_NUF, OUT_F, ST_F, OUT,        \
+					   START, VL)                       \
+    NT vs0 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs1 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs2 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs3 = LD_NF ((void *)START, VL); START += VL;                        \
+    NUT vt0 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt1 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt2 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt3 = LD_NUF ((void *)START, VL); START += VL;                      \
+                                                                            \
+    asm volatile("nop" ::: "memory");                                       \
+                                                                            \
+    WT vd0 = OUT_F (vs0, vt0, VL);                                          \
+    WT vd1 = OUT_F (vs1, vt1, VL);                                          \
+    WT vd2 = OUT_F (vs2, vt2, VL);                                          \
+    WT vd3 = OUT_F (vs3, vt3, VL);                                          \
+                                                                            \
+    asm volatile("nop" ::: "memory");                                       \
+                                                                            \
+    ST_F ((void *)out, vd0, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd1, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd2, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd3, VL); OUT += VL;                                 \
+
+#define LOOP_DUAL_WIDEN_BINARY_BODY_SU_X8(NT, NUT, WT, LD_NF,               \
+					   LD_NUF, OUT_F, ST_F, OUT,        \
+					   START, VL)                       \
+    NT vs0 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs1 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs2 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs3 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs4 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs5 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs6 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs7 = LD_NF ((void *)START, VL); START += VL;                        \
+    NUT vt0 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt1 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt2 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt3 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt4 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt5 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt6 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt7 = LD_NUF ((void *)START, VL); START += VL;                      \
+                                                                            \
+    asm volatile("nop" ::: "memory");                                       \
+                                                                            \
+    WT vd0 = OUT_F (vs0, vt0, VL);                                          \
+    WT vd1 = OUT_F (vs1, vt1, VL);                                          \
+    WT vd2 = OUT_F (vs2, vt2, VL);                                          \
+    WT vd3 = OUT_F (vs3, vt3, VL);                                          \
+    WT vd4 = OUT_F (vs4, vt4, VL);                                          \
+    WT vd5 = OUT_F (vs5, vt5, VL);                                          \
+    WT vd6 = OUT_F (vs6, vt6, VL);                                          \
+    WT vd7 = OUT_F (vs7, vt7, VL);                                          \
+                                                                            \
+    asm volatile("nop" ::: "memory");                                       \
+                                                                            \
+    ST_F ((void *)out, vd0, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd1, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd2, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd3, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd4, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd5, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd6, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd7, VL); OUT += VL;                                 \
+
+#define LOOP_DUAL_WIDEN_BINARY_BODY_SU_X16(NT, NUT, WT, LD_NF,              \
+					   LD_NUF, OUT_F, ST_F, OUT,        \
+					   START, VL)                       \
+    NT vs0 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs1 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs2 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs3 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs4 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs5 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs6 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs7 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs8 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs9 = LD_NF ((void *)START, VL); START += VL;                        \
+    NT vs10 = LD_NF ((void *)START, VL); START += VL;                       \
+    NT vs11 = LD_NF ((void *)START, VL); START += VL;                       \
+    NT vs12 = LD_NF ((void *)START, VL); START += VL;                       \
+    NT vs13 = LD_NF ((void *)START, VL); START += VL;                       \
+    NT vs14 = LD_NF ((void *)START, VL); START += VL;                       \
+    NT vs15 = LD_NF ((void *)START, VL); START += VL;                       \
+    NUT vt0 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt1 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt2 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt3 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt4 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt5 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt6 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt7 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt8 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt9 = LD_NUF ((void *)START, VL); START += VL;                      \
+    NUT vt10 = LD_NUF ((void *)START, VL); START += VL;                     \
+    NUT vt11 = LD_NUF ((void *)START, VL); START += VL;                     \
+    NUT vt12 = LD_NUF ((void *)START, VL); START += VL;                     \
+    NUT vt13 = LD_NUF ((void *)START, VL); START += VL;                     \
+    NUT vt14 = LD_NUF ((void *)START, VL); START += VL;                     \
+    NUT vt15 = LD_NUF ((void *)START, VL); START += VL;                     \
+                                                                            \
+    asm volatile("nop" ::: "memory");                                       \
+                                                                            \
+    WT vd0 = OUT_F (vs0, vt0, VL);                                          \
+    WT vd1 = OUT_F (vs1, vt1, VL);                                          \
+    WT vd2 = OUT_F (vs2, vt2, VL);                                          \
+    WT vd3 = OUT_F (vs3, vt3, VL);                                          \
+    WT vd4 = OUT_F (vs4, vt4, VL);                                          \
+    WT vd5 = OUT_F (vs5, vt5, VL);                                          \
+    WT vd6 = OUT_F (vs6, vt6, VL);                                          \
+    WT vd7 = OUT_F (vs7, vt7, VL);                                          \
+    WT vd8 = OUT_F (vs8, vt8, VL);                                          \
+    WT vd9 = OUT_F (vs9, vt9, VL);                                          \
+    WT vd10 = OUT_F (vs10, vt10, VL);                                       \
+    WT vd11 = OUT_F (vs11, vt11, VL);                                       \
+    WT vd12 = OUT_F (vs12, vt12, VL);                                       \
+    WT vd13 = OUT_F (vs13, vt13, VL);                                       \
+    WT vd14 = OUT_F (vs14, vt14, VL);                                       \
+    WT vd15 = OUT_F (vs15, vt15, VL);                                       \
+                                                                            \
+    asm volatile("nop" ::: "memory");                                       \
+                                                                            \
+    ST_F ((void *)out, vd0, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd1, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd2, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd3, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd4, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd5, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd6, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd7, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd8, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd9, VL); OUT += VL;                                 \
+    ST_F ((void *)out, vd10, VL); OUT += VL;                                \
+    ST_F ((void *)out, vd11, VL); OUT += VL;                                \
+    ST_F ((void *)out, vd12, VL); OUT += VL;                                \
+    ST_F ((void *)out, vd13, VL); OUT += VL;                                \
+    ST_F ((void *)out, vd14, VL); OUT += VL;                                \
+    ST_F ((void *)out, vd15, VL); OUT += VL;                                \
+
 #define DEF_GROUP_OVERLAP_UNARY_0(VL_F, NT, WT, LD_F, OUT_F, ST_F, NAME, \
 				  LOOP_BODY)                             \
   void test_group_overlap_##NAME##_##NT##_unary_0(uint8_t *data,         \
@@ -436,4 +578,20 @@ 
     }                                                                       \
   }
 
+#define DEF_GROUP_OVERLAP_BINARY_2(VL_F, NT, NUT, WT, LD_NF, LD_NUF,        \
+				   OUT_F, ST_F, NAME, LOOP_BODY)            \
+  void test_group_overlap_##NAME##_##NT##_binary_2(uint8_t *data,           \
+						   uint8_t *out,            \
+						   size_t limit)            \
+  {                                                                         \
+    uint8_t *start = data;                                                  \
+    uint8_t *end = data + limit;                                            \
+    size_t vl = VL_F ();                                                    \
+                                                                            \
+    while (start < end) {                                                   \
+      LOOP_BODY (NT, NUT, WT, LD_NF, LD_NUF, OUT_F, ST_F, out, start,       \
+		 vl);                                                       \
+    }                                                                       \
+  }
+
 #endif
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-m1.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-m1.c
new file mode 100644
index 00000000000..81210f52f94
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-m1.c
@@ -0,0 +1,19 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e16m1,
+  vint16m1_t,
+  vuint16m1_t,
+  vint32m2_t,
+  __riscv_vle16_v_i16m1,
+  __riscv_vle16_v_u16m1,
+  __riscv_vwmulsu_vv_i32m2,
+  __riscv_vse32_v_i32m2,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X16)
+
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v2,v1,v3([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v10,v11,v3([^0-9]|$)} 1 } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-m2.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-m2.c
new file mode 100644
index 00000000000..3711878cb90
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-m2.c
@@ -0,0 +1,19 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e16m2,
+  vint16m2_t,
+  vuint16m2_t,
+  vint32m4_t,
+  __riscv_vle16_v_i16m2,
+  __riscv_vle16_v_u16m2,
+  __riscv_vwmulsu_vv_i32m4,
+  __riscv_vse32_v_i32m4,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X8)
+
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v8,v10,v6([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v28,v30,v24([^0-9]|$)} 1 } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-m4.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-m4.c
new file mode 100644
index 00000000000..69b895e2c68
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-m4.c
@@ -0,0 +1,19 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e16m4,
+  vint16m4_t,
+  vuint16m4_t,
+  vint32m8_t,
+  __riscv_vle16_v_i16m4,
+  __riscv_vle16_v_u16m4,
+  __riscv_vwmulsu_vv_i32m8,
+  __riscv_vse32_v_i32m8,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X4)
+
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v8,v4,v12([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v24,v28,v16([^0-9]|$)} 1 } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-mf2.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-mf2.c
new file mode 100644
index 00000000000..49df7b7005f
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-mf2.c
@@ -0,0 +1,21 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e16m1,
+  vint16mf2_t,
+  vuint16mf2_t,
+  vint32m1_t,
+  __riscv_vle16_v_i16mf2,
+  __riscv_vle16_v_u16mf2,
+  __riscv_vwmulsu_vv_i32m1,
+  __riscv_vse32_v_i32m1,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+   register group must not overlap either source at all.  */
+/* { dg-final { scan-assembler-not {vwmulsu\.vv\s+(v[0-9]+),\1,} } } */
+/* { dg-final { scan-assembler-not {vwmulsu\.vv\s+(v[0-9]+),v[0-9]+,\1([^0-9]|$)} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-mf4.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-mf4.c
new file mode 100644
index 00000000000..07d927d87cc
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i16-mf4.c
@@ -0,0 +1,21 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e16m1,
+  vint16mf4_t,
+  vuint16mf4_t,
+  vint32mf2_t,
+  __riscv_vle16_v_i16mf4,
+  __riscv_vle16_v_u16mf4,
+  __riscv_vwmulsu_vv_i32mf2,
+  __riscv_vse32_v_i32mf2,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+   register group must not overlap either source at all.  */
+/* { dg-final { scan-assembler-not {vwmulsu\.vv\s+(v[0-9]+),\1,} } } */
+/* { dg-final { scan-assembler-not {vwmulsu\.vv\s+(v[0-9]+),v[0-9]+,\1([^0-9]|$)} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i32-m1.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i32-m1.c
new file mode 100644
index 00000000000..9c5a435b58e
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i32-m1.c
@@ -0,0 +1,19 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e32m1,
+  vint32m1_t,
+  vuint32m1_t,
+  vint64m2_t,
+  __riscv_vle32_v_i32m1,
+  __riscv_vle32_v_u32m1,
+  __riscv_vwmulsu_vv_i64m2,
+  __riscv_vse64_v_i64m2,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X16)
+
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v2,v1,v3([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v10,v11,v3([^0-9]|$)} 1 } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i32-m2.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i32-m2.c
new file mode 100644
index 00000000000..833b8906083
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i32-m2.c
@@ -0,0 +1,19 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e32m2,
+  vint32m2_t,
+  vuint32m2_t,
+  vint64m4_t,
+  __riscv_vle32_v_i32m2,
+  __riscv_vle32_v_u32m2,
+  __riscv_vwmulsu_vv_i64m4,
+  __riscv_vse64_v_i64m4,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X8)
+
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v8,v10,v6([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v28,v30,v24([^0-9]|$)} 1 } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i32-m4.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i32-m4.c
new file mode 100644
index 00000000000..a769b55966e
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i32-m4.c
@@ -0,0 +1,19 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e32m4,
+  vint32m4_t,
+  vuint32m4_t,
+  vint64m8_t,
+  __riscv_vle32_v_i32m4,
+  __riscv_vle32_v_u32m4,
+  __riscv_vwmulsu_vv_i64m8,
+  __riscv_vse64_v_i64m8,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X4)
+
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v8,v4,v12([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v24,v28,v16([^0-9]|$)} 1 } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i32-mf2.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i32-mf2.c
new file mode 100644
index 00000000000..50b37333588
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i32-mf2.c
@@ -0,0 +1,21 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e32m1,
+  vint32mf2_t,
+  vuint32mf2_t,
+  vint64m1_t,
+  __riscv_vle32_v_i32mf2,
+  __riscv_vle32_v_u32mf2,
+  __riscv_vwmulsu_vv_i64m1,
+  __riscv_vse64_v_i64m1,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+   register group must not overlap either source at all.  */
+/* { dg-final { scan-assembler-not {vwmulsu\.vv\s+(v[0-9]+),\1,} } } */
+/* { dg-final { scan-assembler-not {vwmulsu\.vv\s+(v[0-9]+),v[0-9]+,\1([^0-9]|$)} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-m1.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-m1.c
new file mode 100644
index 00000000000..383d29f26f5
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-m1.c
@@ -0,0 +1,19 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e8m1,
+  vint8m1_t,
+  vuint8m1_t,
+  vint16m2_t,
+  __riscv_vle8_v_i8m1,
+  __riscv_vle8_v_u8m1,
+  __riscv_vwmulsu_vv_i16m2,
+  __riscv_vse16_v_i16m2,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X16)
+
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v2,v1,v3([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v10,v11,v3([^0-9]|$)} 1 } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-m2.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-m2.c
new file mode 100644
index 00000000000..d9de8bae64a
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-m2.c
@@ -0,0 +1,19 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e8m2,
+  vint8m2_t,
+  vuint8m2_t,
+  vint16m4_t,
+  __riscv_vle8_v_i8m2,
+  __riscv_vle8_v_u8m2,
+  __riscv_vwmulsu_vv_i16m4,
+  __riscv_vse16_v_i16m4,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X8)
+
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v8,v10,v6([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v28,v30,v24([^0-9]|$)} 1 } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-m4.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-m4.c
new file mode 100644
index 00000000000..4360d79e313
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-m4.c
@@ -0,0 +1,19 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e8m4,
+  vint8m4_t,
+  vuint8m4_t,
+  vint16m8_t,
+  __riscv_vle8_v_i8m4,
+  __riscv_vle8_v_u8m4,
+  __riscv_vwmulsu_vv_i16m8,
+  __riscv_vse16_v_i16m8,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X4)
+
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v8,v4,v12([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmulsu\.vv\s+v24,v28,v16([^0-9]|$)} 1 } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-mf2.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-mf2.c
new file mode 100644
index 00000000000..6691f4f6b8e
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-mf2.c
@@ -0,0 +1,21 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e8m1,
+  vint8mf2_t,
+  vuint8mf2_t,
+  vint16m1_t,
+  __riscv_vle8_v_i8mf2,
+  __riscv_vle8_v_u8mf2,
+  __riscv_vwmulsu_vv_i16m1,
+  __riscv_vse16_v_i16m1,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+   register group must not overlap either source at all.  */
+/* { dg-final { scan-assembler-not {vwmulsu\.vv\s+(v[0-9]+),\1,} } } */
+/* { dg-final { scan-assembler-not {vwmulsu\.vv\s+(v[0-9]+),v[0-9]+,\1([^0-9]|$)} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-mf4.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-mf4.c
new file mode 100644
index 00000000000..7c0a16085c4
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-mf4.c
@@ -0,0 +1,21 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e8m1,
+  vint8mf4_t,
+  vuint8mf4_t,
+  vint16mf2_t,
+  __riscv_vle8_v_i8mf4,
+  __riscv_vle8_v_u8mf4,
+  __riscv_vwmulsu_vv_i16mf2,
+  __riscv_vse16_v_i16mf2,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+   register group must not overlap either source at all.  */
+/* { dg-final { scan-assembler-not {vwmulsu\.vv\s+(v[0-9]+),\1,} } } */
+/* { dg-final { scan-assembler-not {vwmulsu\.vv\s+(v[0-9]+),v[0-9]+,\1([^0-9]|$)} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-mf8.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-mf8.c
new file mode 100644
index 00000000000..bba41cf3011
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmulsu_vv-i8-mf8.c
@@ -0,0 +1,21 @@ 
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_BINARY_2(
+  __riscv_vsetvlmax_e8m1,
+  vint8mf8_t,
+  vuint8mf8_t,
+  vint16mf4_t,
+  __riscv_vle8_v_i8mf8,
+  __riscv_vle8_v_u8mf8,
+  __riscv_vwmulsu_vv_i16mf4,
+  __riscv_vse16_v_i16mf4,
+  vwmulsu_vv,
+  LOOP_DUAL_WIDEN_BINARY_BODY_SU_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+   register group must not overlap either source at all.  */
+/* { dg-final { scan-assembler-not {vwmulsu\.vv\s+(v[0-9]+),\1,} } } */
+/* { dg-final { scan-assembler-not {vwmulsu\.vv\s+(v[0-9]+),v[0-9]+,\1([^0-9]|$)} } } */