[3/11][POC][Clang][RISCV] Add typedef of the tuple type and define tuple type variant of vlseg2e32

[eopXD]

Summary

For the cover letter of this patch-set, please checkout D146872.

Depends on D146873.

This is the 3rd patch of the patch-set. This patch originates from
D99593.

Note: This patch is a proof-of-concept and will be extended to full
coverage in the future. Currently, the old non-tuple unit-stride
segment load is not removed, and only signed integer unit-strided
segment load of NF=2, EEW=32 is defined here.

When replacing the old intrinsics, the extra IsTuple parameter under
various places will be redundant and removed.

Authored-by: eop Chen <[e-mail]>
Co-Authored-by: Hsiangkai Wang <[e-mail]>

[eopXD]

Change: Add overloaded name vlseg2e32_tuple for the tuple type variant
to avoid naming collision to the existing non-tuple type overloaded
intrinsics of vlseg2e32.

This should resolve CI failure of clang/test/CodeGen/RISCV/rvv-intrinsics-autogenerated/non-policy/overloaded/vlseg2e32.c.

The overloaded name is temporary as the tuple type variant is targeted to replace the current non-tuple type segment load/store.

[craig.topper]

clang/include/clang/AST/ASTContext.h

1486

Use NF or NumFields instead of Tuple?

[eopXD]

Address comments from Craig.

clang/include/clang/AST/ASTContext.h

1486

Changed to NumFields.

[craig.topper]

clang/include/clang/Basic/riscv_vector.td

1758

Use an Offset variable like I suggested on the later patches.

[eopXD]

Address comments from Craig.

[eopXD]

Update code.

[craig.topper]

clang/include/clang/Basic/riscv_vector.td

1757

I don't think we need the Value * variables here. They're only benefit is naming, but I think we can do that with comments.

Something like this

unsigned Offset = IsMasked ? 1 : 0;
Operands.push_back(Ops[Offset]); // Pointer
if (IsMasked)
  Operands.push_back(Ops[0]); // Mask
Operands.push_back(Ops[Offset + 1]); // VL

clang/lib/AST/ASTContext.cpp

4156

Would it be possible to add NumFields to getScalableVectorType maybe with a default value of 1 instead of introducing a new function?

[eopXD]

Address comments from Craig.

[craig.topper]

clang/include/clang/Basic/riscv_vector.td

1760

MaskOperand can still be simplified.

[eopXD]

Removed another variable that is too verbose.

[craig.topper]

LGTM

[eopXD]

Rebase.

[eopXD]

Bump CI.

[eopXD]

Bump CI again due to latest fix (5a61920ed8b3e76d8f9bf39f2c1e18d552fcc976) for D146873 (5e92298f76875e1a89ad58bab042cd7abe9fc004).

[eopXD]

Bump CI again, the previous patch application failed.

[uabelho]

clang/include/clang/Support/RISCVVIntrinsicUtils.h

384

/home/buildbots/docker-RHEL84-buildbot/SetupBot/worker_env/ppc64le-clang-rhel-test/clang-ppc64le-rhel/llvm-project/clang/include/clang/Support/RISCVVIntrinsicUtils.h:384:8: error: private field 'IsTuple' is not used [-Werror,-Wunused-private-field]
  bool IsTuple = false;
       ^
1 error generated.

Seen e.g. here:
https://lab.llvm.org/buildbot/#/builders/57/builds/26989/steps/5/logs/stdio

[eopXD]

clang/include/clang/Support/RISCVVIntrinsicUtils.h

384

Thank you for the report. Resolved this in commit 38140255113297fe3e7926d311cd45004928779e.

clang/include/clang/AST/ASTContext.h

public:
   getBuiltinVectorTypeInfo(const BuiltinType *VecTy) const;
 
   /// Return the unique reference to a scalable vector type of the specified
   /// element type and scalable number of elements.
   ///
   /// \pre \p EltTy must be a built-in type.
   QualType getScalableVectorType(QualType EltTy, unsigned NumElts) const;
 
+  /// Return the unique reference to a scalable vector tuple type of the
+  /// specified element type, scalable number of elements and the number
+  /// of fields.
+  ///
+  /// \pre \p EltTy must be a built-in type.
+  QualType getScalableVectorTupleType(QualType EltTy, unsigned NumElts,
+                                      unsigned Tuple) const;
+
   /// Return a WebAssembly externref type.
   QualType getWebAssemblyExternrefType() const;
 
   /// Return the unique reference to a vector type of the specified
   /// element type and size.
   ///
   /// \pre \p VectorType must be a built-in type.
   QualType getVectorType(QualType VectorType, unsigned NumElts,

clang/include/clang/Basic/riscv_vector.td

 //   (SFixedLog2LMUL:Value): Smaller Fixed Log2LMUL. Given a vector type (SEW
 //      and LMUL), and computes another vector type which only changed LMUL as
 //      given value. The new LMUL should be smaller than the old one. Ignore to
 //      define a new builtin if its equivalent type has illegal lmul.
 //   (LFixedLog2LMUL:Value): Larger Fixed Log2LMUL. Given a vector type (SEW
 //      and LMUL), and computes another vector type which only changed LMUL as
 //      given value. The new LMUL should be larger than the old one. Ignore to
 //      define a new builtin if its equivalent type has illegal lmul.
+//   (Tuple:NF): Given a vector type, transform it to a vector tuple of NF.
 //
 // Following with the example above, if t is "i", then "Ue" will yield unsigned
 // int and "Fv" will yield __rvv_float32m1_t (again assuming LMUL=1), Fw would
 // yield __rvv_float64m2_t, etc.
 //
 // Each builtin is then defined by applying each type in TypeRange against the
 // sequence of type transformers described in Suffix and Prototype.
 //
class RVVBuiltin<string suffix, string prototype, string type_range,
   // Use clang_builtin_alias to save the number of builtins.
   bit HasBuiltinAlias = true;
 
   // Features required to enable for this builtin.
   list<string> RequiredFeatures = [];
 
   // Number of fields for Load/Store Segment instructions.
   int NF = 1;
+
+  bit IsTuple = false;
 }
 
 // This is the code emitted in the header.
 class RVVHeader {
   code HeaderCode;
 }
 
 //===----------------------------------------------------------------------===//

 // 7.8 Vector Load/Store Segment Instructions
 let UnMaskedPolicyScheme = HasPassthruOperand in {
 defm : RVVUnitStridedSegLoad<"vlseg">;
 defm : RVVUnitStridedSegLoadFF<"vlseg">;
 defm : RVVStridedSegLoad<"vlsseg">;
 defm : RVVIndexedSegLoad<"vluxseg">;
 defm : RVVIndexedSegLoad<"vloxseg">;
 }
