[Remarks] Emit optimization remarks for atomics generating CAS loop

[gandhi21299]

Summary

Implements ORE in AtomicExpandPass to report atomics generating a compare and swap loop.

[rampitec]

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

12126

Do not pass it there. Turn reportAtomicExpand into lambda.

[gandhi21299]

• turned reportAtomicExpand(...) into a lambda as requested

[gandhi21299]

refreshing patch

[gandhi21299]

rebase

[gandhi21299]

• emit CAS loop remark only if the instruction is of floating point type

[arsenm]

llvm/lib/CodeGen/AtomicExpandPass.cpp

626–632

This is an AMDGPU specific message/restriction. The floating point operation isn't relevant, and you don't even know the specific reason at this point.

[gandhi21299]

• floating point check is not required for cas loop check, getting rid of it

[gandhi21299]

• restricting remarks emission on AMDGPU targets only

[gandhi21299]

@rampitec comments?

[rampitec]

llvm/lib/CodeGen/AtomicExpandPass.cpp

628

Still the same problem.

[gandhi21299]

requested change

[gandhi21299]

@rampitec besides the remarks, am I missing anything else in the patch?

[rampitec]

In D106891#2938128, @gandhi21299 wrote:

@rampitec besides the remarks, am I missing anything else in the patch?

You should not use AMD specific code in the common code.

[gandhi21299]

• reverted the patch to the previous one

[gandhi21299]

• eliminated unsafe hardware remarks in SIISelLowering.cpp
• updated cas loop remark and corresponding tests

[gandhi21299]

removed AMDGPU check

[rampitec]

In D106891#2940411, @gandhi21299 wrote:

• eliminated unsafe hardware remarks in SIISelLowering.cpp

Most of this patch is not needed now. You do not need to pass ORE to targets, it is a part of the next patch.

clang/test/CodeGenCUDA/fp-atomics-optremarks.cu

10

Need tests for all scopes.

llvm/lib/CodeGen/AtomicExpandPass.cpp

587

I do not see why do you need this function and all its arguments now. You can just call ORE->emit() directly.

631

That does not help with target defined scope names, such as our "one-as" for example.

llvm/test/CodeGen/AMDGPU/fp-atomics-remarks-gfx90a.ll

4

You need to write tests for all scopes.

[rampitec]

You also need to retitle it now, it is not about AMDGPU and not about FP.

llvm/lib/CodeGen/AtomicExpandPass.cpp

629

Need to name the operation.

[gandhi21299]

llvm/lib/CodeGen/AtomicExpandPass.cpp

631

How can I get target defined scope names?

[rampitec]

llvm/lib/CodeGen/AtomicExpandPass.cpp

631

It is right on the instruction:

%result = atomicrmw fadd float addrspace(1)* %ptr, float 4.0 syncscope("one-as") seq_cst

[gandhi21299]

llvm/lib/CodeGen/AtomicExpandPass.cpp

631

Sorry, I meant from the LLVM API.

[rampitec]

llvm/lib/CodeGen/AtomicExpandPass.cpp

631

LLVMContext::getSyncScopeNames()

[gandhi21299]

llvm/lib/CodeGen/AtomicExpandPass.cpp

631

I think that gives me all sync scopes available for the target. If not, which sync scope in the vector corresponds to the instruction I am dealing with?

[rampitec]

llvm/lib/CodeGen/AtomicExpandPass.cpp

631

https://llvm.org/doxygen/MachineOperand_8cpp_source.html#l00474

[gandhi21299]

requested changes from reviewer

• added memory scope tests and updated remarks and tests accordingly
• still working on clang/test/CodeGenCUDA/fp-atomics-optremarks.cu and clang/test/CodeGenOpenCL/atomics-remarks-gfx90a.cl

[rampitec]

llvm/lib/CodeGen/AtomicExpandPass.cpp

626

Only if SSNs.empty().

632

I believe getOpcodeName() will return "atomicrmw" instead of the operation. Also missing space after it.

[gandhi21299]

clang/test/CodeGenCUDA/fp-atomics-optremarks.cu

10

__atomic_fetch_add does not take scope as an argument, how could I add tests with different scopes?

clang/test/CodeGenOpenCL/atomics-remarks-gfx90a.cl

25 ↗

