Hi I am writing a transformation pass to check where on some conditional call instruction I would like to check If fast option is set for that call, instruction and disable it. For example on below call inst, I would like to disable the fast option.
%call3 = call fast fastcc noundef <8 x float> @_ZL13_mm256_mul_psDv8_fS_(<8 x float> noundef %0, <8 x float> noundef %power.0)
Below is my pass, Don’t mind If I miss something trivial, I’m at a very early stage in compiler dev.
Header file contents :
#include "llvm/IR/PassManager.h"
namespace llvm {
class RemoveFastFlagsToPreventFmaPass: public PassInfoMixin<RemoveFastFlagsToPreventFmaPass> {
public:
PreservedAnalyses run(Module &M, ModuleAnalysisManager &AM);
};
} // namespace llvm
#include "llvm/Transforms/Scalar/RemoveFastFlagsToPreventFma.h"
#include "llvm/IR/Instructions.h"
#define Status(b) outs()<<(b?"Yes\n":"No\n")
using namespace llvm;
using namespace std;
bool fastflags_for_intrfma = false;
bool IsArithmeticIntrinsicCall(const Instruction &i){
return 1;
}
PreservedAnalyses RemoveFastFlagsToPreventFmaPass::run(Module &M, ModuleAnalysisManager &AM){
for(Function &F : M){
for(BasicBlock &b : F){
for(Instruction &i : b){
CallInst *call = dyn_cast<CallInst>(&i);
if(fastflags_for_intrfma==false and (call!=nullptr) and IsArithmeticIntrinsicCall(i)){ // Ignore the working of IsArithmeticIntrinsicCall I wrote it for completeness sake.
if(i.isFast()) i.setFast(false);
}
}
}
}
for(Function &F : M){
for(BasicBlock &b : F){
for(Instruction &i : b){
CallInst *call = dyn_cast<CallInst>(&i);
if(fastflags_for_intrfma==false and (call!=nullptr) and IsArithmeticIntrinsicCall(i)){
Status(i.isFast());
}
}
}
}
return PreservedAnalyses::all();
}
The Assertion failure :
opt: llvm-project/llvm/lib/IR/Instruction.cpp:250: bool llvm::Instruction::isFast() const: Assertion `isa<FPMathOperator>(this) && "getting fast-math flag on invalid op"' failed.
The langrefLLVM LangRef suggests that call instruction may use fast flag, but how should I be certain of it, I am assuming this assertion failure is caused because of incompatibility.
Edit :
I figured out what went wrong. Thanks for looking into this. Below is the fix.
for(Instruction &i : b){
auto FMAop = dyn_cast<FPMathOperator>(&i);
auto callinst = dyn_cast<CallInst>(&i);
if(FMAop!=nullptr and callinst!=nullptr and FMAop->isFast() and IsArithmeticIntrinsicCall(i)){
i.setFast(false);
}
}