+
+multiclass RVVUnitStridedSegLoadTuple<string op> {
+  foreach type = ["i"] in {
+    defvar eew = !cond(!eq(type, "i") : "32");
+      foreach nf = [2] in {
+        let Name = op # nf # "e" # eew # "_v_tuple",
+            IRName = op # nf,
+            MaskedIRName = op # nf # "_mask",
+            NF = nf,
+            ManualCodegen = [{
+    {
+      assert(((IsMasked && (PolicyAttrs & RVV_VTA) && (PolicyAttrs & RVV_VMA)) ||
+            (!IsMasked && (PolicyAttrs & RVV_VTA))) &&
+            "FIXME: Only handling default policy (TAMA) for now");
+
+      llvm::Type *ElementVectorType = cast<StructType>(ResultType)->elements()[0];
+      IntrinsicTypes = {ElementVectorType, Ops.back()->getType()};
+      SmallVector<llvm::Value*, 12> Operands;
+
+      Operands.append(NF, llvm::PoisonValue::get(ElementVectorType));
+
+      llvm::Value *MaskOperand = IsMasked ? Ops[0] : nullptr;
+      llvm::Value *PtrOperand = IsMasked ? Ops[1] : Ops[0];
+      llvm::Value *VLOperand = IsMasked ? Ops[2] : Ops[1];
+      Operands.push_back(PtrOperand);
+      if (MaskOperand)
+        Operands.push_back(MaskOperand);
+      Operands.push_back(VLOperand);
+      if (IsMasked)
+        Operands.push_back(ConstantInt::get(Ops.back()->getType(), PolicyAttrs));
+
+      llvm::Function *F = CGM.getIntrinsic(ID, IntrinsicTypes);
+
+      llvm::Value *LoadValue = Builder.CreateCall(F, Operands, "");
+      if (ReturnValue.isNull())
+        return LoadValue;
+      else
+        return Builder.CreateStore(LoadValue, ReturnValue.getValue());
+    }
+    }] in {
+          defvar T = "(Tuple:" # nf # ")";
+          def : RVVBuiltin<"v", T # "vPCe", type>;
+      }
+    }
+  }
+}
+// TODO: Extend for policy
+let UnMaskedPolicyScheme = NonePolicy,
+    MaskedPolicyScheme = NonePolicy,
+    IsTuple = true in {
+defm : RVVUnitStridedSegLoadTuple<"vlseg">;
+}
+
 let UnMaskedPolicyScheme = NonePolicy,
     MaskedPolicyScheme = NonePolicy in {
 defm : RVVUnitStridedSegStore<"vsseg">;
 defm : RVVStridedSegStore<"vssseg">;
 defm : RVVIndexedSegStore<"vsuxseg">;
 defm : RVVIndexedSegStore<"vsoxseg">;
 }
 

clang/include/clang/Support/RISCVVIntrinsicUtils.h

