190 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; Check that various LLVM idioms get lowered to NVPTX as expected.3 4; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 | FileCheck %s5; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s6; RUN: %if ptxas-ptr32 %{ llc < %s -mtriple=nvptx -mcpu=sm_20 | %ptxas-verify %}7; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}8 9%struct.S16 = type { i16, i16 }10%struct.S32 = type { i32, i32 }11 12define i16 @abs_i16(i16 %a) {13; CHECK-LABEL: abs_i16(14; CHECK: {15; CHECK-NEXT: .reg .b16 %rs<3>;16; CHECK-NEXT: .reg .b32 %r<2>;17; CHECK-EMPTY:18; CHECK-NEXT: // %bb.0:19; CHECK-NEXT: ld.param.b16 %rs1, [abs_i16_param_0];20; CHECK-NEXT: abs.s16 %rs2, %rs1;21; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;22; CHECK-NEXT: st.param.b32 [func_retval0], %r1;23; CHECK-NEXT: ret;24 %neg = sub i16 0, %a25 %abs.cond = icmp sge i16 %a, 026 %abs = select i1 %abs.cond, i16 %a, i16 %neg27 ret i16 %abs28}29 30define i32 @abs_i32(i32 %a) {31; CHECK-LABEL: abs_i32(32; CHECK: {33; CHECK-NEXT: .reg .b32 %r<3>;34; CHECK-EMPTY:35; CHECK-NEXT: // %bb.0:36; CHECK-NEXT: ld.param.b32 %r1, [abs_i32_param_0];37; CHECK-NEXT: abs.s32 %r2, %r1;38; CHECK-NEXT: st.param.b32 [func_retval0], %r2;39; CHECK-NEXT: ret;40 %neg = sub i32 0, %a41 %abs.cond = icmp sge i32 %a, 042 %abs = select i1 %abs.cond, i32 %a, i32 %neg43 ret i32 %abs44}45 46define i64 @abs_i64(i64 %a) {47; CHECK-LABEL: abs_i64(48; CHECK: {49; CHECK-NEXT: .reg .b64 %rd<3>;50; CHECK-EMPTY:51; CHECK-NEXT: // %bb.0:52; CHECK-NEXT: ld.param.b64 %rd1, [abs_i64_param_0];53; CHECK-NEXT: abs.s64 %rd2, %rd1;54; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;55; CHECK-NEXT: ret;56 %neg = sub i64 0, %a57 %abs.cond = icmp sge i64 %a, 058 %abs = select i1 %abs.cond, i64 %a, i64 %neg59 ret i64 %abs60}61 62define %struct.S16 @i32_to_2xi16(i32 noundef %in) {63; CHECK-LABEL: i32_to_2xi16(64; CHECK: {65; CHECK-NEXT: .reg .b32 %r<3>;66; CHECK-EMPTY:67; CHECK-NEXT: // %bb.0:68; CHECK-NEXT: ld.param.b32 %r1, [i32_to_2xi16_param_0];69; CHECK-NEXT: shr.u32 %r2, %r1, 16;70; CHECK-NEXT: st.param.b16 [func_retval0], %r1;71; CHECK-NEXT: st.param.b16 [func_retval0+2], %r2;72; CHECK-NEXT: ret;73 %low = trunc i32 %in to i1674 %high32 = lshr i32 %in, 1675 %high = trunc i32 %high32 to i1676 %s1 = insertvalue %struct.S16 poison, i16 %low, 077 %s = insertvalue %struct.S16 %s1, i16 %high, 178 ret %struct.S16 %s79}80 81; Same as above, but with rearranged order of low/high parts.82define %struct.S16 @i32_to_2xi16_lh(i32 noundef %in) {83; CHECK-LABEL: i32_to_2xi16_lh(84; CHECK: {85; CHECK-NEXT: .reg .b32 %r<3>;86; CHECK-EMPTY:87; CHECK-NEXT: // %bb.0:88; CHECK-NEXT: ld.param.b32 %r1, [i32_to_2xi16_lh_param_0];89; CHECK-NEXT: shr.u32 %r2, %r1, 16;90; CHECK-NEXT: st.param.b16 [func_retval0], %r1;91; CHECK-NEXT: st.param.b16 [func_retval0+2], %r2;92; CHECK-NEXT: ret;93 %high32 = lshr i32 %in, 1694 %high = trunc i32 %high32 to i1695 %low = trunc i32 %in to i1696 %s1 = insertvalue %struct.S16 poison, i16 %low, 097 %s = insertvalue %struct.S16 %s1, i16 %high, 198 ret %struct.S16 %s99}100 101 102define %struct.S16 @i32_to_2xi16_not(i32 noundef %in) {103; CHECK-LABEL: i32_to_2xi16_not(104; CHECK: {105; CHECK-NEXT: .reg .b32 %r<3>;106; CHECK-EMPTY:107; CHECK-NEXT: // %bb.0:108; CHECK-NEXT: ld.param.b32 %r1, [i32_to_2xi16_not_param_0];109; CHECK-NEXT: shr.u32 %r2, %r1, 15;110; CHECK-NEXT: st.param.b16 [func_retval0], %r1;111; CHECK-NEXT: st.param.b16 [func_retval0+2], %r2;112; CHECK-NEXT: ret;113 %low = trunc i32 %in to i16114 ; Shift by any value other than 16 blocks the conversiopn to mov.115 %high32 = lshr i32 %in, 15116 %high = trunc i32 %high32 to i16117 %s1 = insertvalue %struct.S16 poison, i16 %low, 0118 %s = insertvalue %struct.S16 %s1, i16 %high, 1119 ret %struct.S16 %s120}121 122define %struct.S32 @i64_to_2xi32(i64 noundef %in) {123; CHECK-LABEL: i64_to_2xi32(124; CHECK: {125; CHECK-NEXT: .reg .b64 %rd<3>;126; CHECK-EMPTY:127; CHECK-NEXT: // %bb.0:128; CHECK-NEXT: ld.param.b64 %rd1, [i64_to_2xi32_param_0];129; CHECK-NEXT: shr.u64 %rd2, %rd1, 32;130; CHECK-NEXT: st.param.b32 [func_retval0], %rd1;131; CHECK-NEXT: st.param.b32 [func_retval0+4], %rd2;132; CHECK-NEXT: ret;133 %low = trunc i64 %in to i32134 %high64 = lshr i64 %in, 32135 %high = trunc i64 %high64 to i32136 %s1 = insertvalue %struct.S32 poison, i32 %low, 0137 %s = insertvalue %struct.S32 %s1, i32 %high, 1138 ret %struct.S32 %s139}140 141define %struct.S32 @i64_to_2xi32_not(i64 noundef %in) {142; CHECK-LABEL: i64_to_2xi32_not(143; CHECK: {144; CHECK-NEXT: .reg .b64 %rd<3>;145; CHECK-EMPTY:146; CHECK-NEXT: // %bb.0:147; CHECK-NEXT: ld.param.b64 %rd1, [i64_to_2xi32_not_param_0];148; CHECK-NEXT: shr.u64 %rd2, %rd1, 31;149; CHECK-NEXT: st.param.b32 [func_retval0], %rd1;150; CHECK-NEXT: st.param.b32 [func_retval0+4], %rd2;151; CHECK-NEXT: ret;152 %low = trunc i64 %in to i32153 ; Shift by any value other than 32 blocks the conversiopn to mov.154 %high64 = lshr i64 %in, 31155 %high = trunc i64 %high64 to i32156 %s1 = insertvalue %struct.S32 poison, i32 %low, 0157 %s = insertvalue %struct.S32 %s1, i32 %high, 1158 ret %struct.S32 %s159}160 161; Make sure we do not get confused when our input itself is [al]shr.162define %struct.S16 @i32_to_2xi16_shr(i32 noundef %i){163; CHECK-LABEL: i32_to_2xi16_shr(164; CHECK: {165; CHECK-NEXT: .reg .b32 %r<4>;166; CHECK-EMPTY:167; CHECK-NEXT: // %bb.0:168; CHECK-NEXT: ld.param.b32 %r1, [i32_to_2xi16_shr_param_0];169; CHECK-NEXT: { // callseq 0, 0170; CHECK-NEXT: .param .b32 param0;171; CHECK-NEXT: st.param.b32 [param0], %r1;172; CHECK-NEXT: call.uni escape_int, (param0);173; CHECK-NEXT: } // callseq 0174; CHECK-NEXT: shr.s32 %r2, %r1, 16;175; CHECK-NEXT: shr.u32 %r3, %r2, 16;176; CHECK-NEXT: st.param.b16 [func_retval0+2], %r3;177; CHECK-NEXT: st.param.b16 [func_retval0], %r2;178; CHECK-NEXT: ret;179 call void @escape_int(i32 %i); // Force %i to be loaded completely.180 %i1 = ashr i32 %i, 16181 %l = trunc i32 %i1 to i16182 %h32 = ashr i32 %i1, 16183 %h = trunc i32 %h32 to i16184 %s0 = insertvalue %struct.S16 poison, i16 %l, 0185 %s1 = insertvalue %struct.S16 %s0, i16 %h, 1186 ret %struct.S16 %s1187}188declare dso_local void @escape_int(i32 noundef)189 190