49 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_100a -mattr=+ptx86 | FileCheck --check-prefixes=CHECK %s3; RUN: llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | FileCheck --check-prefixes=CHECK %s4; RUN: llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | FileCheck --check-prefixes=CHECK %s5; RUN: llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | FileCheck --check-prefixes=CHECK %s6; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %}7; RUN: %if ptxas-sm_103a && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | %ptxas-verify -arch=sm_103a %}8; RUN: %if ptxas-sm_100f && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | %ptxas-verify -arch=sm_100f %}9; RUN: %if ptxas-sm_110f && ptxas-isa-9.0 %{ llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | %ptxas-verify -arch=sm_110f %}10 11declare void @llvm.nvvm.tcgen05.fence.before.thread.sync()12declare void @llvm.nvvm.tcgen05.fence.after.thread.sync()13declare void @llvm.nvvm.tcgen05.wait.ld()14declare void @llvm.nvvm.tcgen05.wait.st()15 16; CHECK-LABEL: test_tcgen05_fence17define void @test_tcgen05_fence() {18; CHECK-LABEL: test_tcgen05_fence(19; CHECK: {20; CHECK-EMPTY:21; CHECK-EMPTY:22; CHECK-NEXT: // %bb.0:23; CHECK-NEXT: tcgen05.fence::before_thread_sync;24; CHECK-NEXT: tcgen05.fence::after_thread_sync;25; CHECK-NEXT: ret;26 call void @llvm.nvvm.tcgen05.fence.before.thread.sync()27 28 call void @llvm.nvvm.tcgen05.fence.after.thread.sync()29 30 ret void31}32 33; CHECK-LABEL: test_tcgen05_wait34define void @test_tcgen05_wait() {35; CHECK-LABEL: test_tcgen05_wait(36; CHECK: {37; CHECK-EMPTY:38; CHECK-EMPTY:39; CHECK-NEXT: // %bb.0:40; CHECK-NEXT: tcgen05.wait::ld.sync.aligned;41; CHECK-NEXT: tcgen05.wait::st.sync.aligned;42; CHECK-NEXT: ret;43 call void @llvm.nvvm.tcgen05.wait.ld()44 45 call void @llvm.nvvm.tcgen05.wait.st()46 47 ret void48}49