94 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -verify-machineinstrs -nvptx-prec-divf32=0 | FileCheck %s --check-prefix=APPROX3; RUN: llc < %s -verify-machineinstrs -nvptx-prec-divf32=1 | FileCheck %s --check-prefix=FULL4; RUN: llc < %s -verify-machineinstrs -nvptx-prec-divf32=2 | FileCheck %s --check-prefixes=IEEE,FTZ5; RUN: llc < %s -verify-machineinstrs -nvptx-prec-divf32=3 | FileCheck %s --check-prefixes=IEEE,NOFTZ6 7target triple = "nvptx64-nvidia-cuda"8 9define float @div_ftz(float %a, float %b) "denormal-fp-math-f32" = "preserve-sign" {10; APPROX-LABEL: div_ftz(11; APPROX: {12; APPROX-NEXT: .reg .b32 %r<4>;13; APPROX-EMPTY:14; APPROX-NEXT: // %bb.0:15; APPROX-NEXT: ld.param.b32 %r1, [div_ftz_param_0];16; APPROX-NEXT: ld.param.b32 %r2, [div_ftz_param_1];17; APPROX-NEXT: div.approx.ftz.f32 %r3, %r1, %r2;18; APPROX-NEXT: st.param.b32 [func_retval0], %r3;19; APPROX-NEXT: ret;20;21; FULL-LABEL: div_ftz(22; FULL: {23; FULL-NEXT: .reg .b32 %r<4>;24; FULL-EMPTY:25; FULL-NEXT: // %bb.0:26; FULL-NEXT: ld.param.b32 %r1, [div_ftz_param_0];27; FULL-NEXT: ld.param.b32 %r2, [div_ftz_param_1];28; FULL-NEXT: div.full.ftz.f32 %r3, %r1, %r2;29; FULL-NEXT: st.param.b32 [func_retval0], %r3;30; FULL-NEXT: ret;31;32; FTZ-LABEL: div_ftz(33; FTZ: {34; FTZ-NEXT: .reg .b32 %r<4>;35; FTZ-EMPTY:36; FTZ-NEXT: // %bb.0:37; FTZ-NEXT: ld.param.b32 %r1, [div_ftz_param_0];38; FTZ-NEXT: ld.param.b32 %r2, [div_ftz_param_1];39; FTZ-NEXT: div.rn.ftz.f32 %r3, %r1, %r2;40; FTZ-NEXT: st.param.b32 [func_retval0], %r3;41; FTZ-NEXT: ret;42;43; NOFTZ-LABEL: div_ftz(44; NOFTZ: {45; NOFTZ-NEXT: .reg .b32 %r<4>;46; NOFTZ-EMPTY:47; NOFTZ-NEXT: // %bb.0:48; NOFTZ-NEXT: ld.param.b32 %r1, [div_ftz_param_0];49; NOFTZ-NEXT: ld.param.b32 %r2, [div_ftz_param_1];50; NOFTZ-NEXT: div.rn.f32 %r3, %r1, %r2;51; NOFTZ-NEXT: st.param.b32 [func_retval0], %r3;52; NOFTZ-NEXT: ret;53 %val = fdiv float %a, %b54 ret float %val55}56 57 58define float @div(float %a, float %b) {59; APPROX-LABEL: div(60; APPROX: {61; APPROX-NEXT: .reg .b32 %r<4>;62; APPROX-EMPTY:63; APPROX-NEXT: // %bb.0:64; APPROX-NEXT: ld.param.b32 %r1, [div_param_0];65; APPROX-NEXT: ld.param.b32 %r2, [div_param_1];66; APPROX-NEXT: div.approx.f32 %r3, %r1, %r2;67; APPROX-NEXT: st.param.b32 [func_retval0], %r3;68; APPROX-NEXT: ret;69;70; FULL-LABEL: div(71; FULL: {72; FULL-NEXT: .reg .b32 %r<4>;73; FULL-EMPTY:74; FULL-NEXT: // %bb.0:75; FULL-NEXT: ld.param.b32 %r1, [div_param_0];76; FULL-NEXT: ld.param.b32 %r2, [div_param_1];77; FULL-NEXT: div.full.f32 %r3, %r1, %r2;78; FULL-NEXT: st.param.b32 [func_retval0], %r3;79; FULL-NEXT: ret;80;81; IEEE-LABEL: div(82; IEEE: {83; IEEE-NEXT: .reg .b32 %r<4>;84; IEEE-EMPTY:85; IEEE-NEXT: // %bb.0:86; IEEE-NEXT: ld.param.b32 %r1, [div_param_0];87; IEEE-NEXT: ld.param.b32 %r2, [div_param_1];88; IEEE-NEXT: div.rn.f32 %r3, %r1, %r2;89; IEEE-NEXT: st.param.b32 [func_retval0], %r3;90; IEEE-NEXT: ret;91 %val = fdiv float %a, %b92 ret float %val93}94