[clang] [llvm] [AMDGPU] Add OpenCL-specific fence address space masks (PR #78572)

[Matt Arsenault via cfe-commits]

================
@@ -18319,6 +18320,26 @@ Value *CodeGenFunction::EmitHLSLBuiltinExpr(unsigned BuiltinID,
   return nullptr;
 }
 
+void CodeGenFunction::AddAMDGCNAddressSpaceMMRA(llvm::Instruction *Inst,
+                                                llvm::Value *ASMask) {
+  constexpr const char *Tag = "opencl-fence-mem";
+
+  uint64_t Mask = cast<llvm::ConstantInt>(ASMask)->getZExtValue();
+  if (Mask == 0)
+    return;
+
+  // 3 bits can be set: local, global, image in that order.
+  LLVMContext &Ctx = Inst->getContext();
+  SmallVector<MMRAMetadata::TagT, 3> MMRAs;
+  if (Mask & (1 << 0))
----------------
arsenm wrote:

Space separated is weird. I meant more __builtin_amdgcn_something_fence("somesyncscope", ordering, "addrspace0", "addrspace1", ...) 

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