127 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -fp-contract=fast -verify-machineinstrs | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -fp-contract=fast -verify-machineinstrs | %ptxas-verify %}4 5declare float @dummy_f32(float, float) #06declare double @dummy_f64(double, double) #07 8define ptx_device float @t1_f32(float %x, float %y, float %z) {9; CHECK-LABEL: t1_f32(10; CHECK: {11; CHECK-NEXT: .reg .b32 %r<5>;12; CHECK-EMPTY:13; CHECK-NEXT: // %bb.0:14; CHECK-NEXT: ld.param.b32 %r1, [t1_f32_param_0];15; CHECK-NEXT: ld.param.b32 %r2, [t1_f32_param_1];16; CHECK-NEXT: ld.param.b32 %r3, [t1_f32_param_2];17; CHECK-NEXT: fma.rn.f32 %r4, %r1, %r2, %r3;18; CHECK-NEXT: st.param.b32 [func_retval0], %r4;19; CHECK-NEXT: ret;20 %a = fmul float %x, %y21 %b = fadd float %a, %z22 ret float %b23}24 25define ptx_device float @t2_f32(float %x, float %y, float %z, float %w) {26; CHECK-LABEL: t2_f32(27; CHECK: {28; CHECK-NEXT: .reg .b32 %r<8>;29; CHECK-EMPTY:30; CHECK-NEXT: // %bb.0:31; CHECK-NEXT: ld.param.b32 %r1, [t2_f32_param_0];32; CHECK-NEXT: ld.param.b32 %r2, [t2_f32_param_1];33; CHECK-NEXT: ld.param.b32 %r3, [t2_f32_param_2];34; CHECK-NEXT: fma.rn.f32 %r4, %r1, %r2, %r3;35; CHECK-NEXT: ld.param.b32 %r5, [t2_f32_param_3];36; CHECK-NEXT: fma.rn.f32 %r6, %r1, %r2, %r5;37; CHECK-NEXT: { // callseq 0, 038; CHECK-NEXT: .param .b32 param0;39; CHECK-NEXT: .param .b32 param1;40; CHECK-NEXT: .param .b32 retval0;41; CHECK-NEXT: st.param.b32 [param1], %r6;42; CHECK-NEXT: st.param.b32 [param0], %r4;43; CHECK-NEXT: call.uni (retval0), dummy_f32, (param0, param1);44; CHECK-NEXT: ld.param.b32 %r7, [retval0];45; CHECK-NEXT: } // callseq 046; CHECK-NEXT: st.param.b32 [func_retval0], %r7;47; CHECK-NEXT: ret;48 %a = fmul float %x, %y49 %b = fadd float %a, %z50 %c = fadd float %a, %w51 %d = call float @dummy_f32(float %b, float %c)52 ret float %d53}54 55define ptx_device double @t1_f64(double %x, double %y, double %z) {56; CHECK-LABEL: t1_f64(57; CHECK: {58; CHECK-NEXT: .reg .b64 %rd<5>;59; CHECK-EMPTY:60; CHECK-NEXT: // %bb.0:61; CHECK-NEXT: ld.param.b64 %rd1, [t1_f64_param_0];62; CHECK-NEXT: ld.param.b64 %rd2, [t1_f64_param_1];63; CHECK-NEXT: ld.param.b64 %rd3, [t1_f64_param_2];64; CHECK-NEXT: fma.rn.f64 %rd4, %rd1, %rd2, %rd3;65; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;66; CHECK-NEXT: ret;67 %a = fmul double %x, %y68 %b = fadd double %a, %z69 ret double %b70}71 72define ptx_device double @t2_f64(double %x, double %y, double %z, double %w) {73; CHECK-LABEL: t2_f64(74; CHECK: {75; CHECK-NEXT: .reg .b64 %rd<8>;76; CHECK-EMPTY:77; CHECK-NEXT: // %bb.0:78; CHECK-NEXT: ld.param.b64 %rd1, [t2_f64_param_0];79; CHECK-NEXT: ld.param.b64 %rd2, [t2_f64_param_1];80; CHECK-NEXT: ld.param.b64 %rd3, [t2_f64_param_2];81; CHECK-NEXT: fma.rn.f64 %rd4, %rd1, %rd2, %rd3;82; CHECK-NEXT: ld.param.b64 %rd5, [t2_f64_param_3];83; CHECK-NEXT: fma.rn.f64 %rd6, %rd1, %rd2, %rd5;84; CHECK-NEXT: { // callseq 1, 085; CHECK-NEXT: .param .b64 param0;86; CHECK-NEXT: .param .b64 param1;87; CHECK-NEXT: .param .b64 retval0;88; CHECK-NEXT: st.param.b64 [param1], %rd6;89; CHECK-NEXT: st.param.b64 [param0], %rd4;90; CHECK-NEXT: call.uni (retval0), dummy_f64, (param0, param1);91; CHECK-NEXT: ld.param.b64 %rd7, [retval0];92; CHECK-NEXT: } // callseq 193; CHECK-NEXT: st.param.b64 [func_retval0], %rd7;94; CHECK-NEXT: ret;95 %a = fmul double %x, %y96 %b = fadd double %a, %z97 %c = fadd double %a, %w98 %d = call double @dummy_f64(double %b, double %c)99 ret double %d100}101 102define ptx_device float @f32_iir(float %x) {103; CHECK-LABEL: f32_iir(104; CHECK: {105; CHECK-NEXT: .reg .b32 %r<3>;106; CHECK-EMPTY:107; CHECK-NEXT: // %bb.0:108; CHECK-NEXT: ld.param.b32 %r1, [f32_iir_param_0];109; CHECK-NEXT: fma.rn.f32 %r2, 0f52E8D4A5, 0f4A52FC54, %r1;110; CHECK-NEXT: st.param.b32 [func_retval0], %r2;111; CHECK-NEXT: ret;112 %r = call float @llvm.fma.f32(float 499999997952.0, float 3456789.0, float %x)113 ret float %r114}115 116define ptx_device float @f32_iii(float %x) {117; CHECK-LABEL: f32_iii(118; CHECK: {119; CHECK-EMPTY:120; CHECK-EMPTY:121; CHECK-NEXT: // %bb.0:122; CHECK-NEXT: st.param.b32 [func_retval0], 1092616192;123; CHECK-NEXT: ret;124 %r = call float @llvm.fma.f32(float 2.0, float 3.0, float 4.0)125 ret float %r126}127