[llvm] [NVPTX] Add intrinsics for wgmma.fence PTX instructions (PR #120523)

[Justin Fargnoli via llvm-commits]

================
@@ -0,0 +1,37 @@
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_90a -mattr=+ptx80 | FileCheck %s
+; RUN: %if ptxas-12.0 %{ llc < %s -march=nvptx64 -mcpu=sm_90a -mattr=+ptx80 | %ptxas-verify -arch=sm_90a %}
+
+target triple = "nvptx64-nvidia-cuda"
+
+declare void @llvm.nvvm.wgmma.fence.sync.aligned()
+
+define void @test_wgmma_fence_sync_aligned() {
+  ; CHECK-LABEL:  test_wgmma_fence_sync_aligned(
+  ; CHECK:        // %bb.0:
+  ; CHECK-NEXT:     wgmma.fence.sync.aligned;
+  ; CHECK-NEXT:     ret;
+  call void @llvm.nvvm.wgmma.fence.sync.aligned()
+  ret void
+}
+
+declare void @llvm.nvvm.wgmma.commit_group.sync.aligned()
+
+define void @test_wgmma_commit_group_sync_aligned() {
+  ; CHECK-LABEL:  test_wgmma_commit_group_sync_aligned(
+  ; CHECK:        // %bb.0:
+  ; CHECK-NEXT:     wgmma.commit_group.sync.aligned;
+  ; CHECK-NEXT:     ret;
+  call void @llvm.nvvm.wgmma.commit_group.sync.aligned()
+  ret void
+}
+
+declare void @llvm.nvvm.wgmma.wait_group.sync.aligned(i64)
+
+define void @test_wgmma_wait_group_sync_aligned() {
+  ; CHECK-LABEL:  test_wgmma_wait_group_sync_aligned(
+  ; CHECK:        // %bb.0:
+  ; CHECK-NEXT:     wgmma.wait_group.sync.aligned   10;
+  ; CHECK-NEXT:     ret;
+  call void @llvm.nvvm.wgmma.wait_group.sync.aligned(i64 10)
+  ret void
+}
----------------
justinfargnoli wrote:

Nit: please add a newline. 

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