enum class VectorTypeModifier : uint8_t {
   LFixedLog2LMUL3,
   SFixedLog2LMULN3,
   SFixedLog2LMULN2,
   SFixedLog2LMULN1,
   SFixedLog2LMUL0,
   SFixedLog2LMUL1,
   SFixedLog2LMUL2,
   SFixedLog2LMUL3,
+  Tuple2,
 };
 
 // Similar to basic type but used to describe what's kind of type related to
 // basic vector type, used to compute type info of arguments.
 enum class BaseTypeModifier : uint8_t {
   Invalid,
   Scalar,
   Vector,
class RVVType {
   bool IsPointer = false;
   // IsConstant indices are "int", but have the constant expression.
   bool IsImmediate = false;
   // Const qualifier for pointer to const object or object of const type.
   bool IsConstant = false;
   unsigned ElementBitwidth = 0;
   VScaleVal Scale = 0;
   bool Valid;
+  bool IsTuple = false;
+  unsigned NF = 0;
 
   std::string BuiltinStr;
   std::string ClangBuiltinStr;
   std::string Str;
   std::string ShortStr;
 
   enum class FixedLMULType { LargerThan, SmallerThan };
 
public:
   bool isFloatVector(unsigned Width) const {
     return isVector() && isFloat() && ElementBitwidth == Width;
   }
   bool isFloat(unsigned Width) const {
     return isFloat() && ElementBitwidth == Width;
   }
   bool isConstant() const { return IsConstant; }
   bool isPointer() const { return IsPointer; }
+  bool isTuple() const { return IsTuple; }
   unsigned getElementBitwidth() const { return ElementBitwidth; }
 
   ScalarTypeKind getScalarType() const { return ScalarType; }
   VScaleVal getScale() const { return Scale; }
+  unsigned getNF() const {
+    assert(NF > 1 && NF < 8 && "Only legal NF should be fetched");
+    return NF;
+  }
 
 private:
   // Verify RVV vector type and set Valid.
   bool verifyType() const;
 
   // Creates a type based on basic types of TypeRange
   void applyBasicType();
 
private:
   std::string ManualCodegen;
   RVVTypePtr OutputType; // Builtin output type
   RVVTypes InputTypes;   // Builtin input types
   // The types we use to obtain the specific LLVM intrinsic. They are index of
   // InputTypes. -1 means the return type.
   std::vector<int64_t> IntrinsicTypes;
   unsigned NF = 1;
   Policy PolicyAttrs;
+  bool IsTuple = false;
 
 public:
   RVVIntrinsic(llvm::StringRef Name, llvm::StringRef Suffix,
                llvm::StringRef OverloadedName, llvm::StringRef OverloadedSuffix,
                llvm::StringRef IRName, bool IsMasked, bool HasMaskedOffOperand,
                bool HasVL, PolicyScheme Scheme, bool SupportOverloading,
                bool HasBuiltinAlias, llvm::StringRef ManualCodegen,
                const RVVTypes &Types,
                const std::vector<int64_t> &IntrinsicTypes,
                const std::vector<llvm::StringRef> &RequiredFeatures,
-               unsigned NF, Policy PolicyAttrs);
+               unsigned NF, Policy PolicyAttrs, bool IsTuple);
   ~RVVIntrinsic() = default;
 
   RVVTypePtr getOutputType() const { return OutputType; }
   const RVVTypes &getInputTypes() const { return InputTypes; }
   llvm::StringRef getBuiltinName() const { return BuiltinName; }
   llvm::StringRef getName() const { return Name; }
   llvm::StringRef getOverloadedName() const { return OverloadedName; }
   bool hasMaskedOffOperand() const { return HasMaskedOffOperand; }
public:
   static std::string
   getSuffixStr(RVVTypeCache &TypeCache, BasicType Type, int Log2LMUL,
                llvm::ArrayRef<PrototypeDescriptor> PrototypeDescriptors);
 
   static llvm::SmallVector<PrototypeDescriptor>
   computeBuiltinTypes(llvm::ArrayRef<PrototypeDescriptor> Prototype,
                       bool IsMasked, bool HasMaskedOffOperand, bool HasVL,
                       unsigned NF, PolicyScheme DefaultScheme,
-                      Policy PolicyAttrs);
+                      Policy PolicyAttrs, bool IsTuple);
 
   static llvm::SmallVector<Policy> getSupportedUnMaskedPolicies();
   static llvm::SmallVector<Policy>
       getSupportedMaskedPolicies(bool HasTailPolicy, bool HasMaskPolicy);
 
   static void updateNamesAndPolicy(bool IsMasked, bool HasPolicy,
                                    std::string &Name, std::string &BuiltinName,
                                    std::string &OverloadedName,
struct RVVIntrinsicRecord {
   // Number of fields, greater than 1 if it's segment load/store.
   uint8_t NF;
 
   bool HasMasked : 1;
   bool HasVL : 1;
   bool HasMaskedOffOperand : 1;
   bool HasTailPolicy : 1;
   bool HasMaskPolicy : 1;
+  bool IsTuple : 1;
   uint8_t UnMaskedPolicyScheme : 2;
   uint8_t MaskedPolicyScheme : 2;
 };
 
 llvm::raw_ostream &operator<<(llvm::raw_ostream &OS,
                               const RVVIntrinsicRecord &RVVInstrRecord);
 
 LLVM_ENABLE_BITMASK_ENUMS_IN_NAMESPACE();
 } // end namespace RISCV
 
 } // end namespace clang
 
 #endif // CLANG_SUPPORT_RISCVVINTRINSICUTILS_H

clang/lib/AST/ASTContext.cpp

 #define RVV_PREDICATE_TYPE(Name, Id, SingletonId, NumEls)                      \
     if (EltTy->isBooleanType() && NumElts == NumEls)                           \
       return SingletonId;
 #include "clang/Basic/RISCVVTypes.def"
   }
   return QualType();
 }
 