(On Diff #366112)

For some reason, remarks are not emitted here. The command to run looks right above...

[rampitec]

clang/test/CodeGenCUDA/fp-atomics-optremarks.cu

10

At least in the IR test.

[gandhi21299]

clang/test/CodeGenCUDA/fp-atomics-optremarks.cu

10

What do you mean by that?

llvm/lib/CodeGen/AtomicExpandPass.cpp

626

Sorry, what do you mean? SSN will be empty at that point.

632

getOpcodeName() returns atomicrmwoperation, as per the tests the spacing looks correct to me.

[rampitec]

clang/test/CodeGenCUDA/fp-atomics-optremarks.cu

10

You need to test all of that. If you cannot write a proper .cu test, then write an IR test and run llc.

llvm/lib/CodeGen/AtomicExpandPass.cpp

626

I thought want to cache it. But really just declare it here.

632

The operation to report is AI->getOperation(). Spacing is wrong, "operation" is your text.

[gandhi21299]

clang/test/CodeGenCUDA/fp-atomics-optremarks.cu

10

Should I discard this test then since the test fp-atomics-remarks-gfx90a.ll already satisfies?

[gandhi21299]

• corrected remarks by replacing the operation name and updated tests accordingly
• code format

[gandhi21299]

• corrected atomics-remarks-gfx90a.cl test to emit remark as well

[rampitec]

clang/test/CodeGenCUDA/fp-atomics-optremarks.cu

10

CU test is still needed. You also need it in the .cl test below.

llvm/lib/CodeGen/AtomicExpandPass.cpp

595

What should this "Passed" do and why wouldn't just declare it where you use it?

[gandhi21299]

no way to pass memory_scope in __atomic_fetch_add(...), discarded the test.

[gandhi21299]

clang/test/CodeGenCUDA/fp-atomics-optremarks.cu

10

Alright, I am not sure how I can test for the other scopes though.

llvm/lib/CodeGen/AtomicExpandPass.cpp

595

https://llvm.org/docs/Remarks.html

Since this is an informative pass and not that pass failed to optimize, the "Passed" argument is used. I will move it downwards, I thought it might be useful in the future for other operations. Its better below for now anyways.

[gandhi21299]

llvm/lib/CodeGen/AtomicExpandPass.cpp

595

Actually I am getting a runtime error at the line where I declare Remark when I bring it down.

[gandhi21299]

• added clang/test/CodeGenCUDA/fp-atomics-optremarks.cu back
• moved Remark declaration into the else block

[gandhi21299]

• rebased against main branch
• cleaned up code

[rampitec]

Please retitle it without AMDGPU and remove the changes to pass ORE to targets. It is not a part of this change, it is a part of the folloup target specific change.

[arsenm]

llvm/lib/CodeGen/AtomicExpandPass.cpp

632

Missing word atomic?

[gandhi21299]

llvm/lib/CodeGen/AtomicExpandPass.cpp

632

Its already part of the OperationName

[gandhi21299]

• eliminated irrelevant changes to this patch

[rampitec]

Please restore opencl test.

llvm/lib/CodeGen/AtomicExpandPass.cpp

632

Matt is right, missing "atomic" word.

[gandhi21299]

• replaced the OpenCL test
• renamed filenames
• added 'atomic' to the remark and tests

[gandhi21299]

• removed StringExtras.h inclusion

[rampitec]

clang/test/CodeGenOpenCL/atomics-remarks-gfx90a.cl

32 ↗

(On Diff #366343)

It is not system as test name suggests. Just rename to atomic_cas and add calls with all other scopes into the same function.

[gandhi21299]

• adding more tests in clang/test/CodeGenOpenCL/atomics-remarks-gfx90a.cl for various scopes, memory_scope_work_item is called out as invalid by the compiler so excluded that

[rampitec]

clang/test/CodeGenOpenCL/atomics-remarks-gfx90a.cl

33 ↗

(On Diff #366349)

Just combine all the calls into a single function.

[gandhi21299]

• combined all tests into one

[rampitec]

LGTM

[gandhi21299]

Thanks a lot for the review! I will merge this patch in as soon as the CI passes.

[xbolva00]

http://llvm-compile-time-tracker.com/compare.php?from=1f2d40c47f5f8fd01d91d73a1f52044fe1c83225&to=c4e5425aa579d21530ef1766d7144b38a347f247&stat=instructions

Compile time regressions especially for -O0 -g are higher than expected with this patch.

[gandhi21299]

• changed type of ORE from OptimizationRemarkEmitter* to std::shared_ptr<OptimizationRemarkEmitter> and construct it within AtomicExpandPass, this solution is implemented to address for the regressions in many backends due to prerequisite passes

[gandhi21299]

• fixed breaking tests

[arsenm]

I don’t think constructing in the pass is the solution. Why exactly is this introducing such a big slowdown?

llvm/lib/CodeGen/AtomicExpandPass.cpp

184

There’s basically never a reason to use shared_ptr over unique_ptr

[gandhi21299]

@xbolva00 is concerned about Dominator Tree Construction

[gandhi21299]

I will actually revert my changes back with only the tests updated to see if the times are reasonable.

[gandhi21299]

Also, @nikic suggested to construct ORE here if we cannot usefully preserve them. I am not sure if preserving the information is useful though.

[gandhi21299]

reverting changes back to declaring ORE using getAnalysis

[gandhi21299]

@xbolva00 I timed X86/opt-pipeline.ll passes and DTC executed in 0.1% of the total compile time.

[xbolva00]

http://llvm-compile-time-tracker.com/compare.php?from=530aa7e4da14fb22493ab7e175f8c34dd10333d3&to=435785214f73ff0c92e97f2ade6356e3ba3bf661&stat=instructions

Still same problem with -O0 -g.
Please wait a bit for more people, dont rush.

[xbolva00]

sqlite3 +1.3% increase of compile time for -O0 -g is simply not acceptable.

[nikic]

In D106891#2945342, @arsenm wrote:

I don’t think constructing in the pass is the solution. Why exactly is this introducing such a big slowdown?

The reason are the additional DominatorTree and LoopInfo calculations. These have a big impact at O0. These analysis calculations are caused by a deficiency in the legacy pass manager, which is still used for the codegen pipeline: Even though these analyses are only needed if diagnostic hotness is used, the pass manager requires them to be scheduled unconditionally.

The way to avoid this is to construct ORE without going through the pass manager. Grep for OptimizationRemarkEmitter ORE for various passes that already do this (though the reason is different in some cases -- there are some analysis preservation concerns with loop passes).

[gandhi21299]

Okay, sorry about that. Thanks for reverting my commit. I will use a unique_ptr and wait for another approval.

[gandhi21299]

reverting type of ORE from OptimizationRemarkEmitter * back to std::unique_ptr<OptimizationRemarkEmitter> and constructing it in AtomicExpand to avoid DTC and LI overhead.

[nikic]

You also need to drop the getAnalysisUsage() change, otherwise the pass manager will still create the ORE itself.

PS: This revision still has incorrectly configured permissions, which prevents me from leaving line comments. Revisions on this Phabricator instance should always be publicly writable.

[gandhi21299]

• removing analysis requirement as requested

+ nikic

[gandhi21299]

• fixing breaking tests by eliminating passes that are no longer in the pass pipelines

[xbolva00]

Please drop a change in PowerPC/O3-pipeline.ll

[gandhi21299]

• eliminated changes in PowerPC/O3-pipeline.ll, as requested

[gandhi21299]

Any ideas on what could be causing the failure in windows pre-merge checks?

[xbolva00]

Not related to your patch, feel free to ignore

[gandhi21299]

Alright, please let me know if this patch is good for merge at your convenience.

[rampitec]

llvm/lib/CodeGen/AtomicExpandPass.cpp

184

Is there a reason to construct it upfront and not just use a local variable only when needed? Like in StackProtector.cpp for example.

[gandhi21299]

llvm/lib/CodeGen/AtomicExpandPass.cpp

184

We can certainly implement it as a local variable as long as we have access to the function this pass is operating on. I was thinking of its potential use throughout this pass in the future.

[rampitec]

llvm/lib/CodeGen/AtomicExpandPass.cpp

184

You have access to the function, AI->getParent()->getParent().
You also will not need to pass ORE everywhere in the subsequent patch, just construct it in target in place.

[gandhi21299]

llvm/lib/CodeGen/AtomicExpandPass.cpp

184

Sounds like a plan!

[gandhi21299]

• ORE does not need to be a pointer anymore, it is constructed as local variable with this patch as requested by reviewer.

[rampitec]

LGTM, but please wait for others too.

[gandhi21299]

Will do, thanks!

[arsenm]

LGTM

clang/test/CodeGenCUDA/fp-atomics-optremarks.cu

+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN:   -target-cpu gfx90a -Rpass=atomic-expand -S -o - 2>&1 | \
+// RUN:   FileCheck %s --check-prefix=GFX90A-CAS
+
+// REQUIRES: amdgpu-registered-target
+
+#include "Inputs/cuda.h"
+#include <stdatomic.h>
+
+// GFX90A-CAS: A compare and swap loop was generated for an atomic operation at system memory scope
+// GFX90A-CAS-LABEL: _Z14atomic_add_casPf
+// GFX90A-CAS:  flat_atomic_cmpswap v0, v[2:3], v[4:5] glc
+// GFX90A-CAS:  s_cbranch_execnz
+__device__ float atomic_add_cas(float *p) {
+  return __atomic_fetch_add(p, 1.0f, memory_order_relaxed);
+}

llvm/include/llvm/CodeGen/TargetLowering.h

 #define LLVM_CODEGEN_TARGETLOWERING_H
 
 #include "llvm/ADT/APInt.h"
 #include "llvm/ADT/ArrayRef.h"
 #include "llvm/ADT/DenseMap.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/StringRef.h"
+#include "llvm/Analysis/OptimizationRemarkEmitter.h"
 #include "llvm/CodeGen/DAGCombine.h"
 #include "llvm/CodeGen/ISDOpcodes.h"
 #include "llvm/CodeGen/RuntimeLibcalls.h"
 #include "llvm/CodeGen/SelectionDAG.h"
 #include "llvm/CodeGen/SelectionDAGNodes.h"
 #include "llvm/CodeGen/TargetCallingConv.h"
 #include "llvm/CodeGen/ValueTypes.h"
 #include "llvm/IR/Attributes.h"
public:
   /// AtomicExpand pass.
   virtual AtomicExpansionKind
   shouldExpandAtomicCmpXchgInIR(AtomicCmpXchgInst *AI) const {
     return AtomicExpansionKind::None;
   }
 
   /// Returns how the IR-level AtomicExpand pass should expand the given
   /// AtomicRMW, if at all. Default is to never expand.
-  virtual AtomicExpansionKind shouldExpandAtomicRMWInIR(AtomicRMWInst *RMW) const {
+  virtual AtomicExpansionKind
+  shouldExpandAtomicRMWInIR(AtomicRMWInst *RMW,
+                            OptimizationRemarkEmitter *ORE) const {
     return RMW->isFloatingPointOperation() ?
       AtomicExpansionKind::CmpXChg : AtomicExpansionKind::None;
   }
 
   /// On some platforms, an AtomicRMW that never actually modifies the value
   /// (such as fetch_add of 0) can be turned into a fence followed by an
   /// atomic load. This may sound useless, but it makes it possible for the
   /// processor to keep the cacheline shared, dramatically improving

llvm/lib/CodeGen/AtomicExpandPass.cpp

 // include the use of (intrinsic-based) load-linked/store-conditional loops,
 // AtomicCmpXchg, or type coercions.
 //
 //===----------------------------------------------------------------------===//
 
 #include "llvm/ADT/ArrayRef.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/SmallVector.h"
+#include "llvm/Analysis/OptimizationRemarkEmitter.h"
 #include "llvm/CodeGen/AtomicExpandUtils.h"
 #include "llvm/CodeGen/RuntimeLibcalls.h"
 #include "llvm/CodeGen/TargetLowering.h"
 #include "llvm/CodeGen/TargetPassConfig.h"
 #include "llvm/CodeGen/TargetSubtargetInfo.h"
 #include "llvm/CodeGen/ValueTypes.h"
 #include "llvm/IR/Attributes.h"
 #include "llvm/IR/BasicBlock.h"

 using namespace llvm;
 
 #define DEBUG_TYPE "atomic-expand"
 
 namespace {
 
   class AtomicExpand: public FunctionPass {
     const TargetLowering *TLI = nullptr;
+    OptimizationRemarkEmitter *ORE;
 
   public:
     static char ID; // Pass identification, replacement for typeid
 
     AtomicExpand() : FunctionPass(ID) {
       initializeAtomicExpandPass(*PassRegistry::getPassRegistry());
     }
 
     bool runOnFunction(Function &F) override;
 
   private:
+    void getAnalysisUsage(AnalysisUsage &AU) const override;
+    TargetLowering::AtomicExpansionKind
+    emitAtomicExpansionRemarks(AtomicRMWInst *RMW,
+                               TargetLowering::AtomicExpansionKind Kind,
+                               OptimizationRemark Remark);
     bool bracketInstWithFences(Instruction *I, AtomicOrdering Order);
     IntegerType *getCorrespondingIntegerType(Type *T, const DataLayout &DL);
     LoadInst *convertAtomicLoadToIntegerType(LoadInst *LI);
     bool tryExpandAtomicLoad(LoadInst *LI);
     bool expandAtomicLoadToLL(LoadInst *LI);
     bool expandAtomicLoadToCmpXchg(LoadInst *LI);
     StoreInst *convertAtomicStoreToIntegerType(StoreInst *SI);
     bool expandAtomicStore(StoreInst *SI);

 template <typename Inst>
 static bool atomicSizeSupported(const TargetLowering *TLI, Inst *I) {
   unsigned Size = getAtomicOpSize(I);
   Align Alignment = I->getAlign();
   return Alignment >= Size &&
          Size <= TLI->getMaxAtomicSizeInBitsSupported() / 8;
 }
 
+void AtomicExpand::getAnalysisUsage(AnalysisUsage &AU) const {
+  AU.addRequired<OptimizationRemarkEmitterWrapperPass>();
+}
+
 bool AtomicExpand::runOnFunction(Function &F) {
   auto *TPC = getAnalysisIfAvailable<TargetPassConfig>();
   if (!TPC)
     return false;
 
+  ORE = &getAnalysis<OptimizationRemarkEmitterWrapperPass>().getORE();
   auto &TM = TPC->getTM<TargetMachine>();
   if (!TM.getSubtargetImpl(F)->enableAtomicExpand())
     return false;
   TLI = TM.getSubtargetImpl(F)->getTargetLowering();
 
   SmallVector<Instruction *, 1> AtomicInsts;
 
   // Changing control-flow while iterating through it is a bad idea, so gather a
case AtomicRMWInst::FAdd:
     return Builder.CreateFAdd(Loaded, Inc, "new");
   case AtomicRMWInst::FSub:
     return Builder.CreateFSub(Loaded, Inc, "new");
   default:
     llvm_unreachable("Unknown atomic op");
   }
 }
 
+TargetLowering::AtomicExpansionKind AtomicExpand::emitAtomicExpansionRemarks(
+    AtomicRMWInst *RMW, TargetLowering::AtomicExpansionKind Kind,
+    OptimizationRemark Remark) {
+  ORE->emit([&]() { return Remark; });
+  return Kind;
+}
+
 bool AtomicExpand::tryExpandAtomicRMW(AtomicRMWInst *AI) {
-  switch (TLI->shouldExpandAtomicRMWInIR(AI)) {
+  TargetLowering::AtomicExpansionKind Kind =
+      TLI->shouldExpandAtomicRMWInIR(AI, ORE);
+  OptimizationRemark Remark(DEBUG_TYPE, "Passed", AI->getFunction());
+  switch (Kind) {
   case TargetLoweringBase::AtomicExpansionKind::None:
     return false;
   case TargetLoweringBase::AtomicExpansionKind::LLSC: {
     unsigned MinCASSize = TLI->getMinCmpXchgSizeInBits() / 8;
     unsigned ValueSize = getAtomicOpSize(AI);
     if (ValueSize < MinCASSize) {
       expandPartwordAtomicRMW(AI,
                               TargetLoweringBase::AtomicExpansionKind::LLSC);
     } else {
       auto PerformOp = [&](IRBuilder<> &Builder, Value *Loaded) {
         return performAtomicOp(AI->getOperation(), Builder, Loaded,
                                AI->getValOperand());
       };
if (ValueSize < MinCASSize) {
       // TODO: Handle atomicrmw fadd/fsub
       if (AI->getType()->isFloatingPointTy())
         return false;
 
       expandPartwordAtomicRMW(AI,
                               TargetLoweringBase::AtomicExpansionKind::CmpXChg);
     } else {
       expandAtomicRMWToCmpXchg(AI, createCmpXchgInstFun);
+      emitAtomicExpansionRemarks(
+          AI, Kind,
+          Remark << "A compare and swap loop was generated for an atomic "
+                    "operation "
+                    "at "
+                 << (AI->getSyncScopeID() ? "system" : "single thread")
+                 << " memory scope");
     }
     return true;
   }
   case TargetLoweringBase::AtomicExpansionKind::MaskedIntrinsic: {
     expandAtomicRMWToMaskedIntrinsic(AI);
     return true;
   }
   default:

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.h

 /// Interface definition of the TargetLowering class that is common
 /// to all AMD GPUs.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_LIB_TARGET_AMDGPU_AMDGPUISELLOWERING_H
 #define LLVM_LIB_TARGET_AMDGPU_AMDGPUISELLOWERING_H
 
+#include "llvm/Analysis/OptimizationRemarkEmitter.h"
 #include "llvm/CodeGen/CallingConvLower.h"
 #include "llvm/CodeGen/TargetLowering.h"
 
 namespace llvm {
 
 class AMDGPUMachineFunction;
 class AMDGPUSubtarget;
 struct ArgDescriptor;
public:
   /// type of implicit parameter.
   uint32_t getImplicitParameterOffset(const MachineFunction &MF,
                                       const ImplicitParameter Param) const;
 
   MVT getFenceOperandTy(const DataLayout &DL) const override {
     return MVT::i32;
   }
 
-  AtomicExpansionKind shouldExpandAtomicRMWInIR(AtomicRMWInst *) const override;
+  AtomicExpansionKind
+  shouldExpandAtomicRMWInIR(AtomicRMWInst *,
+                            OptimizationRemarkEmitter *ORE) const override;
 
   bool isConstantUnsignedBitfieldExtactLegal(unsigned Opc, LLT Ty1,
                                              LLT Ty2) const override;
 };
 
 namespace AMDGPUISD {
 
 enum NodeType : unsigned {

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp

 //===----------------------------------------------------------------------===//
 
 #include "AMDGPUISelLowering.h"
 #include "AMDGPU.h"
 #include "AMDGPUInstrInfo.h"
 #include "AMDGPUMachineFunction.h"
 #include "GCNSubtarget.h"
 #include "SIMachineFunctionInfo.h"
+#include "llvm/Analysis/OptimizationRemarkEmitter.h"
 #include "llvm/CodeGen/Analysis.h"
 #include "llvm/IR/DiagnosticInfo.h"
 #include "llvm/IR/IntrinsicsAMDGPU.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/KnownBits.h"
 #include "llvm/Target/TargetMachine.h"
 
 using namespace llvm;
case ISD::INTRINSIC_WO_CHAIN: {
     }
   }
   default:
     return false;
   }
 }
 
 TargetLowering::AtomicExpansionKind
-AMDGPUTargetLowering::shouldExpandAtomicRMWInIR(AtomicRMWInst *RMW) const {
+AMDGPUTargetLowering::shouldExpandAtomicRMWInIR(
+    AtomicRMWInst *RMW, OptimizationRemarkEmitter *ORE) const {
   switch (RMW->getOperation()) {
   case AtomicRMWInst::Nand:
   case AtomicRMWInst::FAdd:
   case AtomicRMWInst::FSub:
     return AtomicExpansionKind::CmpXChg;
   default:
     return AtomicExpansionKind::None;
   }
 }
 
 bool AMDGPUTargetLowering::isConstantUnsignedBitfieldExtactLegal(
     unsigned Opc, LLT Ty1, LLT Ty2) const {
   return Ty1 == Ty2 && (Ty1 == LLT::scalar(32) || Ty1 == LLT::scalar(64));
 }

llvm/lib/Target/AMDGPU/SIISelLowering.h

bool isCanonicalized(Register Reg, MachineFunction &MF,
                        unsigned MaxDepth = 5) const;
   bool denormalsEnabledForType(const SelectionDAG &DAG, EVT VT) const;
   bool denormalsEnabledForType(LLT Ty, MachineFunction &MF) const;
 
   bool isKnownNeverNaNForTargetNode(SDValue Op,
                                     const SelectionDAG &DAG,
                                     bool SNaN = false,
                                     unsigned Depth = 0) const override;
-  AtomicExpansionKind shouldExpandAtomicRMWInIR(AtomicRMWInst *) const override;
+  AtomicExpansionKind
+  shouldExpandAtomicRMWInIR(AtomicRMWInst *,
+                            OptimizationRemarkEmitter *ORE) const override;
 
   virtual const TargetRegisterClass *
   getRegClassFor(MVT VT, bool isDivergent) const override;
   virtual bool requiresUniformRegister(MachineFunction &MF,
                                        const Value *V) const override;
   Align getPrefLoopAlignment(MachineLoop *ML) const override;
 
   void allocateHSAUserSGPRs(CCState &CCInfo,

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

 /// Custom DAG lowering for SI
 //
 //===----------------------------------------------------------------------===//
 
 #include "SIISelLowering.h"
 #include "AMDGPU.h"
 #include "AMDGPUInstrInfo.h"
 #include "AMDGPUTargetMachine.h"
+#include "GCNSubtarget.h"
 #include "SIMachineFunctionInfo.h"
 #include "SIRegisterInfo.h"
 #include "llvm/ADT/Statistic.h"
 #include "llvm/Analysis/LegacyDivergenceAnalysis.h"
+#include "llvm/Analysis/OptimizationRemarkEmitter.h"
 #include "llvm/BinaryFormat/ELF.h"
 #include "llvm/CodeGen/Analysis.h"
 #include "llvm/CodeGen/FunctionLoweringInfo.h"
 #include "llvm/CodeGen/GlobalISel/GISelKnownBits.h"
 #include "llvm/CodeGen/MachineFunction.h"
 #include "llvm/CodeGen/MachineLoopInfo.h"
 #include "llvm/IR/DiagnosticInfo.h"
 #include "llvm/IR/IntrinsicInst.h"

 static bool fpModeMatchesGlobalFPAtomicMode(const AtomicRMWInst *RMW) {
   const fltSemantics &Flt = RMW->getType()->getScalarType()->getFltSemantics();
   auto DenormMode = RMW->getParent()->getParent()->getDenormalMode(Flt);
   if (&Flt == &APFloat::IEEEsingle())
     return DenormMode == DenormalMode::getPreserveSign();
   return DenormMode == DenormalMode::getIEEE();
 }
 
-TargetLowering::AtomicExpansionKind
-SITargetLowering::shouldExpandAtomicRMWInIR(AtomicRMWInst *RMW) const {
+TargetLowering::AtomicExpansionKind SITargetLowering::shouldExpandAtomicRMWInIR(
+    AtomicRMWInst *RMW, OptimizationRemarkEmitter *ORE) const {
   switch (RMW->getOperation()) {
   case AtomicRMWInst::FAdd: {
     Type *Ty = RMW->getType();
 
     // We don't have a way to support 16-bit atomics now, so just leave them
     // as-is.
     if (Ty->isHalfTy())
       return AtomicExpansionKind::None;
 
     if (!Ty->isFloatTy() && (!Subtarget->hasGFX90AInsts() || !Ty->isDoubleTy()))
       return AtomicExpansionKind::CmpXChg;
 
     unsigned AS = RMW->getPointerAddressSpace();
 
     if ((AS == AMDGPUAS::GLOBAL_ADDRESS || AS == AMDGPUAS::FLAT_ADDRESS) &&
          Subtarget->hasAtomicFaddInsts()) {
       // The amdgpu-unsafe-fp-atomics attribute enables generation of unsafe
       // floating point atomic instructions. May generate more efficient code,
       // but may not respect rounding and denormal modes, and may give incorrect
       // results for certain memory destinations.
       if (RMW->getFunction()
               ->getFnAttribute("amdgpu-unsafe-fp-atomics")
               .getValueAsString() != "true")
         return AtomicExpansionKind::CmpXChg;
 
       if (Subtarget->hasGFX90AInsts()) {
         if (Ty->isFloatTy() && AS == AMDGPUAS::FLAT_ADDRESS)
           return AtomicExpansionKind::CmpXChg;
 
         auto SSID = RMW->getSyncScopeID();
         if (SSID == SyncScope::System ||
             SSID == RMW->getContext().getOrInsertSyncScopeID("one-as"))
           return AtomicExpansionKind::CmpXChg;
 
         return AtomicExpansionKind::None;
       }
 
       if (AS == AMDGPUAS::FLAT_ADDRESS)
         return AtomicExpansionKind::CmpXChg;
 
       return RMW->use_empty() ? AtomicExpansionKind::None
                               : AtomicExpansionKind::CmpXChg;
     }
 
     // DS FP atomics do repect the denormal mode, but the rounding mode is fixed
     // to round-to-nearest-even.
     // The only exception is DS_ADD_F64 which never flushes regardless of mode.
     if (AS == AMDGPUAS::LOCAL_ADDRESS && Subtarget->hasLDSFPAtomics()) {
       if (!Ty->isDoubleTy())
         return AtomicExpansionKind::None;
 
       return (fpModeMatchesGlobalFPAtomicMode(RMW) ||
               RMW->getFunction()
                       ->getFnAttribute("amdgpu-unsafe-fp-atomics")
                       .getValueAsString() == "true")
                  ? AtomicExpansionKind::None
                  : AtomicExpansionKind::CmpXChg;
     }
 
     return AtomicExpansionKind::CmpXChg;
   }
   default:
     break;
   }
 
-  return AMDGPUTargetLowering::shouldExpandAtomicRMWInIR(RMW);
+  return AMDGPUTargetLowering::shouldExpandAtomicRMWInIR(RMW, ORE);
 }
 
 const TargetRegisterClass *
 SITargetLowering::getRegClassFor(MVT VT, bool isDivergent) const {
   const TargetRegisterClass *RC = TargetLoweringBase::getRegClassFor(VT, false);
   const SIRegisterInfo *TRI = Subtarget->getRegisterInfo();
   if (RC == &AMDGPU::VReg_1RegClass && !isDivergent)
     return Subtarget->getWavefrontSize() == 64 ? &AMDGPU::SReg_64RegClass

llvm/lib/Target/X86/X86ISelLowering.h

 //===-- X86ISelLowering.h - X86 DAG Lowering Interface ----------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines the interfaces that X86 uses to lower LLVM code into a
 // selection DAG.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_LIB_TARGET_X86_X86ISELLOWERING_H
 #define LLVM_LIB_TARGET_X86_X86ISELLOWERING_H
 
+#include "llvm/Analysis/OptimizationRemarkEmitter.h"
 #include "llvm/CodeGen/MachineFunction.h"
 #include "llvm/CodeGen/TargetLowering.h"
 
 namespace llvm {
   class X86Subtarget;
   class X86TargetMachine;
 
   namespace X86ISD {
bool CanLowerReturn(CallingConv::ID CallConv, MachineFunction &MF,
                         LLVMContext &Context) const override;
 
     const MCPhysReg *getScratchRegisters(CallingConv::ID CC) const override;
 
     TargetLoweringBase::AtomicExpansionKind
     shouldExpandAtomicLoadInIR(LoadInst *LI) const override;
     bool shouldExpandAtomicStoreInIR(StoreInst *SI) const override;
     TargetLoweringBase::AtomicExpansionKind
-    shouldExpandAtomicRMWInIR(AtomicRMWInst *AI) const override;
+    shouldExpandAtomicRMWInIR(AtomicRMWInst *AI,
+                              OptimizationRemarkEmitter *ORE) const override;
 
     LoadInst *
     lowerIdempotentRMWIntoFencedLoad(AtomicRMWInst *AI) const override;
 
     bool lowerAtomicStoreAsStoreSDNode(const StoreInst &SI) const override;
     bool lowerAtomicLoadAsLoadSDNode(const LoadInst &LI) const override;
 
     bool needsCmpXchgNb(Type *MemType) const;

llvm/lib/Target/X86/X86ISelLowering.cpp

 #include "llvm/ADT/SmallBitVector.h"
 #include "llvm/ADT/SmallSet.h"
 #include "llvm/ADT/Statistic.h"
 #include "llvm/ADT/StringExtras.h"
 #include "llvm/ADT/StringSwitch.h"
 #include "llvm/Analysis/BlockFrequencyInfo.h"
 #include "llvm/Analysis/EHPersonalities.h"
 #include "llvm/Analysis/ObjCARCUtil.h"
+#include "llvm/Analysis/OptimizationRemarkEmitter.h"
 #include "llvm/Analysis/ProfileSummaryInfo.h"
 #include "llvm/Analysis/VectorUtils.h"
 #include "llvm/CodeGen/IntrinsicLowering.h"
 #include "llvm/CodeGen/MachineFrameInfo.h"
 #include "llvm/CodeGen/MachineFunction.h"
 #include "llvm/CodeGen/MachineInstrBuilder.h"
 #include "llvm/CodeGen/MachineJumpTableInfo.h"
 #include "llvm/CodeGen/MachineLoopInfo.h"
if (MemType->getPrimitiveSizeInBits() == 64 && !Subtarget.is64Bit() &&
       (Subtarget.hasSSE1() || Subtarget.hasX87()))
     return AtomicExpansionKind::None;
 
   return needsCmpXchgNb(MemType) ? AtomicExpansionKind::CmpXChg
                                  : AtomicExpansionKind::None;
 }
 
 TargetLowering::AtomicExpansionKind
-X86TargetLowering::shouldExpandAtomicRMWInIR(AtomicRMWInst *AI) const {
+X86TargetLowering::shouldExpandAtomicRMWInIR(
+    AtomicRMWInst *AI, OptimizationRemarkEmitter *ORE) const {
   unsigned NativeWidth = Subtarget.is64Bit() ? 64 : 32;
   Type *MemType = AI->getType();
 
   // If the operand is too big, we must see if cmpxchg8/16b is available
   // and default to library calls otherwise.
   if (MemType->getPrimitiveSizeInBits() > NativeWidth) {
     return needsCmpXchgNb(MemType) ? AtomicExpansionKind::CmpXChg
                                    : AtomicExpansionKind::None;

llvm/test/CodeGen/AMDGPU/fp-atomics-remarks-gfx90a.ll

+; RUN: llc -march=amdgcn -mcpu=gfx90a -verify-machineinstrs --pass-remarks=atomic-expand \
+; RUN:      %s -o - 2>&1 | FileCheck %s --check-prefix=GFX90A-CAS
+
+; GFX90A-CAS: A compare and swap loop was generated for an atomic operation at system memory scope
+; GFX90A-CAS-LABEL: _Z14atomic_add_casPf:
+; GFX90A-CAS:    flat_atomic_cmpswap v2, v[0:1], v[2:3] glc
+; GFX90A-CAS:    s_cbranch_execnz
+define dso_local void @_Z14atomic_add_casPf(float* %p) {
+entry:
+  %ret = atomicrmw fadd float* %p, float 7.0 monotonic, align 4
+  ret void
+}

llvm/test/CodeGen/AMDGPU/llc-pipeline.ll

 ; GCN-O0-NEXT:    AMDGPU Inline All Functions
 ; GCN-O0-NEXT:    CallGraph Construction
 ; GCN-O0-NEXT:    Call Graph SCC Pass Manager
 ; GCN-O0-NEXT:      Inliner for always_inline functions
 ; GCN-O0-NEXT:    A No-Op Barrier Pass
 ; GCN-O0-NEXT:    Lower OpenCL enqueued blocks
 ; GCN-O0-NEXT:    Lower uses of LDS variables from non-kernel functions
 ; GCN-O0-NEXT:    FunctionPass Manager
+; GCN-O0-NEXT:      Dominator Tree Construction
+; GCN-O0-NEXT:      Natural Loop Information
+; GCN-O0-NEXT:      Lazy Branch Probability Analysis
+; GCN-O0-NEXT:      Lazy Block Frequency Analysis
+; GCN-O0-NEXT:      Optimization Remark Emitter
 ; GCN-O0-NEXT:      Expand Atomic instructions
 ; GCN-O0-NEXT:      Lower constant intrinsics
 ; GCN-O0-NEXT:      Remove unreachable blocks from the CFG
 ; GCN-O0-NEXT:      Expand vector predication intrinsics
 ; GCN-O0-NEXT:      Scalarize Masked Memory Intrinsics
 ; GCN-O0-NEXT:      Expand reduction intrinsics
 ; GCN-O0-NEXT:    CallGraph Construction
 ; GCN-O0-NEXT:    Call Graph SCC Pass Manager

 ; GCN-O1-NEXT:    CallGraph Construction
 ; GCN-O1-NEXT:    Call Graph SCC Pass Manager
 ; GCN-O1-NEXT:      Inliner for always_inline functions
 ; GCN-O1-NEXT:    A No-Op Barrier Pass
 ; GCN-O1-NEXT:    Lower OpenCL enqueued blocks
 ; GCN-O1-NEXT:    Lower uses of LDS variables from non-kernel functions
 ; GCN-O1-NEXT:    FunctionPass Manager
 ; GCN-O1-NEXT:      Infer address spaces
+; GCN-O1-NEXT:      Dominator Tree Construction
+; GCN-O1-NEXT:      Natural Loop Information
+; GCN-O1-NEXT:      Lazy Branch Probability Analysis
+; GCN-O1-NEXT:      Lazy Block Frequency Analysis
+; GCN-O1-NEXT:      Optimization Remark Emitter
 ; GCN-O1-NEXT:      Expand Atomic instructions
 ; GCN-O1-NEXT:      AMDGPU Promote Alloca
 ; GCN-O1-NEXT:      Dominator Tree Construction
 ; GCN-O1-NEXT:      SROA
 ; GCN-O1-NEXT:      Post-Dominator Tree Construction
 ; GCN-O1-NEXT:      Natural Loop Information
 ; GCN-O1-NEXT:      Legacy Divergence Analysis
 ; GCN-O1-NEXT:      AMDGPU IR optimizations

 ; GCN-O1-OPTS-NEXT:    CallGraph Construction
 ; GCN-O1-OPTS-NEXT:    Call Graph SCC Pass Manager
 ; GCN-O1-OPTS-NEXT:      Inliner for always_inline functions
 ; GCN-O1-OPTS-NEXT:    A No-Op Barrier Pass
 ; GCN-O1-OPTS-NEXT:    Lower OpenCL enqueued blocks
 ; GCN-O1-OPTS-NEXT:    Lower uses of LDS variables from non-kernel functions
 ; GCN-O1-OPTS-NEXT:    FunctionPass Manager
 ; GCN-O1-OPTS-NEXT:      Infer address spaces
+; GCN-O1-OPTS-NEXT:      Dominator Tree Construction
+; GCN-O1-OPTS-NEXT:      Natural Loop Information
+; GCN-O1-OPTS-NEXT:      Lazy Branch Probability Analysis
+; GCN-O1-OPTS-NEXT:      Lazy Block Frequency Analysis
+; GCN-O1-OPTS-NEXT:      Optimization Remark Emitter
 ; GCN-O1-OPTS-NEXT:      Expand Atomic instructions
 ; GCN-O1-OPTS-NEXT:      AMDGPU Promote Alloca
 ; GCN-O1-OPTS-NEXT:      Dominator Tree Construction
 ; GCN-O1-OPTS-NEXT:      SROA
 ; GCN-O1-OPTS-NEXT:      Basic Alias Analysis (stateless AA impl)
 ; GCN-O1-OPTS-NEXT:      Function Alias Analysis Results
 ; GCN-O1-OPTS-NEXT:      Memory SSA
 ; GCN-O1-OPTS-NEXT:      Natural Loop Information

 ; GCN-O2-NEXT:    CallGraph Construction
 ; GCN-O2-NEXT:    Call Graph SCC Pass Manager
 ; GCN-O2-NEXT:      Inliner for always_inline functions
 ; GCN-O2-NEXT:    A No-Op Barrier Pass
 ; GCN-O2-NEXT:    Lower OpenCL enqueued blocks
 ; GCN-O2-NEXT:    Lower uses of LDS variables from non-kernel functions
 ; GCN-O2-NEXT:    FunctionPass Manager
 ; GCN-O2-NEXT:      Infer address spaces
+; GCN-O2-NEXT:      Dominator Tree Construction
+; GCN-O2-NEXT:      Natural Loop Information
+; GCN-O2-NEXT:      Lazy Branch Probability Analysis
+; GCN-O2-NEXT:      Lazy Block Frequency Analysis
+; GCN-O2-NEXT:      Optimization Remark Emitter
 ; GCN-O2-NEXT:      Expand Atomic instructions
 ; GCN-O2-NEXT:      AMDGPU Promote Alloca
 ; GCN-O2-NEXT:      Dominator Tree Construction
 ; GCN-O2-NEXT:      SROA
 ; GCN-O2-NEXT:      Basic Alias Analysis (stateless AA impl)
 ; GCN-O2-NEXT:      Function Alias Analysis Results
 ; GCN-O2-NEXT:      Memory SSA
 ; GCN-O2-NEXT:      Natural Loop Information

 ; GCN-O3-NEXT:    CallGraph Construction
 ; GCN-O3-NEXT:    Call Graph SCC Pass Manager
 ; GCN-O3-NEXT:      Inliner for always_inline functions
 ; GCN-O3-NEXT:    A No-Op Barrier Pass
 ; GCN-O3-NEXT:    Lower OpenCL enqueued blocks
 ; GCN-O3-NEXT:    Lower uses of LDS variables from non-kernel functions
 ; GCN-O3-NEXT:    FunctionPass Manager
 ; GCN-O3-NEXT:      Infer address spaces
+; GCN-O3-NEXT:      Dominator Tree Construction
+; GCN-O3-NEXT:      Natural Loop Information
+; GCN-O3-NEXT:      Lazy Branch Probability Analysis
+; GCN-O3-NEXT:      Lazy Block Frequency Analysis
+; GCN-O3-NEXT:      Optimization Remark Emitter
 ; GCN-O3-NEXT:      Expand Atomic instructions
 ; GCN-O3-NEXT:      AMDGPU Promote Alloca
 ; GCN-O3-NEXT:      Dominator Tree Construction
 ; GCN-O3-NEXT:      SROA
 ; GCN-O3-NEXT:      Basic Alias Analysis (stateless AA impl)
 ; GCN-O3-NEXT:      Function Alias Analysis Results
 ; GCN-O3-NEXT:      Memory SSA
 ; GCN-O3-NEXT:      Natural Loop Information

llvm/test/CodeGen/X86/O0-pipeline.ll

 ; When EXPENSIVE_CHECKS are enabled, the machine verifier appears between each
 ; pass. Ignore it with 'grep -v'.
 ; RUN: llc -mtriple=x86_64-- -O0 -debug-pass=Structure < %s -o /dev/null 2>&1 \
 ; RUN:   | grep -v 'Verify generated machine code' | FileCheck %s
 
 ; REQUIRES: asserts
 
 ; CHECK-LABEL: Pass Arguments:
 ; CHECK-NEXT: Target Library Information
 ; CHECK-NEXT: Target Pass Configuration
 ; CHECK-NEXT: Machine Module Information
 ; CHECK-NEXT: Target Transform Information
+; CHECK-NEXT: Profile summary info
 ; CHECK-NEXT: Create Garbage Collector Module Metadata
 ; CHECK-NEXT: Assumption Cache Tracker
-; CHECK-NEXT: Profile summary info
 ; CHECK-NEXT: Machine Branch Probability Analysis
 ; CHECK-NEXT:   ModulePass Manager
 ; CHECK-NEXT:     Pre-ISel Intrinsic Lowering
 ; CHECK-NEXT:     FunctionPass Manager
+; CHECK-NEXT:       Dominator Tree Construction
+; CHECK-NEXT:       Natural Loop Information
+; CHECK-NEXT:       Lazy Branch Probability Analysis
+; CHECK-NEXT:       Lazy Block Frequency Analysis
+; CHECK-NEXT:       Optimization Remark Emitter
 ; CHECK-NEXT:       Expand Atomic instructions
 ; CHECK-NEXT:       Lower AMX intrinsics
 ; CHECK-NEXT:       Lower AMX type for load/store
 ; CHECK-NEXT:       Pre AMX Tile Config
 ; CHECK-NEXT:       Module Verifier
 ; CHECK-NEXT:       Lower Garbage Collection Instructions
 ; CHECK-NEXT:       Shadow Stack GC Lowering
 ; CHECK-NEXT:       Lower constant intrinsics

llvm/test/CodeGen/X86/opt-pipeline.ll

 
 ; REQUIRES: asserts
 
 ; CHECK-LABEL: Pass Arguments:
 ; CHECK-NEXT: Target Library Information
 ; CHECK-NEXT: Target Pass Configuration
 ; CHECK-NEXT: Machine Module Information
 ; CHECK-NEXT: Target Transform Information
+; CHECK-NEXT: Profile summary info
 ; CHECK-NEXT: Type-Based Alias Analysis
 ; CHECK-NEXT: Scoped NoAlias Alias Analysis
 ; CHECK-NEXT: Assumption Cache Tracker
-; CHECK-NEXT: Profile summary info
 ; CHECK-NEXT: Create Garbage Collector Module Metadata
 ; CHECK-NEXT: Machine Branch Probability Analysis
 ; CHECK-NEXT:   ModulePass Manager
 ; CHECK-NEXT:     Pre-ISel Intrinsic Lowering
 ; CHECK-NEXT:     FunctionPass Manager
+; CHECK-NEXT:       Dominator Tree Construction
+; CHECK-NEXT:       Natural Loop Information
+; CHECK-NEXT:       Lazy Branch Probability Analysis
+; CHECK-NEXT:       Lazy Block Frequency Analysis
+; CHECK-NEXT:       Optimization Remark Emitter
 ; CHECK-NEXT:       Expand Atomic instructions
 ; CHECK-NEXT:       Lower AMX intrinsics
 ; CHECK-NEXT:       Lower AMX type for load/store
 ; CHECK-NEXT:       Module Verifier
 ; CHECK-NEXT:       Dominator Tree Construction
 ; CHECK-NEXT:       Basic Alias Analysis (stateless AA impl)
 ; CHECK-NEXT:       Natural Loop Information
 ; CHECK-NEXT:       Canonicalize natural loops