diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1783,6 +1783,15 @@ let Documentation = [RISCVInterruptDocs]; } +def RISCVVPolicy : InheritableAttr, TargetSpecificAttr { + 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. // diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td --- a/clang/include/clang/Basic/AttrDocs.td +++ b/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)"; diff --git a/clang/include/clang/Basic/riscv_vector.td b/clang/include/clang/Basic/riscv_vector.td --- a/clang/include/clang/Basic/riscv_vector.td +++ b/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(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(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(), diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/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(); + /* 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 IntrinsicTypes; diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/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(S, D, AL); break; + + case ParsedAttr::AT_RISCVVPolicy: + handleRISCVVPolicyAttr(S, D, AL); + break; } } diff --git a/clang/test/CodeGen/RISCV/rvv-intrinsics/vadd-policy.c b/clang/test/CodeGen/RISCV/rvv-intrinsics/vadd-policy.c new file mode 100644 --- /dev/null +++ b/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 + +// CHECK-RV64-LABEL: @test_vadd_vv_i8m1_tama( +// CHECK-RV64-NEXT: entry: +// CHECK-RV64-NEXT: [[TMP0:%.*]] = call @llvm.riscv.vadd.mask.nxv8i8.nxv8i8.i64( [[MASKEDOFF:%.*]], [[OP1:%.*]], [[OP2:%.*]], [[MASK:%.*]], i64 [[VL:%.*]], i64 3) +// CHECK-RV64-NEXT: ret [[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 @llvm.riscv.vadd.mask.nxv8i8.nxv8i8.i64( [[MASKEDOFF:%.*]], [[OP1:%.*]], [[OP2:%.*]], [[MASK:%.*]], i64 [[VL:%.*]], i64 1) +// CHECK-RV64-NEXT: ret [[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 @llvm.riscv.vadd.mask.nxv8i8.nxv8i8.i64( [[MASKEDOFF:%.*]], [[OP1:%.*]], [[OP2:%.*]], [[MASK:%.*]], i64 [[VL:%.*]], i64 2) +// CHECK-RV64-NEXT: ret [[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 @llvm.riscv.vadd.mask.nxv8i8.nxv8i8.i64( [[MASKEDOFF:%.*]], [[OP1:%.*]], [[OP2:%.*]], [[MASK:%.*]], i64 [[VL:%.*]], i64 0) +// CHECK-RV64-NEXT: ret [[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); +} diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/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) diff --git a/clang/utils/TableGen/RISCVVEmitter.cpp b/clang/utils/TableGen/RISCVVEmitter.cpp --- a/clang/utils/TableGen/RISCVVEmitter.cpp +++ b/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"; diff --git a/llvm/lib/Target/RISCV/RISCVInsertVSETVLI.cpp b/llvm/lib/Target/RISCV/RISCVInsertVSETVLI.cpp --- a/llvm/lib/Target/RISCV/RISCVInsertVSETVLI.cpp +++ b/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; }