+/// getScalableVectorTupleType - Return the unique reference to a scalable
+/// vector tuple type of the specified element type, size and the number of
+/// fields.
+QualType ASTContext::getScalableVectorTupleType(QualType EltTy,
+                                                unsigned NumElts,
+                                                unsigned Tuple) const {
+  if (Target->hasRISCVVTypes()) {
+    uint64_t EltTySize = getTypeSize(EltTy);
+#define RVV_TYPE(Name, Id, SingletonId)
+#define RVV_VECTOR_TYPE(Name, Id, SingletonId, NumEls, ElBits, NF, IsSigned,   \
+                        IsFP)                                                  \
+  if (!EltTy->isBooleanType() &&                                               \
+      ((EltTy->hasIntegerRepresentation() &&                                   \
+        EltTy->hasSignedIntegerRepresentation() == IsSigned) ||                \
+       (EltTy->hasFloatingRepresentation() && IsFP)) &&                        \
+      EltTySize == ElBits && NumElts == NumEls && Tuple == NF)                 \
+    return SingletonId;
+#include "clang/Basic/RISCVVTypes.def"
+  }
+  return QualType();
+}
+
 /// getVectorType - Return the unique reference to a vector type of
 /// the specified element type and size. VectorType must be a built-in type.
 QualType ASTContext::getVectorType(QualType vecType, unsigned NumElts,
                                    VectorType::VectorKind VecKind) const {
   assert(vecType->isBuiltinType() ||
          (vecType->isBitIntType() &&
           // Only support _BitInt elements with byte-sized power of 2 NumBits.
           llvm::isPowerOf2_32(vecType->getAs<BitIntType>()->getNumBits()) &&

clang/lib/Sema/SemaRISCVVectorLookup.cpp

case 16:
       break;
     default:
       llvm_unreachable("Unsupported floating point width.");
     }
     break;
   case Invalid:
     llvm_unreachable("Unhandled type.");
   }
-  if (Type->isVector())
+  if (Type->isVector()) {
+    if (Type->isTuple())
+      QT = Context.getScalableVectorTupleType(QT, *Type->getScale(),
+                                              Type->getNF());
+    else
       QT = Context.getScalableVectorType(QT, *Type->getScale());
+  }
 
   if (Type->isConstant())
     QT = Context.getConstType(QT);
 
   // Transform the type to a pointer as the last step, if necessary.
   if (Type->isPointer())
     QT = Context.getPointerType(QT);
 
for (auto &Record : RVVIntrinsicRecords) {
     PolicyScheme UnMaskedPolicyScheme =
         static_cast<PolicyScheme>(Record.UnMaskedPolicyScheme);
     PolicyScheme MaskedPolicyScheme =
         static_cast<PolicyScheme>(Record.MaskedPolicyScheme);
 
     const Policy DefaultPolicy;
 
     llvm::SmallVector<PrototypeDescriptor> ProtoSeq =
-        RVVIntrinsic::computeBuiltinTypes(BasicProtoSeq, /*IsMasked=*/false,
-                                          /*HasMaskedOffOperand=*/false,
-                                          Record.HasVL, Record.NF,
-                                          UnMaskedPolicyScheme, DefaultPolicy);
+        RVVIntrinsic::computeBuiltinTypes(
+            BasicProtoSeq, /*IsMasked=*/false,
+            /*HasMaskedOffOperand=*/false, Record.HasVL, Record.NF,
+            UnMaskedPolicyScheme, DefaultPolicy, Record.IsTuple);
 
     llvm::SmallVector<PrototypeDescriptor> ProtoMaskSeq =
         RVVIntrinsic::computeBuiltinTypes(
             BasicProtoSeq, /*IsMasked=*/true, Record.HasMaskedOffOperand,
-            Record.HasVL, Record.NF, MaskedPolicyScheme, DefaultPolicy);
+            Record.HasVL, Record.NF, MaskedPolicyScheme, DefaultPolicy,
+            Record.IsTuple);
 
     bool UnMaskedHasPolicy = UnMaskedPolicyScheme != PolicyScheme::SchemeNone;
     bool MaskedHasPolicy = MaskedPolicyScheme != PolicyScheme::SchemeNone;
     SmallVector<Policy> SupportedUnMaskedPolicies =
         RVVIntrinsic::getSupportedUnMaskedPolicies();
     SmallVector<Policy> SupportedMaskedPolicies =
         RVVIntrinsic::getSupportedMaskedPolicies(Record.HasTailPolicy,
                                                  Record.HasMaskPolicy);
for (unsigned int TypeRangeMaskShift = 0;
 
         // Create non-masked policy intrinsic.
         if (Record.UnMaskedPolicyScheme != PolicyScheme::SchemeNone) {
           for (auto P : SupportedUnMaskedPolicies) {
             llvm::SmallVector<PrototypeDescriptor> PolicyPrototype =
                 RVVIntrinsic::computeBuiltinTypes(
                     BasicProtoSeq, /*IsMasked=*/false,
                     /*HasMaskedOffOperand=*/false, Record.HasVL, Record.NF,
-                    UnMaskedPolicyScheme, P);
+                    UnMaskedPolicyScheme, P, Record.IsTuple);
             std::optional<RVVTypes> PolicyTypes = TypeCache.computeTypes(
                 BaseType, Log2LMUL, Record.NF, PolicyPrototype);
             InitRVVIntrinsic(Record, SuffixStr, OverloadedSuffixStr,
                              /*IsMask=*/false, *PolicyTypes, UnMaskedHasPolicy,
                              P);
           }
         }
         if (!Record.HasMasked)
           continue;
         // Create masked intrinsic.
         std::optional<RVVTypes> MaskTypes =
             TypeCache.computeTypes(BaseType, Log2LMUL, Record.NF, ProtoMaskSeq);
+
         InitRVVIntrinsic(Record, SuffixStr, OverloadedSuffixStr, true,
                          *MaskTypes, MaskedHasPolicy, DefaultPolicy);
         if (Record.MaskedPolicyScheme == PolicyScheme::SchemeNone)
           continue;
         // Create masked policy intrinsic.
         for (auto P : SupportedMaskedPolicies) {
           llvm::SmallVector<PrototypeDescriptor> PolicyPrototype =
               RVVIntrinsic::computeBuiltinTypes(
                   BasicProtoSeq, /*IsMasked=*/true, Record.HasMaskedOffOperand,
-                  Record.HasVL, Record.NF, MaskedPolicyScheme, P);
+                  Record.HasVL, Record.NF, MaskedPolicyScheme, P,
+                  Record.IsTuple);
           std::optional<RVVTypes> PolicyTypes = TypeCache.computeTypes(
               BaseType, Log2LMUL, Record.NF, PolicyPrototype);
           InitRVVIntrinsic(Record, SuffixStr, OverloadedSuffixStr,
                            /*IsMask=*/true, *PolicyTypes, MaskedHasPolicy, P);
         }
       } // End for different LMUL
     } // End for different TypeRange
   }

clang/lib/Support/RISCVVIntrinsicUtils.cpp

