108 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 @szext_wrap_u32(i32 %a, i32 %b) {8; CHECK-LABEL: szext_wrap_u32(9; CHECK: {10; CHECK-NEXT: .reg .b32 %r<4>;11; CHECK-EMPTY:12; CHECK-NEXT: // %bb.0:13; CHECK-NEXT: ld.param.b32 %r1, [szext_wrap_u32_param_0];14; CHECK-NEXT: ld.param.b32 %r2, [szext_wrap_u32_param_1];15; CHECK-NEXT: szext.wrap.u32 %r3, %r1, %r2;16; CHECK-NEXT: st.param.b32 [func_retval0], %r3;17; CHECK-NEXT: ret;18 %c = call i32 @llvm.nvvm.zext.wrap(i32 %a, i32 %b)19 ret i32 %c20}21 22define i32 @szext_clamp_u32(i32 %a, i32 %b) {23; CHECK-LABEL: szext_clamp_u32(24; CHECK: {25; CHECK-NEXT: .reg .b32 %r<4>;26; CHECK-EMPTY:27; CHECK-NEXT: // %bb.0:28; CHECK-NEXT: ld.param.b32 %r1, [szext_clamp_u32_param_0];29; CHECK-NEXT: ld.param.b32 %r2, [szext_clamp_u32_param_1];30; CHECK-NEXT: szext.clamp.u32 %r3, %r1, %r2;31; CHECK-NEXT: st.param.b32 [func_retval0], %r3;32; CHECK-NEXT: ret;33 %c = call i32 @llvm.nvvm.zext.clamp(i32 %a, i32 %b)34 ret i32 %c35}36 37define i32 @szext_wrap_s32(i32 %a, i32 %b) {38; CHECK-LABEL: szext_wrap_s32(39; CHECK: {40; CHECK-NEXT: .reg .b32 %r<4>;41; CHECK-EMPTY:42; CHECK-NEXT: // %bb.0:43; CHECK-NEXT: ld.param.b32 %r1, [szext_wrap_s32_param_0];44; CHECK-NEXT: ld.param.b32 %r2, [szext_wrap_s32_param_1];45; CHECK-NEXT: szext.wrap.s32 %r3, %r1, %r2;46; CHECK-NEXT: st.param.b32 [func_retval0], %r3;47; CHECK-NEXT: ret;48 %c = call i32 @llvm.nvvm.sext.wrap(i32 %a, i32 %b)49 ret i32 %c50}51 52define i32 @szext_clamp_s32(i32 %a, i32 %b) {53; CHECK-LABEL: szext_clamp_s32(54; CHECK: {55; CHECK-NEXT: .reg .b32 %r<4>;56; CHECK-EMPTY:57; CHECK-NEXT: // %bb.0:58; CHECK-NEXT: ld.param.b32 %r1, [szext_clamp_s32_param_0];59; CHECK-NEXT: ld.param.b32 %r2, [szext_clamp_s32_param_1];60; CHECK-NEXT: szext.clamp.s32 %r3, %r1, %r2;61; CHECK-NEXT: st.param.b32 [func_retval0], %r3;62; CHECK-NEXT: ret;63 %c = call i32 @llvm.nvvm.sext.clamp(i32 %a, i32 %b)64 ret i32 %c65}66 67define i32 @szext_clamp_s32_ii() {68; CHECK-LABEL: szext_clamp_s32_ii(69; CHECK: {70; CHECK-NEXT: .reg .b32 %r<3>;71; CHECK-EMPTY:72; CHECK-NEXT: // %bb.0:73; CHECK-NEXT: mov.b32 %r1, 3;74; CHECK-NEXT: szext.clamp.s32 %r2, %r1, 4;75; CHECK-NEXT: st.param.b32 [func_retval0], %r2;76; CHECK-NEXT: ret;77 %c = call i32 @llvm.nvvm.sext.clamp(i32 3, i32 4)78 ret i32 %c79}80 81define i32 @szext_wrap_s32_ir(i32 %a) {82; CHECK-LABEL: szext_wrap_s32_ir(83; CHECK: {84; CHECK-NEXT: .reg .b32 %r<3>;85; CHECK-EMPTY:86; CHECK-NEXT: // %bb.0:87; CHECK-NEXT: ld.param.b32 %r1, [szext_wrap_s32_ir_param_0];88; CHECK-NEXT: szext.wrap.s32 %r2, 5, %r1;89; CHECK-NEXT: st.param.b32 [func_retval0], %r2;90; CHECK-NEXT: ret;91 %c = call i32 @llvm.nvvm.sext.wrap(i32 5, i32 %a)92 ret i32 %c93}94 95define i32 @szext_clamp_u32_ri(i32 %a) {96; CHECK-LABEL: szext_clamp_u32_ri(97; CHECK: {98; CHECK-NEXT: .reg .b32 %r<3>;99; CHECK-EMPTY:100; CHECK-NEXT: // %bb.0:101; CHECK-NEXT: ld.param.b32 %r1, [szext_clamp_u32_ri_param_0];102; CHECK-NEXT: szext.clamp.u32 %r2, %r1, 7;103; CHECK-NEXT: st.param.b32 [func_retval0], %r2;104; CHECK-NEXT: ret;105 %c = call i32 @llvm.nvvm.zext.clamp(i32 %a, i32 7)106 ret i32 %c107}108