[libclc] [libclc] update __clc_mem_fence: add MemorySemantic arg and use __builtin_amdgcn_fence for AMDGPU (PR #152275)

[Wenju He via cfe-commits]

================
@@ -0,0 +1,21 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MEM_FENCE_CLC_MEM_SEMANTIC_H__
+#define __CLC_MEM_FENCE_CLC_MEM_SEMANTIC_H__
+
+// The memory or address space to which the memory ordering is applied.
+typedef enum MemorySemantic {
+  MEMORY_PRIVATE = 1 << 0,
+  MEMORY_GLOBAL = 1 << 1,
+  MEMORY_CONSTANT = 1 << 2,
+  MEMORY_LOCAL = 1 << 3,
+  MEMORY_GENERIC = 1 << 4,
+} MemorySemantic;
----------------
wenju-he wrote:

> Is this MemorySemantic name copied from somewhere else? I don't think the name is quite right. 

Yes, it is copied from our closed source code where the enum macros are defined as `memory semantic`.

It is also copied from https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#Memory_Semantics_-id-
But the enum values here only models `SubgroupMemory`, `WorkgroupMemory` and `CrossWorkgroupMemory` in SPIR-V Memory_Semantics. So the name is likely not quite right.

> It's a mask of fenced address spaces?

Yes, IIUC it should be a mask of fenced address space for AMDGPU.

In term of the usage in Intel GPU and modeling of SPIR-V spec, it is the memory that the fence operation affects and the memory ordering applies.

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