[PoC][RISCV] Use an attribute to declare C intrinsics with different policy.

[HsiangKai]

Summary

This patch provides a proof-of-concept implementation of the proposal. https://github.com/riscv-non-isa/rvv-intrinsic-doc/issues/120

In this patch, we create a new attribute rvv_policy to annotate C
intrinsics with its tail/inactive elements policy. The syntax is

__attribute__((rvv_policy(tama)))
vint32m1_t vadd_vv_i8m1_tama(...);

The possible policies are tama, tamu, tuma, tumu.
ta: tail agnostic
tu: tail undisturbed
ma: inactive masked-off agnostic
mu: inactive masked-off undisturbed

The attribute is used in riscv_vector.h. This implementation will not increase the number of clang builtins.

[HsiangKai]

Remove redundant test case.

[craig.topper]

clang/lib/CodeGen/CGBuiltin.cpp

18816

Why size_t? This would be the size_t of the host machine that's building/running the compiler and would have no connection to the architecture being targetted.

clang/utils/TableGen/RISCVVEmitter.cpp

843

Do we need to emit this switch for every builtin? Couldn't we assign PolicyValue before including the autogenerated file and only builtins that have a policy would add PolicyValue it to their Ops vector?

916

Does the InputTypes.empty() check provide any value. Looks like it just prevents constructing a ListSeparator that wouldn't be used, but I would think that's cheap to construct. I know this was copied from the function above, so my question applies there too.

[HsiangKai]

Address @craig.topper's comments.

[craig.topper]

I think the concept seems good to me. I'd like @aaron.ballman to review the attribute code.

clang/lib/CodeGen/CGBuiltin.cpp

18812–18832

Is constant still used?

clang/utils/TableGen/RISCVVEmitter.cpp

883

Capitalize policySuffix and make it

static const char *const PolicySuffix[]

[frasercrmck]

Just nits from me at this stage.

clang/include/clang/Basic/AttrDocs.td

2171

Nit, but the use of could seems out of place in this documentation. Is can or may perhaps more common?

clang/utils/TableGen/RISCVVEmitter.cpp

895

This variable i shadowing the outer loop's induction variable is a little odd. Perhaps the outer loop could use a range-based for?

[luke957]

So sorry for my bad herald script.

[HsiangKai]

clang/lib/CodeGen/CGBuiltin.cpp

18812–18832

Yes, it is still used in ManualCodegenMask in clang/include/clang/Basic/riscv_vector.td.

[HsiangKai]

clang/lib/CodeGen/CGBuiltin.cpp

18812–18832

Sorry, after reviewing these usage, I think I could remove it.

[HsiangKai]

Address @craig.topper and @frasercrmck's comments.

[HsiangKai]

Fix build errors.

[HsiangKai]

Update attribute test cases.

[HsiangKai]

In riscv-insert-vsetvli, use the policy argument. No use implicit-def maskedoff to adjust the setting.

clang/include/clang/Basic/Attr.td

def RISCVInterrupt : InheritableAttr, TargetSpecificAttr<TargetRISCV> {
   let Args = [EnumArgument<"Interrupt", "InterruptType",
                            ["user", "supervisor", "machine"],
                            ["user", "supervisor", "machine"],
                            1>];
   let ParseKind = "Interrupt";
   let Documentation = [RISCVInterruptDocs];
 }
 
+def RISCVVPolicy : InheritableAttr, TargetSpecificAttr<TargetRISCV> {
+  let Spellings = [Clang<"rvv_policy">];
+  let Subjects = SubjectList<[Function]>;
+  let Args = [EnumArgument<"Policy", "PolicyType",
+              ["tumu", "tamu", "tuma", "tama"],
+              ["TUMU", "TAMU", "TUMA", "TAMA"]>];
+  let Documentation = [RISCVVPolicyDocs];
+}
+
 // This is not a TargetSpecificAttr so that is silently accepted and
 // ignored on other targets as encouraged by the OpenCL spec.
 //
 // See OpenCL 1.2 6.11.5: "It is our intention that a particular
 // implementation of OpenCL be free to ignore all attributes and the
 // resulting executable binary will produce the same result."
 //
 // However, only AMD GPU targets will emit the corresponding IR

clang/include/clang/Basic/AttrDocs.td

 Refer to:
 https://gcc.gnu.org/onlinedocs/gcc/RISC-V-Function-Attributes.html
 https://riscv.org/specifications/privileged-isa/
 The RISC-V Instruction Set Manual Volume II: Privileged Architecture
 Version 1.10.
   }];
 }
 
