0

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);
  }
 } 
Pawan Nirpal
  • 565
  • 1
  • 12
  • 29

0 Answers0