bool RVVType::verifyType() const {
   if (ScalarType == Invalid)
     return false;
   if (isScalar())
     return true;
   if (!Scale)
     return false;
   if (isFloat() && ElementBitwidth == 8)
     return false;
+  if (IsTuple && (NF == 1 || NF > 8))
+    return false;
   unsigned V = *Scale;
   switch (ElementBitwidth) {
   case 1:
   case 8:
     // Check Scale is 1,2,4,8,16,32,64
     return (V <= 64 && isPowerOf2_32(V));
   case 16:
     // Check Scale is 1,2,4,8,16,32
if (isScalar()) {
     return;
   }
   BuiltinStr = "q" + utostr(*Scale) + BuiltinStr;
   // Pointer to vector types. Defined for segment load intrinsics.
   // segment load intrinsics have pointer type arguments to store the loaded
   // vector values.
   if (IsPointer)
     BuiltinStr += "*";
+
+  if (IsTuple)
+    BuiltinStr = "T" + utostr(NF) + BuiltinStr;
 }
 
 void RVVType::initClangBuiltinStr() {
   assert(isValid() && "RVVType is invalid");
   assert(isVector() && "Handle Vector type only");
 
   ClangBuiltinStr = "__rvv_";
   switch (ScalarType) {
   case ScalarTypeKind::Boolean:
     ClangBuiltinStr += "bool" + utostr(64 / *Scale) + "_t";
     return;
   case ScalarTypeKind::Float:
     ClangBuiltinStr += "float";
     break;
   case ScalarTypeKind::SignedInteger:
     ClangBuiltinStr += "int";
     break;
   case ScalarTypeKind::UnsignedInteger:
     ClangBuiltinStr += "uint";
     break;
   default:
     llvm_unreachable("ScalarTypeKind is invalid");
   }
-  ClangBuiltinStr += utostr(ElementBitwidth) + LMUL.str() + "_t";
+  ClangBuiltinStr += utostr(ElementBitwidth) + LMUL.str() +
+                     (IsTuple ? "x" + utostr(NF) : "") + "_t";
 }
 
 void RVVType::initTypeStr() {
   assert(isValid() && "RVVType is invalid");
 
   if (IsConstant)
     Str += "const ";
 
   auto getTypeString = [&](StringRef TypeStr) {
     if (isScalar())
       return Twine(TypeStr + Twine(ElementBitwidth) + "_t").str();
-    return Twine("v" + TypeStr + Twine(ElementBitwidth) + LMUL.str() + "_t")
+    return Twine("v" + TypeStr + Twine(ElementBitwidth) + LMUL.str() +
+                 (IsTuple ? "x" + utostr(NF) : "") + "_t")
         .str();
   };
 
   switch (ScalarType) {
   case ScalarTypeKind::Void:
     Str = "void";
     return;
   case ScalarTypeKind::Size_t:
if (ComplexTT.first == "Log2EEW") {
       case 3:
         VTM = VectorTypeModifier::SFixedLog2LMUL3;
         break;
       default:
         llvm_unreachable("Invalid LFixedLog2LMUL value, should be [-3, 3]");
         return std::nullopt;
       }
 
+    } else if (ComplexTT.first == "Tuple") {
+      unsigned NF = 0;
+      if (ComplexTT.second.getAsInteger(10, NF)) {
+        llvm_unreachable("Invalid NF value!");
+        return std::nullopt;
+      }
+      switch (NF) {
+      case 2:
+        VTM = VectorTypeModifier::Tuple2;
+        break;
+      default:
+        llvm_unreachable("Unhandled NF");
+      }
     } else {
       llvm_unreachable("Illegal complex type transformers!");
     }
   }
   PD.VTM = static_cast<uint8_t>(VTM);
 
   // Compute the remain type transformers
   TypeModifier TM = TypeModifier::NoModifier;
case VectorTypeModifier::SFixedLog2LMUL1:
     applyFixedLog2LMUL(1, FixedLMULType::SmallerThan);
     break;
   case VectorTypeModifier::SFixedLog2LMUL2:
     applyFixedLog2LMUL(2, FixedLMULType::SmallerThan);
     break;
   case VectorTypeModifier::SFixedLog2LMUL3:
     applyFixedLog2LMUL(3, FixedLMULType::SmallerThan);
     break;
+  case VectorTypeModifier::Tuple2: {
+    IsTuple = true;
+    NF = 2;
+    break;
+  }
   case VectorTypeModifier::NoModifier:
     break;
   }
 
   for (unsigned TypeModifierMaskShift = 0;
        TypeModifierMaskShift <= static_cast<unsigned>(TypeModifier::MaxOffset);
        ++TypeModifierMaskShift) {
     unsigned TypeModifierMask = 1 << TypeModifierMaskShift;
RVVIntrinsic::RVVIntrinsic(StringRef NewName, StringRef Suffix,
                            StringRef NewOverloadedName,
                            StringRef OverloadedSuffix, StringRef IRName,
                            bool IsMasked, bool HasMaskedOffOperand, bool HasVL,
                            PolicyScheme Scheme, bool SupportOverloading,
                            bool HasBuiltinAlias, StringRef ManualCodegen,
                            const RVVTypes &OutInTypes,
                            const std::vector<int64_t> &NewIntrinsicTypes,
                            const std::vector<StringRef> &RequiredFeatures,
-                           unsigned NF, Policy NewPolicyAttrs)
+                           unsigned NF, Policy NewPolicyAttrs, bool IsTuple)
     : IRName(IRName), IsMasked(IsMasked),
       HasMaskedOffOperand(HasMaskedOffOperand), HasVL(HasVL), Scheme(Scheme),
       SupportOverloading(SupportOverloading), HasBuiltinAlias(HasBuiltinAlias),
