diff options
Diffstat (limited to 'llvm/lib/IR/Verifier.cpp')
-rw-r--r-- | llvm/lib/IR/Verifier.cpp | 63 |
1 files changed, 63 insertions, 0 deletions
diff --git a/llvm/lib/IR/Verifier.cpp b/llvm/lib/IR/Verifier.cpp index 9bd573e..3ff9895 100644 --- a/llvm/lib/IR/Verifier.cpp +++ b/llvm/lib/IR/Verifier.cpp @@ -2979,6 +2979,16 @@ void Verifier::visitFunction(const Function &F) { "perfect forwarding!", &F); break; + case CallingConv::AMDGPU_Gfx_WholeWave: + Check(!F.arg_empty() && F.arg_begin()->getType()->isIntegerTy(1), + "Calling convention requires first argument to be i1", &F); + Check(!F.arg_begin()->hasInRegAttr(), + "Calling convention requires first argument to not be inreg", &F); + Check(!F.isVarArg(), + "Calling convention does not support varargs or " + "perfect forwarding!", + &F); + break; } // Check that the argument values match the function type for this function... @@ -6658,6 +6668,54 @@ void Verifier::visitIntrinsicCall(Intrinsic::ID ID, CallBase &Call) { "invalid vector type for format", &Call, Src1, Call.getArgOperand(5)); break; } + case Intrinsic::amdgcn_wmma_f32_16x16x128_f8f6f4: { + Value *Src0 = Call.getArgOperand(1); + Value *Src1 = Call.getArgOperand(3); + + unsigned FmtA = cast<ConstantInt>(Call.getArgOperand(0))->getZExtValue(); + unsigned FmtB = cast<ConstantInt>(Call.getArgOperand(2))->getZExtValue(); + Check(FmtA <= 4, "invalid value for matrix format", Call, + Call.getArgOperand(0)); + Check(FmtB <= 4, "invalid value for matrix format", Call, + Call.getArgOperand(2)); + + // AMDGPU::MatrixFMT values + auto getFormatNumRegs = [](unsigned FormatVal) { + switch (FormatVal) { + case 0: + case 1: + return 16u; + case 2: + case 3: + return 12u; + case 4: + return 8u; + default: + llvm_unreachable("invalid format value"); + } + }; + + auto isValidSrcASrcBVector = [](FixedVectorType *Ty) { + if (!Ty || !Ty->getElementType()->isIntegerTy(32)) + return false; + unsigned NumElts = Ty->getNumElements(); + return NumElts == 16 || NumElts == 12 || NumElts == 8; + }; + + auto *Src0Ty = dyn_cast<FixedVectorType>(Src0->getType()); + auto *Src1Ty = dyn_cast<FixedVectorType>(Src1->getType()); + Check(isValidSrcASrcBVector(Src0Ty), + "operand 1 must be 8, 12 or 16 element i32 vector", &Call, Src0); + Check(isValidSrcASrcBVector(Src1Ty), + "operand 3 must be 8, 12 or 16 element i32 vector", &Call, Src1); + + // Permit excess registers for the format. + Check(Src0Ty->getNumElements() >= getFormatNumRegs(FmtA), + "invalid vector type for format", &Call, Src0, Call.getArgOperand(0)); + Check(Src1Ty->getNumElements() >= getFormatNumRegs(FmtB), + "invalid vector type for format", &Call, Src1, Call.getArgOperand(2)); + break; + } case Intrinsic::nvvm_setmaxnreg_inc_sync_aligned_u32: case Intrinsic::nvvm_setmaxnreg_dec_sync_aligned_u32: { Value *V = Call.getArgOperand(0); @@ -6710,6 +6768,11 @@ void Verifier::visitIntrinsicCall(Intrinsic::ID ID, CallBase &Call) { "llvm.threadlocal.address operand isThreadLocal() must be true"); break; } + case Intrinsic::lifetime_start: + case Intrinsic::lifetime_end: + Check(isa<AllocaInst>(Call.getArgOperand(1)), + "llvm.lifetime.start/end can only be used on alloca", &Call); + break; }; // Verify that there aren't any unmediated control transfers between funclets. |