correctly use unique + comments

Author: Pierre-vh

clang/docs/LanguageExtensions.rst

@@ -4432,7 +4432,8 @@ The address spaces arguments must be string literals with known values, such as:
 * ``"image"``
 
 If one or more address space name are provided, the code generator will attempt
-to emit potentially faster instructions that only fence those address spaces.
+to emit potentially faster instructions that order access to at least those
+address spaces.
 Emitting such instructions may not always be possible and the compiler is free
 to fence more aggressively.
 

clang/lib/CodeGen/CGBuiltin.cpp

@@ -18328,7 +18328,7 @@ Value *CodeGenFunction::EmitHLSLBuiltinExpr(unsigned BuiltinID,
   return nullptr;
 }
 
-void CodeGenFunction::AddAMDGCNFenceAddressSpaceMMRA(llvm::Instruction *Inst,
+void CodeGenFunction::AddAMDGPUFenceAddressSpaceMMRA(llvm::Instruction *Inst,
                                                      const CallExpr *E) {
   constexpr const char *Tag = "amdgpu-as";
 
@@ -18346,7 +18346,8 @@ void CodeGenFunction::AddAMDGCNFenceAddressSpaceMMRA(llvm::Instruction *Inst,
               "expected an address space name as a string literal");
   }
 
-  llvm::unique(MMRAs);
+  llvm::sort(MMRAs);
+  MMRAs.erase(llvm::unique(MMRAs), MMRAs.end());
   Inst->setMetadata(LLVMContext::MD_mmra, MMRAMetadata::getMD(Ctx, MMRAs));
 }
 
@@ -19022,7 +19023,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
                             EmitScalarExpr(E->getArg(1)), AO, SSID);
     FenceInst *Fence = Builder.CreateFence(AO, SSID);
     if (E->getNumArgs() > 2)
-      AddAMDGCNFenceAddressSpaceMMRA(Fence, E);
+      AddAMDGPUFenceAddressSpaceMMRA(Fence, E);
     return Fence;
   }
   case AMDGPU::BI__builtin_amdgcn_atomic_inc32:

clang/lib/CodeGen/CodeGenFunction.h

@@ -4636,7 +4636,7 @@ class CodeGenFunction : public CodeGenTypeCache {
   llvm::Value *EmitRISCVBuiltinExpr(unsigned BuiltinID, const CallExpr *E,
                                     ReturnValueSlot ReturnValue);
 
-  void AddAMDGCNFenceAddressSpaceMMRA(llvm::Instruction *Inst,
+  void AddAMDGPUFenceAddressSpaceMMRA(llvm::Instruction *Inst,
                                       const CallExpr *E);
   void ProcessOrderScopeAMDGCN(llvm::Value *Order, llvm::Value *Scope,
                                llvm::AtomicOrdering &AO,

clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp

@@ -102,12 +102,12 @@ void test_image() {
 //
 void test_mixed() {
   __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "image", "global");
-  __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "image", "local", "global");
+  __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "image", "local", "global", "image", "image");
 }
 //.
 // CHECK: [[META3]] = !{!"amdgpu-as", !"local"}
 // CHECK: [[META4]] = !{!"amdgpu-as", !"global"}
 // CHECK: [[META5]] = !{!"amdgpu-as", !"image"}
-// CHECK: [[META6]] = !{[[META5]], [[META4]]}
-// CHECK: [[META7]] = !{[[META5]], [[META3]], [[META4]]}
+// CHECK: [[META6]] = !{[[META4]], [[META5]]}
+// CHECK: [[META7]] = !{[[META4]], [[META5]], [[META3]]}
 //.

llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp

@@ -18,6 +18,7 @@
 #include "GCNSubtarget.h"
 #include "MCTargetDesc/AMDGPUMCTargetDesc.h"
 #include "llvm/ADT/BitmaskEnum.h"
+#include "llvm/ADT/StringExtras.h"
 #include "llvm/CodeGen/MachineBasicBlock.h"
 #include "llvm/CodeGen/MachineFunctionPass.h"
 #include "llvm/IR/DiagnosticInfo.h"
@@ -691,14 +692,9 @@ void diagnoseUnknownMMRAASName(const MachineInstr &MI, StringRef AS) {
   std::string Str;
   raw_string_ostream OS(Str);
   OS << "unknown address space '" << AS << "'; expected one of ";
-  bool IsFirst = true;
-  for (const auto &[Name, Val] : ASNames) {
-    if (IsFirst)
-      IsFirst = false;
-    else
-      OS << ", ";
-    OS << '\'' << Name << '\'';
-  }
+  ListSeparator LS;
+  for (const auto &[Name, Val] : ASNames)
+    OS << LS << '\'' << Name << '\'';
   DiagnosticInfoUnsupported BadTag(Fn, Str, MI.getDebugLoc(), DS_Warning);
   Fn.getContext().diagnose(BadTag);
 }