-      ManualCodegen(ManualCodegen.str()), NF(NF), PolicyAttrs(NewPolicyAttrs) {
+      ManualCodegen(ManualCodegen.str()), NF(NF), PolicyAttrs(NewPolicyAttrs),
+      IsTuple(IsTuple) {
 
   // Init BuiltinName, Name and OverloadedName
   BuiltinName = NewName.str();
   Name = BuiltinName;
   if (NewOverloadedName.empty())
     OverloadedName = NewName.split("_").first.str();
   else
     OverloadedName = NewOverloadedName.str();
for (auto PD : PrototypeDescriptors) {
     SuffixStrs.push_back((*T)->getShortStr());
   }
   return join(SuffixStrs, "_");
 }
 
 llvm::SmallVector<PrototypeDescriptor> RVVIntrinsic::computeBuiltinTypes(
     llvm::ArrayRef<PrototypeDescriptor> Prototype, bool IsMasked,
     bool HasMaskedOffOperand, bool HasVL, unsigned NF,
-    PolicyScheme DefaultScheme, Policy PolicyAttrs) {
+    PolicyScheme DefaultScheme, Policy PolicyAttrs, bool IsTuple) {
   SmallVector<PrototypeDescriptor> NewPrototype(Prototype.begin(),
                                                 Prototype.end());
   bool HasPassthruOp = DefaultScheme == PolicyScheme::HasPassthruOperand;
   if (IsMasked) {
     // If HasMaskedOffOperand, insert result type as first input operand if
     // need.
     if (HasMaskedOffOperand && !PolicyAttrs.isTAMAPolicy()) {
       if (NF == 1) {
if (HasMaskedOffOperand && !PolicyAttrs.isTAMAPolicy()) {
       }
     }
     if (HasMaskedOffOperand && NF > 1) {
       // Convert
       // (void, op0 address, op1 address, ..., maskedoff0, maskedoff1, ...)
       // to
       // (void, op0 address, op1 address, ..., mask, maskedoff0, maskedoff1,
       // ...)
+      if (IsTuple)
+        NewPrototype.insert(NewPrototype.begin() + 1,
+                            PrototypeDescriptor::Mask);
+      else
         NewPrototype.insert(NewPrototype.begin() + NF + 1,
                             PrototypeDescriptor::Mask);
     } else {
       // If IsMasked, insert PrototypeDescriptor:Mask as first input operand.
       NewPrototype.insert(NewPrototype.begin() + 1, PrototypeDescriptor::Mask);
     }
   } else {
     if (NF == 1) {
       if (PolicyAttrs.isTUPolicy() && HasPassthruOp)
         NewPrototype.insert(NewPrototype.begin(), NewPrototype[0]);
     } else if (PolicyAttrs.isTUPolicy() && HasPassthruOp) {
       // NF > 1 cases for segment load operations.
       // Convert
       // (void, op0 address, op1 address, ...)
       // to
       // (void, op0 address, op1 address, maskedoff0, maskedoff1, ...)
       PrototypeDescriptor MaskoffType = Prototype[1];
       MaskoffType.TM &= ~static_cast<uint8_t>(TypeModifier::Pointer);
       NewPrototype.insert(NewPrototype.begin() + NF + 1, NF, MaskoffType);
     }
  }
 
   // If HasVL, append PrototypeDescriptor:VL to last operand
   if (HasVL)
     NewPrototype.push_back(PrototypeDescriptor::VL);
+  if (IsTuple)
+    NewPrototype[0].VTM = static_cast<uint8_t>(VectorTypeModifier::Tuple2);
   return NewPrototype;
 }
 
 llvm::SmallVector<Policy> RVVIntrinsic::getSupportedUnMaskedPolicies() {
   return {Policy(Policy::PolicyType::Undisturbed)}; // TU
 }
 
 llvm::SmallVector<Policy>
raw_ostream &operator<<(raw_ostream &OS, const RVVIntrinsicRecord &Record) {
   OS << (int)Record.TypeRangeMask << ",";
   OS << (int)Record.Log2LMULMask << ",";
   OS << (int)Record.NF << ",";
   OS << (int)Record.HasMasked << ",";
   OS << (int)Record.HasVL << ",";
   OS << (int)Record.HasMaskedOffOperand << ",";
   OS << (int)Record.HasTailPolicy << ",";
   OS << (int)Record.HasMaskPolicy << ",";
+  OS << (int)Record.IsTuple << ",";
   OS << (int)Record.UnMaskedPolicyScheme << ",";
   OS << (int)Record.MaskedPolicyScheme << ",";
   OS << "},\n";
   return OS;
 }
 
 } // end namespace RISCV
 } // end namespace clang

clang/test/CodeGen/RISCV/rvv-intrinsics-autogenerated/non-policy/non-overloaded/vlseg2e32_tuple.c

+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 2
+// REQUIRES: riscv-registered-target
+// RUN: %clang_cc1 -triple riscv64 -target-feature +v -target-feature +zfh \
+// RUN:   -target-feature +experimental-zvfh -disable-O0-optnone  \
+// RUN:   -emit-llvm %s -o - | opt -S -passes=mem2reg | \
+// RUN:   FileCheck --check-prefix=CHECK-RV64 %s
+#include <riscv_vector.h>
+
+// CHECK-RV64-LABEL: define dso_local { <vscale x 2 x i32>, <vscale x 2 x i32> } @test_vlseg2e32_v_tuple_i32m1
+// CHECK-RV64-SAME: (ptr noundef [[BASE:%.*]], i64 noundef [[VL:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK-RV64-NEXT:  entry:
+// CHECK-RV64-NEXT:    [[TMP0:%.*]] = call { <vscale x 2 x i32>, <vscale x 2 x i32> } @llvm.riscv.vlseg2.nxv2i32.i64(<vscale x 2 x i32> poison, <vscale x 2 x i32> poison, ptr [[BASE]], i64 [[VL]])
+// CHECK-RV64-NEXT:    ret { <vscale x 2 x i32>, <vscale x 2 x i32> } [[TMP0]]
+//
+vint32m1x2_t test_vlseg2e32_v_tuple_i32m1(const int32_t *base, size_t vl) {
+  return __riscv_vlseg2e32_v_tuple_i32m1(base, vl);
+}
+
+// CHECK-RV64-LABEL: define dso_local { <vscale x 2 x i32>, <vscale x 2 x i32> } @test_vlseg2e32_v_tuple_i32m1_m
+// CHECK-RV64-SAME: (<vscale x 2 x i1> [[MASK:%.*]], ptr noundef [[BASE:%.*]], i64 noundef [[VL:%.*]]) #[[ATTR0]] {
+// CHECK-RV64-NEXT:  entry:
+// CHECK-RV64-NEXT:    [[TMP0:%.*]] = call { <vscale x 2 x i32>, <vscale x 2 x i32> } @llvm.riscv.vlseg2.mask.nxv2i32.i64(<vscale x 2 x i32> poison, <vscale x 2 x i32> poison, ptr [[BASE]], <vscale x 2 x i1> [[MASK]], i64 [[VL]], i64 3)
+// CHECK-RV64-NEXT:    ret { <vscale x 2 x i32>, <vscale x 2 x i32> } [[TMP0]]
+//
+vint32m1x2_t test_vlseg2e32_v_tuple_i32m1_m(vbool32_t mask, const int32_t *base, size_t vl) {
+  return __riscv_vlseg2e32_v_tuple_i32m1_m(mask, base, vl);
+}

