39 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_110a -mattr=+ptx90 | FileCheck --check-prefixes=CHECK %s5; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %}6; RUN: %if ptxas-sm_103a && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | %ptxas-verify -arch=sm_103a %}7; RUN: %if ptxas-sm_110a && ptxas-isa-9.0 %{ llc < %s -march=nvptx64 -mcpu=sm_110a -mattr=+ptx90 | %ptxas-verify -arch=sm_110a %}8 9declare void @llvm.nvvm.tcgen05.shift.down.cg1(ptr addrspace(6) %tmem_addr)10declare void @llvm.nvvm.tcgen05.shift.down.cg2(ptr addrspace(6) %tmem_addr)11 12define void @test_tcgen05_shift_cg1(ptr addrspace(6) %tmem_addr) {13; CHECK-LABEL: test_tcgen05_shift_cg1(14; CHECK: {15; CHECK-NEXT: .reg .b32 %r<2>;16; CHECK-EMPTY:17; CHECK-NEXT: // %bb.0:18; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_shift_cg1_param_0];19; CHECK-NEXT: tcgen05.shift.cta_group::1.down [%r1];20; CHECK-NEXT: ret;21 call void @llvm.nvvm.tcgen05.shift.down.cg1(ptr addrspace(6) %tmem_addr)22 23 ret void24}25 26define void @test_tcgen05_shift_cg2(ptr addrspace(6) %tmem_addr) {27; CHECK-LABEL: test_tcgen05_shift_cg2(28; CHECK: {29; CHECK-NEXT: .reg .b32 %r<2>;30; CHECK-EMPTY:31; CHECK-NEXT: // %bb.0:32; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_shift_cg2_param_0];33; CHECK-NEXT: tcgen05.shift.cta_group::2.down [%r1];34; CHECK-NEXT: ret;35 call void @llvm.nvvm.tcgen05.shift.down.cg2(ptr addrspace(6) %tmem_addr)36 37 ret void38}39