1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 2; RUN: llc < %s -march=nvptx64 -mcpu=sm_90a -mattr=+ptx80 | FileCheck %s 3; RUN: %if ptxas-12.0 %{ llc < %s -march=nvptx64 -mcpu=sm_90a -mattr=+ptx80 | %ptxas-verify -arch=sm_90a %} 4 5target triple = "nvptx64-nvidia-cuda" 6 7declare void @llvm.nvvm.wgmma.fence.sync.aligned() 8 9define void @test_wgmma_fence_sync_aligned() { 10; CHECK-LABEL: test_wgmma_fence_sync_aligned( 11; CHECK: { 12; CHECK-EMPTY: 13; CHECK-EMPTY: 14; CHECK-NEXT: // %bb.0: 15; CHECK-NEXT: wgmma.fence.sync.aligned; 16; CHECK-NEXT: ret; 17 call void @llvm.nvvm.wgmma.fence.sync.aligned() 18 ret void 19} 20 21declare void @llvm.nvvm.wgmma.commit_group.sync.aligned() 22 23define void @test_wgmma_commit_group_sync_aligned() { 24; CHECK-LABEL: test_wgmma_commit_group_sync_aligned( 25; CHECK: { 26; CHECK-EMPTY: 27; CHECK-EMPTY: 28; CHECK-NEXT: // %bb.0: 29; CHECK-NEXT: wgmma.commit_group.sync.aligned; 30; CHECK-NEXT: ret; 31 call void @llvm.nvvm.wgmma.commit_group.sync.aligned() 32 ret void 33} 34 35declare void @llvm.nvvm.wgmma.wait_group.sync.aligned(i64) 36 37define void @test_wgmma_wait_group_sync_aligned() { 38; CHECK-LABEL: test_wgmma_wait_group_sync_aligned( 39; CHECK: { 40; CHECK-EMPTY: 41; CHECK-EMPTY: 42; CHECK-NEXT: // %bb.0: 43; CHECK-NEXT: wgmma.wait_group.sync.aligned 10; 44; CHECK-NEXT: ret; 45 call void @llvm.nvvm.wgmma.wait_group.sync.aligned(i64 10) 46 ret void 47} 48