[clang] [Clang][AMDGPU] Add ``amdgcn_av("none")`` attribute for atomic expressions (PR #199622)

[Pierre van Houtryve via cfe-commits]

================
@@ -351,6 +351,61 @@ static Attr *handleMustTailAttr(Sema &S, Stmt *St, const ParsedAttr &A,
   return ::new (S.Context) MustTailAttr(S.Context, A);
 }
 
+/// Return true if St is an atomic operation or fence builtin call.
+static bool isAtomicOp(const Stmt *St) {
+  const Expr *E = dyn_cast<Expr>(St);
+  if (!E)
+    return false;
+
+  E = E->IgnoreParenCasts();
+
+  if (isa<AtomicExpr>(E))
+    return true;
+
+  const CallExpr *CE = dyn_cast<CallExpr>(E);
+  if (!CE)
+    return false;
+
+  unsigned BuiltinID = CE->getBuiltinCallee();
+  switch (BuiltinID) {
+  case Builtin::BI__atomic_thread_fence:
+  case Builtin::BI__atomic_signal_fence:
+  case Builtin::BI__c11_atomic_thread_fence:
+  case Builtin::BI__c11_atomic_signal_fence:
+  case Builtin::BI__scoped_atomic_thread_fence:
+    return true;
+  default:
+    break;
+  }
+
+  // Check for target-specific fence builtins.
+  if (const FunctionDecl *FD = CE->getDirectCallee()) {
----------------
Pierre-vh wrote:

Why can't this just check `BI__builtin_amdgcn_fence` like other built-ins check above ?

https://github.com/llvm/llvm-project/pull/199622