78 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc -o - < %s -mcpu=sm_70 -mattr=+ptx76 | FileCheck %s3; RUN: %if ptxas-sm_70 && ptxas-isa-7.6 %{ llc < %s -mcpu=sm_70 -mattr=+ptx76 | %ptxas-verify -arch=sm_70 %}4 5target triple = "nvptx64-unknown-cuda"6 7define i32 @bmsk_wrap(i32 %a, i32 %b) {8; CHECK-LABEL: bmsk_wrap(9; CHECK: {10; CHECK-NEXT: .reg .b32 %r<4>;11; CHECK-EMPTY:12; CHECK-NEXT: // %bb.0:13; CHECK-NEXT: ld.param.b32 %r1, [bmsk_wrap_param_0];14; CHECK-NEXT: ld.param.b32 %r2, [bmsk_wrap_param_1];15; CHECK-NEXT: bmsk.wrap.b32 %r3, %r1, %r2;16; CHECK-NEXT: st.param.b32 [func_retval0], %r3;17; CHECK-NEXT: ret;18 %c = call i32 @llvm.nvvm.bmsk.wrap(i32 %a, i32 %b)19 ret i32 %c20}21 22define i32 @bmsk_clamp(i32 %a, i32 %b) {23; CHECK-LABEL: bmsk_clamp(24; CHECK: {25; CHECK-NEXT: .reg .b32 %r<4>;26; CHECK-EMPTY:27; CHECK-NEXT: // %bb.0:28; CHECK-NEXT: ld.param.b32 %r1, [bmsk_clamp_param_0];29; CHECK-NEXT: ld.param.b32 %r2, [bmsk_clamp_param_1];30; CHECK-NEXT: bmsk.clamp.b32 %r3, %r1, %r2;31; CHECK-NEXT: st.param.b32 [func_retval0], %r3;32; CHECK-NEXT: ret;33 %c = call i32 @llvm.nvvm.bmsk.clamp(i32 %a, i32 %b)34 ret i32 %c35}36 37define i32 @bmsk_wrap_ii() {38; CHECK-LABEL: bmsk_wrap_ii(39; CHECK: {40; CHECK-NEXT: .reg .b32 %r<3>;41; CHECK-EMPTY:42; CHECK-NEXT: // %bb.0:43; CHECK-NEXT: mov.b32 %r1, 5;44; CHECK-NEXT: bmsk.wrap.b32 %r2, %r1, 6;45; CHECK-NEXT: st.param.b32 [func_retval0], %r2;46; CHECK-NEXT: ret;47 %c = call i32 @llvm.nvvm.bmsk.wrap(i32 5, i32 6)48 ret i32 %c49}50 51define i32 @bmsk_clamp_ir(i32 %a) {52; CHECK-LABEL: bmsk_clamp_ir(53; CHECK: {54; CHECK-NEXT: .reg .b32 %r<3>;55; CHECK-EMPTY:56; CHECK-NEXT: // %bb.0:57; CHECK-NEXT: ld.param.b32 %r1, [bmsk_clamp_ir_param_0];58; CHECK-NEXT: bmsk.clamp.b32 %r2, %r1, 7;59; CHECK-NEXT: st.param.b32 [func_retval0], %r2;60; CHECK-NEXT: ret;61 %c = call i32 @llvm.nvvm.bmsk.clamp(i32 %a, i32 7)62 ret i32 %c63}64 65define i32 @bmsk_wrap_ri(i32 %a) {66; CHECK-LABEL: bmsk_wrap_ri(67; CHECK: {68; CHECK-NEXT: .reg .b32 %r<3>;69; CHECK-EMPTY:70; CHECK-NEXT: // %bb.0:71; CHECK-NEXT: ld.param.b32 %r1, [bmsk_wrap_ri_param_0];72; CHECK-NEXT: bmsk.wrap.b32 %r2, 5, %r1;73; CHECK-NEXT: st.param.b32 [func_retval0], %r2;74; CHECK-NEXT: ret;75 %c = call i32 @llvm.nvvm.bmsk.wrap(i32 5, i32 %a)76 ret i32 %c77}78