+def RISCVVPolicyDocs : Documentation {
+  let Category = DocCatFunction;
+  let Heading = "RISC-V vector tail/mask policy";
+  let Content = [{
+Users can use the attribute to specify the policy of destination tail and
+destination inactive masked-off elements in the vector operations. There are
+two kinds of policies described in the vector specification. One is undisturbed.
+It will retain the value they previously held. Another is agnostic. It will
+retain the value they previously held or are overwritten with 1s. It is intended
+for use only inside ``riscv_*.h``.
+  }];
+}
+
 def AVRInterruptDocs : Documentation {
   let Category = DocCatFunction;
   let Heading = "interrupt (AVR)";
   let Content = [{
 Clang supports the GNU style ``__attribute__((interrupt))`` attribute on
 AVR targets. This attribute may be attached to a function definition and instructs
 the backend to generate appropriate function entry/exit code so that it can be used
 directly as an interrupt service routine.

clang/include/clang/Basic/riscv_vector.td

 let HasNoMaskedOverloaded = false,
     ManualCodegen = [{
       IntrinsicTypes = {ResultType, Ops[1]->getType()};
       Ops[0] = Builder.CreateBitCast(Ops[0], ResultType->getPointerTo());
     }],
     ManualCodegenMask= [{
       // Move mask to right before vl.
       std::rotate(Ops.begin(), Ops.begin() + 1, Ops.end() - 1);
-      Ops.push_back(ConstantInt::get(Ops.back()->getType(), TAIL_UNDISTURBED));
+      Ops.push_back(ConstantInt::get(Ops.back()->getType(), PolicyValue));
       IntrinsicTypes = {ResultType, Ops[3]->getType()};
       Ops[1] = Builder.CreateBitCast(Ops[1], ResultType->getPointerTo());
     }] in {
   multiclass RVVVLEBuiltin<list<string> types> {
     let Name = NAME # "_v",
         IRName = "vle",
         IRNameMask ="vle_mask" in {
       foreach type = types in {
let Name = NAME # "_v",
                             Address(NewVL, Align));
         return V;
       }
       }],
       ManualCodegenMask = [{
       {
         // Move mask to right before vl.
         std::rotate(Ops.begin(), Ops.begin() + 1, Ops.end() - 1);
-        Ops.push_back(ConstantInt::get(Ops.back()->getType(), TAIL_UNDISTURBED));
+        Ops.push_back(ConstantInt::get(Ops.back()->getType(), PolicyValue));
         IntrinsicTypes = {ResultType, Ops[4]->getType()};
         Ops[1] = Builder.CreateBitCast(Ops[1], ResultType->getPointerTo());
         Value *NewVL = Ops[2];
         Ops.erase(Ops.begin() + 2);
         llvm::Function *F = CGM.getIntrinsic(ID, IntrinsicTypes);
         llvm::Value *LoadValue = Builder.CreateCall(F, Ops, "");
         llvm::Value *V = Builder.CreateExtractValue(LoadValue, {0});
         // Store new_vl.
let Name = NAME # "_v",
       HasNoMaskedOverloaded = false,
       ManualCodegen = [{
         IntrinsicTypes = {ResultType, Ops[2]->getType()};
         Ops[0] = Builder.CreateBitCast(Ops[0], ResultType->getPointerTo());
       }],
       ManualCodegenMask= [{
         // Move mask to right before vl.
         std::rotate(Ops.begin(), Ops.begin() + 1, Ops.end() - 1);
-        Ops.push_back(ConstantInt::get(Ops.back()->getType(), TAIL_UNDISTURBED));
+        Ops.push_back(ConstantInt::get(Ops.back()->getType(), PolicyValue));
         IntrinsicTypes = {ResultType, Ops[4]->getType()};
         Ops[1] = Builder.CreateBitCast(Ops[1], ResultType->getPointerTo());
       }] in {
     foreach type = types in {
       def : RVVBuiltin<"v", "vPCet", type>;
       if !not(IsFloat<type>.val) then {
         def : RVVBuiltin<"Uv", "UvPCUet", type>;
       }
     }
   }
 }
 
 multiclass RVVIndexedLoad<string op> {
   let ManualCodegen = [{
         IntrinsicTypes = {ResultType, Ops[1]->getType(), Ops[2]->getType()};
         Ops[0] = Builder.CreateBitCast(Ops[0], ResultType->getPointerTo());
       }],
       ManualCodegenMask = [{
         // Move mask to right before vl.
         std::rotate(Ops.begin(), Ops.begin() + 1, Ops.end() - 1);
-        Ops.push_back(ConstantInt::get(Ops.back()->getType(), TAIL_UNDISTURBED));
+        Ops.push_back(ConstantInt::get(Ops.back()->getType(), PolicyValue));
         IntrinsicTypes = {ResultType, Ops[2]->getType(), Ops[4]->getType()};
         Ops[1] = Builder.CreateBitCast(Ops[1], ResultType->getPointerTo());
       }] in {
       foreach type = TypeList in {
         foreach eew_list = EEWList in {
           defvar eew = eew_list[0];
           defvar eew_type = eew_list[1];
           let Name = op # eew # "_v", IRName = op, IRNameMask = op # "_mask" in {
defvar eew = !cond(!eq(type, "c") : "8",
       IntrinsicTypes = {ConvertType(E->getArg(0)->getType()->getPointeeType()),
                         Ops[2 * NF + 2]->getType()};
       SmallVector<llvm::Value*, 12> Operands;
       for (unsigned I = 0; I < NF; ++I)
         Operands.push_back(Ops[NF + I + 1]);
       Operands.push_back(Ops[2 * NF + 1]);
       Operands.push_back(Ops[NF]);
       Operands.push_back(Ops[2 * NF + 2]);
-      Operands.push_back(ConstantInt::get(Ops.back()->getType(), TAIL_UNDISTURBED));
+      Operands.push_back(ConstantInt::get(Ops.back()->getType(), PolicyValue));
       assert(Operands.size() == NF + 4);
       llvm::Function *F = CGM.getIntrinsic(ID, IntrinsicTypes);
       llvm::Value *LoadValue = Builder.CreateCall(F, Operands, "");
       clang::CharUnits Align =
           CGM.getNaturalPointeeTypeAlignment(E->getArg(0)->getType());
       llvm::Value *V;
       for (unsigned I = 0; I < NF; ++I) {
         V = Builder.CreateStore(Builder.CreateExtractValue(LoadValue, {I}),
defvar eew = !cond(!eq(type, "c") : "8",
       IntrinsicTypes = {ConvertType(E->getArg(0)->getType()->getPointeeType()),
                         Ops[2 * NF + 3]->getType()};
       SmallVector<llvm::Value*, 12> Operands;
       for (unsigned I = 0; I < NF; ++I)
         Operands.push_back(Ops[NF + I + 1]);
       Operands.push_back(Ops[2 * NF + 1]);
       Operands.push_back(Ops[NF]);
       Operands.push_back(Ops[2 * NF + 3]);
-      Operands.push_back(ConstantInt::get(Ops.back()->getType(), TAIL_UNDISTURBED));
+      Operands.push_back(ConstantInt::get(Ops.back()->getType(), PolicyValue));
       Value *NewVL = Ops[2 * NF + 2];
       assert(Operands.size() == NF + 4);
       llvm::Function *F = CGM.getIntrinsic(ID, IntrinsicTypes);
       llvm::Value *LoadValue = Builder.CreateCall(F, Operands, "");
       clang::CharUnits Align =
           CGM.getNaturalPointeeTypeAlignment(E->getArg(0)->getType());
       for (unsigned I = 0; I < NF; ++I) {
         Builder.CreateStore(Builder.CreateExtractValue(LoadValue, {I}),
defvar eew = !cond(!eq(type, "c") : "8",
                         Ops[2 * NF + 3]->getType()};
       SmallVector<llvm::Value*, 12> Operands;
       for (unsigned I = 0; I < NF; ++I)
         Operands.push_back(Ops[NF + I + 1]);
       Operands.push_back(Ops[2 * NF + 1]);
       Operands.push_back(Ops[2 * NF + 2]);
       Operands.push_back(Ops[NF]);
       Operands.push_back(Ops[2 * NF + 3]);
-      Operands.push_back(ConstantInt::get(Ops.back()->getType(), TAIL_UNDISTURBED));
+      Operands.push_back(ConstantInt::get(Ops.back()->getType(), PolicyValue));
       assert(Operands.size() == NF + 5);
       llvm::Function *F = CGM.getIntrinsic(ID, IntrinsicTypes);
       llvm::Value *LoadValue = Builder.CreateCall(F, Operands, "");
       clang::CharUnits Align =
           CGM.getNaturalPointeeTypeAlignment(E->getArg(0)->getType());
       llvm::Value *V;
       for (unsigned I = 0; I < NF; ++I) {
         V = Builder.CreateStore(Builder.CreateExtractValue(LoadValue, {I}),
foreach eew_info = EEWList in {
       // intrinsic: (maskedoff0, ..., ptr, index, mask, vl)
       SmallVector<llvm::Value*, 12> Operands;
       for (unsigned I = 0; I < NF; ++I)
         Operands.push_back(Ops[NF + I + 1]);
       Operands.push_back(Ops[2 * NF + 1]);
       Operands.push_back(Ops[2 * NF + 2]);
       Operands.push_back(Ops[NF]);
       Operands.push_back(Ops[2 * NF + 3]);
-      Operands.push_back(ConstantInt::get(Ops.back()->getType(), TAIL_UNDISTURBED));
+      Operands.push_back(ConstantInt::get(Ops.back()->getType(), PolicyValue));
       assert(Operands.size() == NF + 5);
       llvm::Function *F = CGM.getIntrinsic(ID, IntrinsicTypes);
       llvm::Value *LoadValue = Builder.CreateCall(F, Operands, "");
       clang::CharUnits Align =
           CGM.getNaturalPointeeTypeAlignment(E->getArg(0)->getType());
       llvm::Value *V;
       for (unsigned I = 0; I < NF; ++I) {
         V = Builder.CreateStore(Builder.CreateExtractValue(LoadValue, {I}),
let Name = NAME,
                           Ops[1]->getType()};
         Ops.insert(Ops.begin() + 1, llvm::Constant::getNullValue(IntrinsicTypes[1]));
         break;
       }
       }],
       ManualCodegenMask = [{
       {
         std::rotate(Ops.begin(), Ops.begin() + 1, Ops.end() - 1);
-        Ops.push_back(ConstantInt::get(Ops.back()->getType(), TAIL_UNDISTURBED));
+        Ops.push_back(ConstantInt::get(Ops.back()->getType(), PolicyValue));
         // maskedoff, op1, mask, vl
         IntrinsicTypes = {ResultType,
                           cast<llvm::VectorType>(ResultType)->getElementType(),
                           Ops[3]->getType()};
         Ops.insert(Ops.begin() + 2, llvm::Constant::getNullValue(IntrinsicTypes[1]));
         break;
       }
       }] in {
let Name = NAME,
         Ops.insert(Ops.begin() + 1,
                    llvm::Constant::getAllOnesValue(IntrinsicTypes[1]));
         break;
       }
       }],
       ManualCodegenMask = [{
       {
         std::rotate(Ops.begin(), Ops.begin() + 1, Ops.end() - 1);
-        Ops.push_back(ConstantInt::get(Ops.back()->getType(), TAIL_UNDISTURBED));
+        Ops.push_back(ConstantInt::get(Ops.back()->getType(), PolicyValue));
         // maskedoff, op1, mask, vl
         IntrinsicTypes = {ResultType,
                           cast<llvm::VectorType>(ResultType)->getElementType(),
                           Ops[3]->getType()};
         Ops.insert(Ops.begin() + 2,
                    llvm::Constant::getAllOnesValue(IntrinsicTypes[1]));
         break;
       }
let Name = NAME,
                           Ops[0]->getType(), Ops[1]->getType()};
         Ops.insert(Ops.begin() + 1, Ops[0]);
         break;
       }
       }],
       ManualCodegenMask = [{
       {
         std::rotate(Ops.begin(), Ops.begin() + 1, Ops.end() - 1);
-        Ops.push_back(ConstantInt::get(Ops.back()->getType(), TAIL_UNDISTURBED));
+        Ops.push_back(ConstantInt::get(Ops.back()->getType(), PolicyValue));
         // maskedoff, op1, mask, vl
         IntrinsicTypes = {ResultType,
                           Ops[1]->getType(),
                           Ops[3]->getType()};
         Ops.insert(Ops.begin() + 2, Ops[1]);
         break;
       }
       }] in {
let Name = NAME,
                           Ops[1]->getType()};
         Ops.insert(Ops.begin() + 1, llvm::Constant::getNullValue(IntrinsicTypes[2]));
         break;
       }
       }],
       ManualCodegenMask = [{
       {
         std::rotate(Ops.begin(), Ops.begin() + 1, Ops.end() - 1);
-        Ops.push_back(ConstantInt::get(Ops.back()->getType(), TAIL_UNDISTURBED));
+        Ops.push_back(ConstantInt::get(Ops.back()->getType(), PolicyValue));
         // maskedoff, op1, mask, vl
         IntrinsicTypes = {ResultType,
                           Ops[1]->getType(),
                           cast<llvm::VectorType>(Ops[1]->getType())->getElementType(),
                           Ops[3]->getType()};
         Ops.insert(Ops.begin() + 2, llvm::Constant::getNullValue(IntrinsicTypes[2]));
         break;
       }
let Name = NAME,
                           Ops[1]->getType()};
         Ops.insert(Ops.begin() + 1, llvm::Constant::getNullValue(IntrinsicTypes[2]));
         break;
       }
       }],
       ManualCodegenMask = [{
       {
         std::rotate(Ops.begin(), Ops.begin() + 1, Ops.end() - 1);
-        Ops.push_back(ConstantInt::get(Ops.back()->getType(), TAIL_UNDISTURBED));
+        Ops.push_back(ConstantInt::get(Ops.back()->getType(), PolicyValue));
         // maskedoff, op1, mask, vl
         IntrinsicTypes = {ResultType,
                           Ops[1]->getType(),
                           Ops[3]->getType(),
                           Ops[3]->getType()};
         Ops.insert(Ops.begin() + 2, llvm::Constant::getNullValue(IntrinsicTypes[2]));
         break;
       }

clang/lib/CodeGen/CGBuiltin.cpp

Value *CodeGenFunction::EmitRISCVBuiltinExpr(unsigned BuiltinID,
   SmallVector<Value *, 4> Ops;
   llvm::Type *ResultType = ConvertType(E->getType());
 
   for (unsigned i = 0, e = E->getNumArgs(); i != e; i++)
     Ops.push_back(EmitScalarExpr(E->getArg(i)));
 
   Intrinsic::ID ID = Intrinsic::not_intrinsic;
   unsigned NF = 1;
-  constexpr unsigned TAIL_UNDISTURBED = 0;
+  constexpr unsigned TAIL_AGNOSTIC = 0b01;
+  constexpr unsigned MASK_AGNOSTIC = 0b10;
+  auto *PolicyAttr = E->getCalleeDecl()->getAttr<RISCVVPolicyAttr>();
+  /* Default is tail undisturbed and mask undisturbed. */
+  uint64_t PolicyValue = 0;
+
+  if (PolicyAttr) {
+    switch (PolicyAttr->getPolicy()) {
+    default:
+      break;
+    case RISCVVPolicyAttr::TAMU:
+      PolicyValue = TAIL_AGNOSTIC;
+      break;
+    case RISCVVPolicyAttr::TUMA:
+      PolicyValue = MASK_AGNOSTIC;
+      break;
+    case RISCVVPolicyAttr::TAMA:
+      PolicyValue = MASK_AGNOSTIC | TAIL_AGNOSTIC;
+      break;
+    }
+  }
 
   // Required for overloaded intrinsics.
   llvm::SmallVector<llvm::Type *, 2> IntrinsicTypes;
   switch (BuiltinID) {
   default: llvm_unreachable("unexpected builtin ID");
   case RISCV::BI__builtin_riscv_orc_b_32:
   case RISCV::BI__builtin_riscv_orc_b_64:
   case RISCV::BI__builtin_riscv_clmul:

clang/lib/Sema/SemaDeclAttr.cpp

if ((IsAArch64 && !ArmSveAliasValid(S.Context, BuiltinID, AliasName)) ||
       (!IsAArch64 && !IsARM && !IsRISCV)) {
     S.Diag(AL.getLoc(), diag::err_attribute_builtin_alias) << AL;
     return;
   }
 
   D->addAttr(::new (S.Context) BuiltinAliasAttr(S.Context, AL, Ident));
 }
 
+static void handleRISCVVPolicyAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
+  if (!AL.isArgIdent(0)) {
+    S.Diag(AL.getLoc(), diag::err_attribute_argument_n_type)
+        << AL << 0 << AANT_ArgumentIdentifier;
+    return;
+  }
+
+  RISCVVPolicyAttr::PolicyType Policy;
+  IdentifierInfo *II = AL.getArgAsIdent(0)->Ident;
+  if (!RISCVVPolicyAttr::ConvertStrToPolicyType(II->getName(), Policy)) {
+    S.Diag(AL.getLoc(), diag::warn_attribute_type_not_supported) << AL << II;
+    return;
+  }
+
+  D->addAttr(::new (S.Context) RISCVVPolicyAttr(S.Context, AL, Policy));
+}
+
 //===----------------------------------------------------------------------===//
 // Checker-specific attribute handlers.
 //===----------------------------------------------------------------------===//
 static bool isValidSubjectOfNSReturnsRetainedAttribute(QualType QT) {
   return QT->isDependentType() || QT->isObjCRetainableType();
 }
 
 static bool isValidSubjectOfNSAttribute(QualType QT) {
static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
 
   case ParsedAttr::AT_BuiltinAlias:
     handleBuiltinAliasAttr(S, D, AL);
     break;
 
   case ParsedAttr::AT_UsingIfExists:
     handleSimpleAttribute<UsingIfExistsAttr>(S, D, AL);
     break;
+
+  case ParsedAttr::AT_RISCVVPolicy:
+    handleRISCVVPolicyAttr(S, D, AL);
+    break;
   }
 }
 
 /// ProcessDeclAttributeList - Apply all the decl attributes in the specified
 /// attribute list to the specified decl, ignoring any type attributes.
 void Sema::ProcessDeclAttributeList(Scope *S, Decl *D,
                                     const ParsedAttributesView &AttrList,
                                     bool IncludeCXX11Attributes) {

clang/test/CodeGen/RISCV/rvv-intrinsics/vadd-policy.c

+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
+// REQUIRES: riscv-registered-target
+// RUN: %clang_cc1 -triple riscv64 -target-feature +f -target-feature +d -target-feature +experimental-v \
+// RUN:   -target-feature +experimental-zfh -disable-O0-optnone -emit-llvm %s -o - | opt -S -mem2reg | FileCheck --check-prefix=CHECK-RV64 %s
+
+#include <riscv_vector.h>
+
+// CHECK-RV64-LABEL: @test_vadd_vv_i8m1_tama(
+// CHECK-RV64-NEXT:  entry:
+// CHECK-RV64-NEXT:    [[TMP0:%.*]] = call <vscale x 8 x i8> @llvm.riscv.vadd.mask.nxv8i8.nxv8i8.i64(<vscale x 8 x i8> [[MASKEDOFF:%.*]], <vscale x 8 x i8> [[OP1:%.*]], <vscale x 8 x i8> [[OP2:%.*]], <vscale x 8 x i1> [[MASK:%.*]], i64 [[VL:%.*]], i64 3)
+// CHECK-RV64-NEXT:    ret <vscale x 8 x i8> [[TMP0]]
+//
+vint8m1_t test_vadd_vv_i8m1_tama(vbool8_t mask, vint8m1_t maskedoff, vint8m1_t op1, vint8m1_t op2, size_t vl) {
+  return vadd_vv_i8m1_tama(mask, maskedoff, op1, op2, vl);
+}
+
+// CHECK-RV64-LABEL: @test_vadd_vv_i8m1_tamu(
+// CHECK-RV64-NEXT:  entry:
+// CHECK-RV64-NEXT:    [[TMP0:%.*]] = call <vscale x 8 x i8> @llvm.riscv.vadd.mask.nxv8i8.nxv8i8.i64(<vscale x 8 x i8> [[MASKEDOFF:%.*]], <vscale x 8 x i8> [[OP1:%.*]], <vscale x 8 x i8> [[OP2:%.*]], <vscale x 8 x i1> [[MASK:%.*]], i64 [[VL:%.*]], i64 1)
+// CHECK-RV64-NEXT:    ret <vscale x 8 x i8> [[TMP0]]
+//
+vint8m1_t test_vadd_vv_i8m1_tamu(vbool8_t mask, vint8m1_t maskedoff, vint8m1_t op1, vint8m1_t op2, size_t vl) {
+  return vadd_vv_i8m1_tamu(mask, maskedoff, op1, op2, vl);
+}
+
+// CHECK-RV64-LABEL: @test_vadd_vv_i8m1_tuma(
+// CHECK-RV64-NEXT:  entry:
+// CHECK-RV64-NEXT:    [[TMP0:%.*]] = call <vscale x 8 x i8> @llvm.riscv.vadd.mask.nxv8i8.nxv8i8.i64(<vscale x 8 x i8> [[MASKEDOFF:%.*]], <vscale x 8 x i8> [[OP1:%.*]], <vscale x 8 x i8> [[OP2:%.*]], <vscale x 8 x i1> [[MASK:%.*]], i64 [[VL:%.*]], i64 2)
+// CHECK-RV64-NEXT:    ret <vscale x 8 x i8> [[TMP0]]
+//
+vint8m1_t test_vadd_vv_i8m1_tuma(vbool8_t mask, vint8m1_t maskedoff, vint8m1_t op1, vint8m1_t op2, size_t vl) {
+  return vadd_vv_i8m1_tuma(mask, maskedoff, op1, op2, vl);
+}
+
+// CHECK-RV64-LABEL: @test_vadd_vv_i8m1_tumu(
+// CHECK-RV64-NEXT:  entry:
+// CHECK-RV64-NEXT:    [[TMP0:%.*]] = call <vscale x 8 x i8> @llvm.riscv.vadd.mask.nxv8i8.nxv8i8.i64(<vscale x 8 x i8> [[MASKEDOFF:%.*]], <vscale x 8 x i8> [[OP1:%.*]], <vscale x 8 x i8> [[OP2:%.*]], <vscale x 8 x i1> [[MASK:%.*]], i64 [[VL:%.*]], i64 0)
+// CHECK-RV64-NEXT:    ret <vscale x 8 x i8> [[TMP0]]
+//
+vint8m1_t test_vadd_vv_i8m1_tumu(vbool8_t mask, vint8m1_t maskedoff, vint8m1_t op1, vint8m1_t op2, size_t vl) {
+  return vadd_vv_i8m1_tumu(mask, maskedoff, op1, op2, vl);
+}

clang/test/Misc/pragma-attribute-supported-attributes-list.test

 // CHECK-NEXT: OpenCLNoSVM (SubjectMatchRule_variable)
 // CHECK-NEXT: OptimizeNone (SubjectMatchRule_function, SubjectMatchRule_objc_method)
 // CHECK-NEXT: Overloadable (SubjectMatchRule_function)
 // CHECK-NEXT: Owner (SubjectMatchRule_record_not_is_union)
 // CHECK-NEXT: ParamTypestate (SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: PassObjectSize (SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: PatchableFunctionEntry (SubjectMatchRule_function, SubjectMatchRule_objc_method)
 // CHECK-NEXT: Pointer (SubjectMatchRule_record_not_is_union)
+// CHECK-NEXT: RISCVVPolicy (SubjectMatchRule_function)
 // CHECK-NEXT: ReleaseHandle (SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: RenderScriptKernel (SubjectMatchRule_function)
 // CHECK-NEXT: ReqdWorkGroupSize (SubjectMatchRule_function)
 // CHECK-NEXT: Restrict (SubjectMatchRule_function)
 // CHECK-NEXT: ReturnTypestate (SubjectMatchRule_function, SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: ReturnsNonNull (SubjectMatchRule_objc_method, SubjectMatchRule_function)
 // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function)
 // CHECK-NEXT: ScopedLockable (SubjectMatchRule_record)

clang/utils/TableGen/RISCVVEmitter.cpp

public:
 
   // Emit the code block for switch body in EmitRISCVBuiltinExpr, it should
   // init the RVVIntrinsic ID and IntrinsicTypes.
   void emitCodeGenSwitchBody(raw_ostream &o) const;
 
   // Emit the macros for mapping C/C++ intrinsic function to builtin functions.
   void emitIntrinsicFuncDef(raw_ostream &o) const;
 
+  // Emit the declarations for mapping C/C++ intrinsic function to builtin
+  // functions.
+  void emitIntrinsicWithPolicyFuncDef(raw_ostream &o) const;
+
   // Emit the mangled function definition.
   void emitMangledFuncDef(raw_ostream &o) const;
 };
 
 class RVVEmitter {
 private:
   RecordKeeper &Records;
   std::string HeaderCode;
if (hasManualCodegen()) {
     OS << ManualCodegen;
     OS << "break;\n";
     return;
   }
 
   if (isMask()) {
     if (hasVL()) {
       OS << "  std::rotate(Ops.begin(), Ops.begin() + 1, Ops.end() - 1);\n";
-      if (hasPolicy())
-        OS << "  Ops.push_back(ConstantInt::get(Ops.back()->getType(),"
-                               " TAIL_UNDISTURBED));\n";
+      if (hasPolicy()) {
+        OS << "  Ops.push_back(ConstantInt::get(Ops.back()->getType(), "
+              "PolicyValue));\n";
+      }
     } else {
       OS << "  std::rotate(Ops.begin(), Ops.begin() + 1, Ops.end());\n";
     }
   }
 
   OS << "  IntrinsicTypes = {";
   ListSeparator LS;
   for (const auto &Idx : IntrinsicTypes) {
void RVVIntrinsic::emitCodeGenSwitchBody(raw_ostream &OS) const {
   OS << "  break;\n";
 }
 
 void RVVIntrinsic::emitIntrinsicFuncDef(raw_ostream &OS) const {
   OS << "__attribute__((__clang_builtin_alias__(";
   OS << "__builtin_rvv_" << getBuiltinName() << ")))\n";
   OS << OutputType->getTypeStr() << " " << getName() << "(";
   // Emit function arguments
-  if (!InputTypes.empty()) {
-    ListSeparator LS;
-    for (unsigned i = 0; i < InputTypes.size(); ++i)
-      OS << LS << InputTypes[i]->getTypeStr();
-  }
+  OS << ");\n";
+
+void RVVIntrinsic::emitIntrinsicWithPolicyFuncDef(raw_ostream &OS) const {
+  if (!isMask())
+    return;
+
+  static const char *const PolicySuffix[] = {"tumu", "tamu", "tuma", "tama"};
+
+  for (auto Suffix : PolicySuffix) {
+    OS << "__rvv_ai ";
+    OS << "__attribute__((__clang_builtin_alias__(";
+    OS << "__builtin_rvv_" << getBuiltinName() << ")))\n";
+    OS << "__attribute__((rvv_policy(" << Suffix << ")))\n";
+    StringRef IntrinsicName = getName().substr(0, getName().size() - 2);
+    OS << OutputType->getTypeStr() << " " << IntrinsicName << "_" << Suffix
+       << "(";
+    // Emit function arguments
+    ListSeparator LS;
+    for (unsigned i = 0; i < InputTypes.size(); ++i)
+      OS << LS << InputTypes[i]->getTypeStr();
     OS << ");\n";
   }
+}
 
 void RVVIntrinsic::emitMangledFuncDef(raw_ostream &OS) const {
   OS << "__attribute__((__clang_builtin_alias__(";
   OS << "__builtin_rvv_" << getBuiltinName() << ")))\n";
   OS << OutputType->getTypeStr() << " " << getMangledName() << "(";
   // Emit function arguments
-  if (!InputTypes.empty()) {
-    ListSeparator LS;
-    for (unsigned i = 0; i < InputTypes.size(); ++i)
-      OS << LS << InputTypes[i]->getTypeStr();
-  }
   OS << ");\n";
 }
 
 //===----------------------------------------------------------------------===//
 // RVVEmitter implementation
 //===----------------------------------------------------------------------===//
 void RVVEmitter::createHeader(raw_ostream &OS) {
 
   OS << "/*===---- riscv_vector.h - RISC-V V-extension RVVIntrinsics "
         "-------------------===\n"
         " *\n"
         " *\n"
         " * Part of the LLVM Project, under the Apache License v2.0 with LLVM "
         "Exceptions.\n"
         " * See https://llvm.org/LICENSE.txt for license information.\n"
         " * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception\n"
void RVVEmitter::createHeader(raw_ostream &OS) {
   OS << "#define __rvv_ai static __inline__\n";
 
   // Print intrinsic functions with macro
   emitArchMacroAndBody(Defs, OS, [](raw_ostream &OS, const RVVIntrinsic &Inst) {
     OS << "__rvv_ai ";
     Inst.emitIntrinsicFuncDef(OS);
   });
 
+  emitArchMacroAndBody(Defs, OS, [](raw_ostream &OS, const RVVIntrinsic &Inst) {
+    Inst.emitIntrinsicWithPolicyFuncDef(OS);
+  });
+
   OS << "#undef __rvv_ai\n\n";
 
   OS << "#define __riscv_v_intrinsic_overloading 1\n";
 
   // Print Overloaded APIs
   OS << "#define __rvv_aio static __inline__ "
         "__attribute__((__overloadable__))\n";
 

llvm/lib/Target/RISCV/RISCVInsertVSETVLI.cpp

 
 } // end anonymous namespace
 
 char RISCVInsertVSETVLI::ID = 0;
 
 INITIALIZE_PASS(RISCVInsertVSETVLI, DEBUG_TYPE, RISCV_INSERT_VSETVLI_NAME,
                 false, false)
 
-static MachineInstr *elideCopies(MachineInstr *MI,
-                                 const MachineRegisterInfo *MRI) {
-  while (true) {
-    if (!MI->isFullCopy())
-      return MI;
-    if (!Register::isVirtualRegister(MI->getOperand(1).getReg()))
-      return nullptr;
-    MI = MRI->getVRegDef(MI->getOperand(1).getReg());
-    if (!MI)
-      return nullptr;
-  }
-}
-
 static VSETVLIInfo computeInfoForInstr(const MachineInstr &MI, uint64_t TSFlags,
                                        const MachineRegisterInfo *MRI) {
   VSETVLIInfo InstrInfo;
   unsigned NumOperands = MI.getNumExplicitOperands();
   bool HasPolicy = RISCVII::hasVecPolicyOp(TSFlags);
-
-  // Default to tail agnostic unless the destination is tied to a source.
-  // Unless the source is undef. In that case the user would have some control
-  // over the tail values. Some pseudo instructions force a tail agnostic policy
-  // despite having a tied def.
-  bool ForceTailAgnostic = RISCVII::doesForceTailAgnostic(TSFlags);
   bool TailAgnostic = true;
+  bool MaskAgnostic = false;
   // If the instruction has policy argument, use the argument.
   if (HasPolicy) {
     const MachineOperand &Op = MI.getOperand(MI.getNumExplicitOperands() - 1);
     TailAgnostic = Op.getImm() & 0x1;
-  }
-
-  unsigned UseOpIdx;
-  if (!(ForceTailAgnostic || (HasPolicy && TailAgnostic)) &&
-      MI.isRegTiedToUseOperand(0, &UseOpIdx)) {
-    TailAgnostic = false;
-    // If the tied operand is an IMPLICIT_DEF we can keep TailAgnostic.
-    const MachineOperand &UseMO = MI.getOperand(UseOpIdx);
-    MachineInstr *UseMI = MRI->getVRegDef(UseMO.getReg());
-    if (UseMI) {
-      UseMI = elideCopies(UseMI, MRI);
-      if (UseMI && UseMI->isImplicitDef())
-        TailAgnostic = true;
-    }
+    MaskAgnostic = Op.getImm() & 0x2;
   }
 
   // Remove the tail policy so we can find the SEW and VL.
   if (HasPolicy)
     --NumOperands;
 
   RISCVII::VLMUL VLMul = RISCVII::getLMul(TSFlags);
 
if (VLOp.isImm()) {
         InstrInfo.setAVLReg(RISCV::X0);
       else
         InstrInfo.setAVLImm(Imm);
     } else {
       InstrInfo.setAVLReg(VLOp.getReg());
     }
   } else
     InstrInfo.setAVLReg(RISCV::NoRegister);
-  InstrInfo.setVTYPE(VLMul, SEW, /*TailAgnostic*/ TailAgnostic,
-                     /*MaskAgnostic*/ false, MaskRegOp, StoreOp);
+  InstrInfo.setVTYPE(VLMul, SEW, TailAgnostic, MaskAgnostic, MaskRegOp,
+                     StoreOp);
 
   return InstrInfo;
 }
 
 void RISCVInsertVSETVLI::insertVSETVLI(MachineBasicBlock &MBB, MachineInstr &MI,
                                        const VSETVLIInfo &Info,
                                        const VSETVLIInfo &PrevInfo) {
   DebugLoc DL = MI.getDebugLoc();