111 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 42; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_50 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_50 | %ptxas-verify %}4 5define i16 @test_sad_i16(i16 %x, i16 %y, i16 %z) {6; CHECK-LABEL: test_sad_i16(7; CHECK: {8; CHECK-NEXT: .reg .b16 %rs<5>;9; CHECK-NEXT: .reg .b32 %r<2>;10; CHECK-EMPTY:11; CHECK-NEXT: // %bb.0:12; CHECK-NEXT: ld.param.b16 %rs1, [test_sad_i16_param_0];13; CHECK-NEXT: ld.param.b16 %rs2, [test_sad_i16_param_1];14; CHECK-NEXT: ld.param.b16 %rs3, [test_sad_i16_param_2];15; CHECK-NEXT: sad.s16 %rs4, %rs1, %rs2, %rs3;16; CHECK-NEXT: cvt.u32.u16 %r1, %rs4;17; CHECK-NEXT: st.param.b32 [func_retval0], %r1;18; CHECK-NEXT: ret;19 %1 = call i16 @llvm.nvvm.sad.s(i16 %x, i16 %y, i16 %z)20 ret i16 %121}22 23define i16 @test_sad_u16(i16 %x, i16 %y, i16 %z) {24; CHECK-LABEL: test_sad_u16(25; CHECK: {26; CHECK-NEXT: .reg .b16 %rs<5>;27; CHECK-NEXT: .reg .b32 %r<2>;28; CHECK-EMPTY:29; CHECK-NEXT: // %bb.0:30; CHECK-NEXT: ld.param.b16 %rs1, [test_sad_u16_param_0];31; CHECK-NEXT: ld.param.b16 %rs2, [test_sad_u16_param_1];32; CHECK-NEXT: ld.param.b16 %rs3, [test_sad_u16_param_2];33; CHECK-NEXT: sad.u16 %rs4, %rs1, %rs2, %rs3;34; CHECK-NEXT: cvt.u32.u16 %r1, %rs4;35; CHECK-NEXT: st.param.b32 [func_retval0], %r1;36; CHECK-NEXT: ret;37 %1 = call i16 @llvm.nvvm.sad.us(i16 %x, i16 %y, i16 %z)38 ret i16 %139}40 41define i32 @test_sad_i32(i32 %x, i32 %y, i32 %z) {42; CHECK-LABEL: test_sad_i32(43; CHECK: {44; CHECK-NEXT: .reg .b32 %r<5>;45; CHECK-EMPTY:46; CHECK-NEXT: // %bb.0:47; CHECK-NEXT: ld.param.b32 %r1, [test_sad_i32_param_0];48; CHECK-NEXT: ld.param.b32 %r2, [test_sad_i32_param_1];49; CHECK-NEXT: ld.param.b32 %r3, [test_sad_i32_param_2];50; CHECK-NEXT: sad.s32 %r4, %r1, %r2, %r3;51; CHECK-NEXT: st.param.b32 [func_retval0], %r4;52; CHECK-NEXT: ret;53 %1 = call i32 @llvm.nvvm.sad.i(i32 %x, i32 %y, i32 %z)54 ret i32 %155}56 57define i32 @test_sad_u32(i32 %x, i32 %y, i32 %z) {58; CHECK-LABEL: test_sad_u32(59; CHECK: {60; CHECK-NEXT: .reg .b32 %r<5>;61; CHECK-EMPTY:62; CHECK-NEXT: // %bb.0:63; CHECK-NEXT: ld.param.b32 %r1, [test_sad_u32_param_0];64; CHECK-NEXT: ld.param.b32 %r2, [test_sad_u32_param_1];65; CHECK-NEXT: ld.param.b32 %r3, [test_sad_u32_param_2];66; CHECK-NEXT: sad.u32 %r4, %r1, %r2, %r3;67; CHECK-NEXT: st.param.b32 [func_retval0], %r4;68; CHECK-NEXT: ret;69 %1 = call i32 @llvm.nvvm.sad.ui(i32 %x, i32 %y, i32 %z)70 ret i32 %171}72 73define i64 @test_sad_i64(i64 %x, i64 %y, i64 %z) {74; CHECK-LABEL: test_sad_i64(75; CHECK: {76; CHECK-NEXT: .reg .b64 %rd<5>;77; CHECK-EMPTY:78; CHECK-NEXT: // %bb.0:79; CHECK-NEXT: ld.param.b64 %rd1, [test_sad_i64_param_0];80; CHECK-NEXT: ld.param.b64 %rd2, [test_sad_i64_param_1];81; CHECK-NEXT: ld.param.b64 %rd3, [test_sad_i64_param_2];82; CHECK-NEXT: sad.s64 %rd4, %rd1, %rd2, %rd3;83; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;84; CHECK-NEXT: ret;85 %1 = call i64 @llvm.nvvm.sad.ll(i64 %x, i64 %y, i64 %z)86 ret i64 %187}88 89define i64 @test_sad_u64(i64 %x, i64 %y, i64 %z) {90; CHECK-LABEL: test_sad_u64(91; CHECK: {92; CHECK-NEXT: .reg .b64 %rd<5>;93; CHECK-EMPTY:94; CHECK-NEXT: // %bb.0:95; CHECK-NEXT: ld.param.b64 %rd1, [test_sad_u64_param_0];96; CHECK-NEXT: ld.param.b64 %rd2, [test_sad_u64_param_1];97; CHECK-NEXT: ld.param.b64 %rd3, [test_sad_u64_param_2];98; CHECK-NEXT: sad.u64 %rd4, %rd1, %rd2, %rd3;99; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;100; CHECK-NEXT: ret;101 %1 = call i64 @llvm.nvvm.sad.ull(i64 %x, i64 %y, i64 %z)102 ret i64 %1103}104 105declare i16 @llvm.nvvm.sad.s(i16, i16, i16)106declare i16 @llvm.nvvm.sad.us(i16, i16, i16)107declare i32 @llvm.nvvm.sad.i(i32, i32, i32)108declare i32 @llvm.nvvm.sad.ui(i32, i32, i32)109declare i64 @llvm.nvvm.sad.ll(i64, i64, i64)110declare i64 @llvm.nvvm.sad.ull(i64, i64, i64)111