72 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 | FileCheck %s -check-prefix=CHECK3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -fp-contract=fast | FileCheck %s -check-prefix=CHECK -check-prefix=CHECK-UNSAFE4; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -fp-contract=fast | %ptxas-verify %}5; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -fp-contract=fast | %ptxas-verify %}6 7define ptx_device float @t1_f32(float %x, float %y, float %z,8; CHECK-UNSAFE-LABEL: t1_f32(9; CHECK-UNSAFE: {10; CHECK-UNSAFE-NEXT: .reg .b32 %r<8>;11; CHECK-UNSAFE-EMPTY:12; CHECK-UNSAFE-NEXT: // %bb.0:13; CHECK-UNSAFE-NEXT: ld.param.b32 %r1, [t1_f32_param_0];14; CHECK-UNSAFE-NEXT: ld.param.b32 %r2, [t1_f32_param_1];15; CHECK-UNSAFE-NEXT: ld.param.b32 %r3, [t1_f32_param_2];16; CHECK-UNSAFE-NEXT: ld.param.b32 %r4, [t1_f32_param_3];17; CHECK-UNSAFE-NEXT: ld.param.b32 %r5, [t1_f32_param_4];18; CHECK-UNSAFE-NEXT: fma.rn.f32 %r6, %r4, %r5, %r3;19; CHECK-UNSAFE-NEXT: fma.rn.f32 %r7, %r1, %r2, %r6;20; CHECK-UNSAFE-NEXT: st.param.b32 [func_retval0], %r7;21; CHECK-UNSAFE-NEXT: ret;22 float %u, float %v) {23 %a = fmul reassoc float %x, %y24 %b = fmul reassoc float %u, %v25 %c = fadd reassoc float %a, %b26 %d = fadd reassoc float %c, %z27 ret float %d28}29 30define ptx_device double @t1_f64(double %x, double %y, double %z,31; CHECK-UNSAFE-LABEL: t1_f64(32; CHECK-UNSAFE: {33; CHECK-UNSAFE-NEXT: .reg .b64 %rd<8>;34; CHECK-UNSAFE-EMPTY:35; CHECK-UNSAFE-NEXT: // %bb.0:36; CHECK-UNSAFE-NEXT: ld.param.b64 %rd1, [t1_f64_param_0];37; CHECK-UNSAFE-NEXT: ld.param.b64 %rd2, [t1_f64_param_1];38; CHECK-UNSAFE-NEXT: ld.param.b64 %rd3, [t1_f64_param_2];39; CHECK-UNSAFE-NEXT: ld.param.b64 %rd4, [t1_f64_param_3];40; CHECK-UNSAFE-NEXT: ld.param.b64 %rd5, [t1_f64_param_4];41; CHECK-UNSAFE-NEXT: fma.rn.f64 %rd6, %rd4, %rd5, %rd3;42; CHECK-UNSAFE-NEXT: fma.rn.f64 %rd7, %rd1, %rd2, %rd6;43; CHECK-UNSAFE-NEXT: st.param.b64 [func_retval0], %rd7;44; CHECK-UNSAFE-NEXT: ret;45 double %u, double %v) {46 %a = fmul reassoc double %x, %y47 %b = fmul reassoc double %u, %v48 %c = fadd reassoc double %a, %b49 %d = fadd reassoc double %c, %z50 ret double %d51}52 53define double @two_choices(double %val1, double %val2) {54; CHECK-LABEL: two_choices(55; CHECK: {56; CHECK-NEXT: .reg .b64 %rd<5>;57; CHECK-EMPTY:58; CHECK-NEXT: // %bb.0:59; CHECK-NEXT: ld.param.b64 %rd1, [two_choices_param_0];60; CHECK-NEXT: ld.param.b64 %rd2, [two_choices_param_1];61; CHECK-NEXT: mul.f64 %rd3, %rd1, %rd2;62; CHECK-NEXT: fma.rn.f64 %rd4, %rd3, %rd3, %rd3;63; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;64; CHECK-NEXT: ret;65 %1 = fmul double %val1, %val266 %2 = fmul double %1, %167 %3 = fadd double %1, %268 69 ret double %370}71 72