[AMD] Add memrealtime to gfx11/gfx12 (#7357)

The semantics of s_sendmsg_rtn_b64 should be identical to
s_memrealtime, so we can just reuse the function.
This was changed in gfx11. For details, see
https://www.amd.com/content/dam/amd/en/documents/radeon-tech-docs/instruction-set-architectures/rdna3-shader-instruction-set-architecture-feb-2023_0.pdf

This fixes test_core.py::test_globaltimer on gfx11, gfx12 after
triton-lang/triton#7282

---------

Co-authored-by: Paul Trojahn <[email protected]>

Author: ptrojahn

python/test/unit/language/test_core.py

@@ -5982,7 +5982,11 @@ def kernel(Out1, Out2, func: tl.constexpr):
     if is_cuda():
         assert h.asm["ptx"].count("%globaltimer") == 2
     else:
-        assert h.asm["amdgcn"].count("s_memrealtime") == 2
+        target_arch = triton.runtime.driver.active.get_current_target().arch
+        if "gfx11" in target_arch or "gfx12" in target_arch:
+            assert h.asm["amdgcn"].count("s_sendmsg_rtn_b64") == 2
+        else:
+            assert h.asm["amdgcn"].count("s_memrealtime") == 2
 
 
 def test_smid(device):

third_party/amd/language/hip/utils.py

@@ -6,15 +6,30 @@ def memrealtime(_semantic=None):
     """
     Returns a 64-bit real time-counter value
     """
-    return core.inline_asm_elementwise(
-        """
-        s_memrealtime $0
-        s_waitcnt vmcnt(0)
-        """,
-        "=r",
-        [],
-        dtype=core.int64,
-        is_pure=False,
-        pack=1,
-        _semantic=_semantic,
-    )
+    target_arch = _semantic.builder.options.arch
+    if 'gfx11' in target_arch or 'gfx12' in target_arch:
+        return core.inline_asm_elementwise(
+            """
+            s_sendmsg_rtn_b64 $0, sendmsg(MSG_RTN_GET_REALTIME)
+            s_waitcnt lgkmcnt(0)
+            """,
+            "=r",
+            [],
+            dtype=core.int64,
+            is_pure=False,
+            pack=1,
+            _semantic=_semantic,
+        )
+    else:
+        return core.inline_asm_elementwise(
+            """
+            s_memrealtime $0
+            s_waitcnt vmcnt(0)
+            """,
+            "=r",
+            [],
+            dtype=core.int64,
+            is_pure=False,
+            pack=1,
+            _semantic=_semantic,
+        )