clang/utils/TableGen/RISCVVEmitter.cpp

struct SemaRecord {
   // Number of field, large than 1 if it's segment load/store.
   unsigned NF;
 
   bool HasMasked :1;
   bool HasVL :1;
   bool HasMaskedOffOperand :1;
   bool HasTailPolicy : 1;
   bool HasMaskPolicy : 1;
+  bool IsTuple : 1;
   uint8_t UnMaskedPolicyScheme : 2;
   uint8_t MaskedPolicyScheme : 2;
 };
 
 // Compressed function signature table.
 class SemaSignatureTable {
 private:
   std::vector<PrototypeDescriptor> SignatureTable;
for (int Log2LMUL : Log2LMULs) {
         printType(*T);
         auto UT = TypeCache.computeType(
             BT, Log2LMUL,
             PrototypeDescriptor(BaseTypeModifier::Vector,
                                 VectorTypeModifier::NoModifier,
                                 TypeModifier::UnsignedInteger));
         printType(*UT);
       }
+      // FIXME: Expand more type declaration
+      if (I == 'i' && Log2LMUL == 0) { // vint32m1x2_t
+        auto TupleT = TypeCache.computeType(
+            BT, Log2LMUL,
+            PrototypeDescriptor(BaseTypeModifier::Vector,
+                                VectorTypeModifier::Tuple2,
+                                TypeModifier::SignedInteger));
+        if (TupleT)
+          printType(*TupleT);
+      }
     }
   }
 
   for (int Log2LMUL : Log2LMULs) {
     auto T = TypeCache.computeType(BasicType::Float16, Log2LMUL,
                                    PrototypeDescriptor::Vector);
     if (T)
       printType(*T);
for (auto *R : RV) {
     StringRef ManualCodegen = R->getValueAsString("ManualCodegen");
     std::vector<int64_t> IntrinsicTypes =
         R->getValueAsListOfInts("IntrinsicTypes");
     std::vector<StringRef> RequiredFeatures =
         R->getValueAsListOfStrings("RequiredFeatures");
     StringRef IRName = R->getValueAsString("IRName");
     StringRef MaskedIRName = R->getValueAsString("MaskedIRName");
     unsigned NF = R->getValueAsInt("NF");
+    bool IsTuple = R->getValueAsBit("IsTuple");
 
     const Policy DefaultPolicy;
     SmallVector<Policy> SupportedUnMaskedPolicies =
         RVVIntrinsic::getSupportedUnMaskedPolicies();
     SmallVector<Policy> SupportedMaskedPolicies =
         RVVIntrinsic::getSupportedMaskedPolicies(HasTailPolicy, HasMaskPolicy);
 
     // Parse prototype and create a list of primitive type with transformers
     // (operand) in Prototype. Prototype[0] is output operand.
     SmallVector<PrototypeDescriptor> BasicPrototype =
         parsePrototypes(Prototypes);
 
     SmallVector<PrototypeDescriptor> SuffixDesc = parsePrototypes(SuffixProto);
     SmallVector<PrototypeDescriptor> OverloadedSuffixDesc =
         parsePrototypes(OverloadedSuffixProto);
 
     // Compute Builtin types
     auto Prototype = RVVIntrinsic::computeBuiltinTypes(
         BasicPrototype, /*IsMasked=*/false,
         /*HasMaskedOffOperand=*/false, HasVL, NF, UnMaskedPolicyScheme,
-        DefaultPolicy);
+        DefaultPolicy, IsTuple);
     auto MaskedPrototype = RVVIntrinsic::computeBuiltinTypes(
         BasicPrototype, /*IsMasked=*/true, HasMaskedOffOperand, HasVL, NF,
-        MaskedPolicyScheme, DefaultPolicy);
+        MaskedPolicyScheme, DefaultPolicy, IsTuple);
 
     // Create Intrinsics for each type and LMUL.
     for (char I : TypeRange) {
       for (int Log2LMUL : Log2LMULList) {
         BasicType BT = ParseBasicType(I);
         std::optional<RVVTypes> Types =
             TypeCache.computeTypes(BT, Log2LMUL, NF, Prototype);
         // Ignored to create new intrinsic if there are any illegal types.
         if (!Types)
           continue;
 
         auto SuffixStr =
             RVVIntrinsic::getSuffixStr(TypeCache, BT, Log2LMUL, SuffixDesc);
         auto OverloadedSuffixStr = RVVIntrinsic::getSuffixStr(
             TypeCache, BT, Log2LMUL, OverloadedSuffixDesc);
         // Create a unmasked intrinsic
         Out.push_back(std::make_unique<RVVIntrinsic>(
             Name, SuffixStr, OverloadedName, OverloadedSuffixStr, IRName,
             /*IsMasked=*/false, /*HasMaskedOffOperand=*/false, HasVL,
             UnMaskedPolicyScheme, SupportOverloading, HasBuiltinAlias,
             ManualCodegen, *Types, IntrinsicTypes, RequiredFeatures, NF,
-            DefaultPolicy));
+            DefaultPolicy, IsTuple));
         if (UnMaskedPolicyScheme != PolicyScheme::SchemeNone)
           for (auto P : SupportedUnMaskedPolicies) {
             SmallVector<PrototypeDescriptor> PolicyPrototype =
                 RVVIntrinsic::computeBuiltinTypes(
                     BasicPrototype, /*IsMasked=*/false,
                     /*HasMaskedOffOperand=*/false, HasVL, NF,
-                    UnMaskedPolicyScheme, P);
+                    UnMaskedPolicyScheme, P, IsTuple);
             std::optional<RVVTypes> PolicyTypes =
                 TypeCache.computeTypes(BT, Log2LMUL, NF, PolicyPrototype);
             Out.push_back(std::make_unique<RVVIntrinsic>(
                 Name, SuffixStr, OverloadedName, OverloadedSuffixStr, IRName,
                 /*IsMask=*/false, /*HasMaskedOffOperand=*/false, HasVL,
                 UnMaskedPolicyScheme, SupportOverloading, HasBuiltinAlias,
                 ManualCodegen, *PolicyTypes, IntrinsicTypes, RequiredFeatures,
-                NF, P));
+                NF, P, IsTuple));
           }
         if (!HasMasked)
           continue;
         // Create a masked intrinsic
         std::optional<RVVTypes> MaskTypes =
             TypeCache.computeTypes(BT, Log2LMUL, NF, MaskedPrototype);
         Out.push_back(std::make_unique<RVVIntrinsic>(
             Name, SuffixStr, OverloadedName, OverloadedSuffixStr, MaskedIRName,
             /*IsMasked=*/true, HasMaskedOffOperand, HasVL, MaskedPolicyScheme,
             SupportOverloading, HasBuiltinAlias, ManualCodegen, *MaskTypes,
-            IntrinsicTypes, RequiredFeatures, NF, DefaultPolicy));
+            IntrinsicTypes, RequiredFeatures, NF, DefaultPolicy, IsTuple));
         if (MaskedPolicyScheme == PolicyScheme::SchemeNone)
           continue;
         for (auto P : SupportedMaskedPolicies) {
           SmallVector<PrototypeDescriptor> PolicyPrototype =
               RVVIntrinsic::computeBuiltinTypes(
                   BasicPrototype, /*IsMasked=*/true, HasMaskedOffOperand, HasVL,
-                  NF, MaskedPolicyScheme, P);
+                  NF, MaskedPolicyScheme, P, IsTuple);
           std::optional<RVVTypes> PolicyTypes =
               TypeCache.computeTypes(BT, Log2LMUL, NF, PolicyPrototype);
           Out.push_back(std::make_unique<RVVIntrinsic>(
               Name, SuffixStr, OverloadedName, OverloadedSuffixStr,
               MaskedIRName, /*IsMasked=*/true, HasMaskedOffOperand, HasVL,
               MaskedPolicyScheme, SupportOverloading, HasBuiltinAlias,
               ManualCodegen, *PolicyTypes, IntrinsicTypes, RequiredFeatures, NF,
-              P));
+              P, IsTuple));
         }
       } // End for Log2LMULList
     }   // End for TypeRange
 
     // We don't emit vsetvli and vsetvlimax for SemaRecord.
     // They are written in riscv_vector.td and will emit those marco define in
     // riscv_vector.h
     if (Name == "vsetvli" || Name == "vsetvlimax")
