Introduce DisjointAgents fn attribute/intrinsic property.

This patch attempts to generalise this workaround: intel#1334

Author: jchlanda

llvm/include/llvm/AsmParser/LLToken.h

@@ -188,6 +188,7 @@ enum Kind {
   kw_inalloca,
   kw_cold,
   kw_convergent,
+  kw_disjoint_agents,
   kw_dereferenceable,
   kw_dereferenceable_or_null,
   kw_disable_sanitizer_instrumentation,

llvm/include/llvm/Bitcode/LLVMBitCodes.h

@@ -672,6 +672,7 @@ enum AttributeKindCodes {
   ATTR_KIND_NO_SANITIZE_COVERAGE = 76,
   ATTR_KIND_ELEMENTTYPE = 77,
   ATTR_KIND_DISABLE_SANITIZER_INSTRUMENTATION = 78,
+  ATTR_KIND_DISJOINT_AGENTS = 79,
 };
 
 enum ComdatSelectionKindCodes {

llvm/include/llvm/IR/Attributes.td

@@ -76,6 +76,9 @@ def Cold : EnumAttr<"cold", [FnAttr]>;
 /// Can only be moved to control-equivalent blocks.
 def Convergent : EnumAttr<"convergent", [FnAttr]>;
 
+/// TODO: JKB: Document
+def DisjointAgents : EnumAttr<"disjoint_agents", [FnAttr]>;
+
 /// Marks function as being in a hot path and frequently called.
 def Hot: EnumAttr<"hot", [FnAttr]>;
 

llvm/include/llvm/IR/Function.h

@@ -589,6 +589,17 @@ class LLVM_EXTERNAL_VISIBILITY Function : public GlobalObject,
     removeFnAttr(Attribute::Convergent);
   }
 
+  /// Determine if the call has disjoint agents.
+  bool isDisjointAgents() const {
+    return hasFnAttribute(Attribute::DisjointAgents);
+  }
+  void setDisjointAgents() {
+    addFnAttr(Attribute::DisjointAgents);
+  }
+  void setNotDisjointAgents() {
+    removeFnAttr(Attribute::DisjointAgents);
+  }
+
   /// Determine if the call has sideeffects.
   bool isSpeculatable() const {
     return hasFnAttribute(Attribute::Speculatable);

llvm/include/llvm/IR/Intrinsics.td

@@ -159,6 +159,9 @@ def IntrSpeculatable : IntrinsicProperty;
 // defined by the hasSideEffects property of the TableGen Instruction class.
 def IntrHasSideEffects : IntrinsicProperty;
 
+// TODO: JKB: Document.
+def IntrDisjointAgents : IntrinsicProperty;
+
 //===----------------------------------------------------------------------===//
 // Types used by intrinsics.
 //===----------------------------------------------------------------------===//

llvm/include/llvm/IR/IntrinsicsNVVM.td

@@ -1251,7 +1251,7 @@ let TargetPrefix = "nvvm" in {
   // The builtin for "bar.sync 0" is called __syncthreads.  Unlike most of the
   // intrinsics in this file, this one is a user-facing API.
   def int_nvvm_barrier0 : GCCBuiltin<"__syncthreads">,
-      Intrinsic<[], [], [IntrConvergent]>;
+      Intrinsic<[], [], [IntrConvergent, IntrDisjointAgents]>;
   // Synchronize all threads in the CTA at barrier 'n'.
   def int_nvvm_barrier_n : GCCBuiltin<"__nvvm_bar_n">,
       Intrinsic<[], [llvm_i32_ty], [IntrConvergent]>;

llvm/include/llvm/Target/Target.td

@@ -550,6 +550,7 @@ class Instruction : InstructionEncoding {
   bit hasCtrlDep   = false;     // Does this instruction r/w ctrl-flow chains?
   bit isNotDuplicable = false;  // Is it unsafe to duplicate this instruction?
   bit isConvergent = false;     // Is this instruction convergent?
+  bit isDisjointAgents = false; // TODO: JKB: Document.
   bit isAuthenticated = false;  // Does this instruction authenticate a pointer?
   bit isAsCheapAsAMove = false; // As cheap (or cheaper) than a move instruction.
   bit hasExtraSrcRegAllocReq = false; // Sources have special regalloc requirement?

llvm/lib/Analysis/GlobalsModRef.cpp

@@ -533,18 +533,14 @@ void GlobalsAAResult::AnalyzeCallGraph(CallGraph &CG, Module &M) {
           FI.addModRefInfo(ModRefInfo::ModRef);
           if (!F->onlyAccessesArgMemory())
             FI.setMayReadAnyGlobal();
-          if (!F->isIntrinsic()) {
-            KnowNothing = true;
-            break;
-          } else if (F->getName().contains("nvvm.barrier") or
-                     F->getName().contains("nvvm.membar")) {
+          if (!F->isIntrinsic() || F->isDisjointAgents()) {
             // Even if it is an intrinsic, consider that nothing is known for
-            // NVVM barrier itrinsics to prevent illegal optimizations.
-            // This is a workaround for the bug on PTX target: barrier
-            // intrinsics are implemented as llvm intrinsics, as result there
-            // are cases when globals alias analysis can produce a result that
-            // barrier doesn't modify internal global which causes illegal
-            // reordering of memory accesses.
+            // calls that interact with disjoint agents, such as NVVM barrier
+            // itrinsics to prevent illegal optimizations.
+            // For context in PTX  barriers are implemented as llvm intrinsics,
+            // as result there are cases when globals alias analysis can
+            // produce a result that barrier doesn't modify internal global
+            // which causes illegal reordering of memory accesses.
             KnowNothing = true;
             break;
           }

llvm/lib/Bitcode/Reader/BitcodeReader.cpp

@@ -1228,6 +1228,7 @@ static uint64_t getRawAttributeMask(Attribute::AttrKind Val) {
   case Attribute::Naked:           return 1 << 24;
   case Attribute::InlineHint:      return 1 << 25;
   case Attribute::StackAlignment:  return 7 << 26;
+  case Attribute::DisjointAgents:  return 1ULL << 28;
   case Attribute::ReturnsTwice:    return 1 << 29;
   case Attribute::UWTable:         return 1 << 30;
   case Attribute::NonLazyBind:     return 1U << 31;
@@ -1393,6 +1394,8 @@ static Attribute::AttrKind getAttrFromCode(uint64_t Code) {
     return Attribute::Cold;
   case bitc::ATTR_KIND_CONVERGENT:
     return Attribute::Convergent;
+  case bitc::ATTR_KIND_DISJOINT_AGENTS:
+    return Attribute::DisjointAgents;
   case bitc::ATTR_KIND_DISABLE_SANITIZER_INSTRUMENTATION:
     return Attribute::DisableSanitizerInstrumentation;
   case bitc::ATTR_KIND_ELEMENTTYPE:

llvm/lib/Bitcode/Writer/BitcodeWriter.cpp

@@ -622,6 +622,8 @@ static uint64_t getAttrKindEncoding(Attribute::AttrKind Kind) {
     return bitc::ATTR_KIND_BY_VAL;
   case Attribute::Convergent:
     return bitc::ATTR_KIND_CONVERGENT;
+  case Attribute::DisjointAgents:
+    return bitc::ATTR_KIND_DISJOINT_AGENTS;
   case Attribute::InAlloca:
     return bitc::ATTR_KIND_IN_ALLOCA;
   case Attribute::Cold: