48 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -march=nvptx64 -mcpu=sm_90a -mattr=+ptx80 | FileCheck %s3; RUN: %if ptxas-sm_90a && ptxas-isa-8.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 void19}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 void33}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 void47}48