for (auto *R : RV) {
     SR.HasMaskedOffOperand = HasMaskedOffOperand;
     SR.HasTailPolicy = HasTailPolicy;
     SR.HasMaskPolicy = HasMaskPolicy;
     SR.UnMaskedPolicyScheme = static_cast<uint8_t>(UnMaskedPolicyScheme);
     SR.MaskedPolicyScheme = static_cast<uint8_t>(MaskedPolicyScheme);
     SR.Prototype = std::move(BasicPrototype);
     SR.Suffix = parsePrototypes(SuffixProto);
     SR.OverloadedSuffix = parsePrototypes(OverloadedSuffixProto);
+    SR.IsTuple = IsTuple;
 
     SemaRecords->push_back(SR);
   }
 }
 
 void RVVEmitter::printHeaderCode(raw_ostream &OS) {
   std::vector<Record *> RVVHeaders =
       Records.getAllDerivedDefinitions("RVVHeader");
for (const auto &SR : SemaRecords) {
     R.NF = SR.NF;
     R.HasMasked = SR.HasMasked;
     R.HasVL = SR.HasVL;
     R.HasMaskedOffOperand = SR.HasMaskedOffOperand;
     R.HasTailPolicy = SR.HasTailPolicy;
     R.HasMaskPolicy = SR.HasMaskPolicy;
     R.UnMaskedPolicyScheme = SR.UnMaskedPolicyScheme;
     R.MaskedPolicyScheme = SR.MaskedPolicyScheme;
+    R.IsTuple = SR.IsTuple;
 
     assert(R.PrototypeIndex !=
            static_cast<uint16_t>(SemaSignatureTable::INVALID_INDEX));
     assert(R.SuffixIndex !=
            static_cast<uint16_t>(SemaSignatureTable::INVALID_INDEX));
     assert(R.OverloadedSuffixIndex !=
            static_cast<uint16_t>(SemaSignatureTable::INVALID_INDEX));
   }