HsiangKai updated this revision to Diff 395365.
HsiangKai added a comment.
Herald added subscribers: llvm-commits, hiraditya.
Herald added a project: LLVM.
In riscv-insert-vsetvli, use the policy argument. No use implicit-def maskedoff
to adjust the setting.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D112534/new/
https://reviews.llvm.org/D112534
Files:
clang/include/clang/Basic/Attr.td
clang/include/clang/Basic/AttrDocs.td
clang/include/clang/Basic/riscv_vector.td
clang/lib/CodeGen/CGBuiltin.cpp
clang/lib/Sema/SemaDeclAttr.cpp
clang/test/CodeGen/RISCV/rvv-intrinsics/vadd-policy.c
clang/test/Misc/pragma-attribute-supported-attributes-list.test
clang/utils/TableGen/RISCVVEmitter.cpp
llvm/lib/Target/RISCV/RISCVInsertVSETVLI.cpp
Index: llvm/lib/Target/RISCV/RISCVInsertVSETVLI.cpp
===================================================================
--- llvm/lib/Target/RISCV/RISCVInsertVSETVLI.cpp
+++ llvm/lib/Target/RISCV/RISCVInsertVSETVLI.cpp
@@ -401,49 +401,18 @@
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.
@@ -476,8 +445,8 @@
}
} 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;
}
Index: clang/utils/TableGen/RISCVVEmitter.cpp
===================================================================
--- clang/utils/TableGen/RISCVVEmitter.cpp
+++ clang/utils/TableGen/RISCVVEmitter.cpp
@@ -204,6 +204,10 @@
// 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;
};
@@ -835,9 +839,10 @@
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";
}
@@ -865,12 +870,32 @@
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";
}
- OS << ");\n";
}
void RVVIntrinsic::emitMangledFuncDef(raw_ostream &OS) const {
@@ -878,11 +903,9 @@
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();
- }
+ ListSeparator LS;
+ for (unsigned i = 0; i < InputTypes.size(); ++i)
+ OS << LS << InputTypes[i]->getTypeStr();
OS << ");\n";
}
@@ -989,6 +1012,10 @@
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";
Index: clang/test/Misc/pragma-attribute-supported-attributes-list.test
===================================================================
--- clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -148,6 +148,7 @@
// 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)
Index: clang/test/CodeGen/RISCV/rvv-intrinsics/vadd-policy.c
===================================================================
--- /dev/null
+++ clang/test/CodeGen/RISCV/rvv-intrinsics/vadd-policy.c
@@ -0,0 +1,42 @@
+// 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);
+}
Index: clang/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- clang/lib/Sema/SemaDeclAttr.cpp
+++ clang/lib/Sema/SemaDeclAttr.cpp
@@ -5580,6 +5580,23 @@
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.
//===----------------------------------------------------------------------===//
@@ -8737,6 +8754,10 @@
case ParsedAttr::AT_UsingIfExists:
handleSimpleAttribute<UsingIfExistsAttr>(S, D, AL);
break;
+
+ case ParsedAttr::AT_RISCVVPolicy:
+ handleRISCVVPolicyAttr(S, D, AL);
+ break;
}
}
Index: clang/lib/CodeGen/CGBuiltin.cpp
===================================================================
--- clang/lib/CodeGen/CGBuiltin.cpp
+++ clang/lib/CodeGen/CGBuiltin.cpp
@@ -18809,7 +18809,27 @@
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;
Index: clang/include/clang/Basic/riscv_vector.td
===================================================================
--- clang/include/clang/Basic/riscv_vector.td
+++ clang/include/clang/Basic/riscv_vector.td
@@ -599,7 +599,7 @@
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 {
@@ -643,7 +643,7 @@
{
// 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];
@@ -681,7 +681,7 @@
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 {
@@ -702,7 +702,7 @@
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 {
@@ -870,7 +870,7 @@
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, "");
@@ -943,7 +943,7 @@
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);
@@ -1017,7 +1017,7 @@
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, "");
@@ -1084,7 +1084,7 @@
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, "");
@@ -1274,7 +1274,7 @@
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(),
@@ -1305,7 +1305,7 @@
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(),
@@ -1353,7 +1353,7 @@
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(),
@@ -1386,7 +1386,7 @@
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(),
@@ -1422,7 +1422,7 @@
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(),
Index: clang/include/clang/Basic/AttrDocs.td
===================================================================
--- clang/include/clang/Basic/AttrDocs.td
+++ clang/include/clang/Basic/AttrDocs.td
@@ -2164,6 +2164,19 @@
}];
}
+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)";
Index: clang/include/clang/Basic/Attr.td
===================================================================
--- clang/include/clang/Basic/Attr.td
+++ clang/include/clang/Basic/Attr.td
@@ -1783,6 +1783,15 @@
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.
//
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
