113 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_100 -fp-contract=fast | FileCheck %s --check-prefixes=CHECK,FAST3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 | FileCheck %s --check-prefixes=CHECK,DEFAULT4; RUN: %if ptxas-sm_100 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -fp-contract=fast | %ptxas-verify -arch sm_100 %}5; RUN: %if ptxas-sm_100 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 | %ptxas-verify -arch sm_100 %}6 7target triple = "nvptx64-unknown-cuda"8 9;; FAST-LABEL: @t010;; DEFAULT-LABEL: @t011define <2 x float> @t0(<2 x float> %a, <2 x float> %b, <2 x float> %c) {12; FAST-LABEL: t0(13; FAST: {14; FAST-NEXT: .reg .b64 %rd<5>;15; FAST-EMPTY:16; FAST-NEXT: // %bb.0:17; FAST-NEXT: ld.param.b64 %rd1, [t0_param_0];18; FAST-NEXT: ld.param.b64 %rd2, [t0_param_1];19; FAST-NEXT: ld.param.b64 %rd3, [t0_param_2];20; FAST-NEXT: fma.rn.f32x2 %rd4, %rd1, %rd2, %rd3;21; FAST-NEXT: st.param.b64 [func_retval0], %rd4;22; FAST-NEXT: ret;23;24; DEFAULT-LABEL: t0(25; DEFAULT: {26; DEFAULT-NEXT: .reg .b64 %rd<6>;27; DEFAULT-EMPTY:28; DEFAULT-NEXT: // %bb.0:29; DEFAULT-NEXT: ld.param.b64 %rd1, [t0_param_0];30; DEFAULT-NEXT: ld.param.b64 %rd2, [t0_param_1];31; DEFAULT-NEXT: mul.rn.f32x2 %rd3, %rd1, %rd2;32; DEFAULT-NEXT: ld.param.b64 %rd4, [t0_param_2];33; DEFAULT-NEXT: add.rn.f32x2 %rd5, %rd3, %rd4;34; DEFAULT-NEXT: st.param.b64 [func_retval0], %rd5;35; DEFAULT-NEXT: ret;36 %v0 = fmul <2 x float> %a, %b37 %v1 = fadd <2 x float> %v0, %c38 ret <2 x float> %v139}40 41;; We cannot form an fma here, but make sure we explicitly emit add.rn.f32x242;; to prevent ptxas from fusing this with anything else.43define <2 x float> @t1(<2 x float> %a, <2 x float> %b) {44; FAST-LABEL: t1(45; FAST: {46; FAST-NEXT: .reg .b64 %rd<6>;47; FAST-EMPTY:48; FAST-NEXT: // %bb.0:49; FAST-NEXT: ld.param.b64 %rd1, [t1_param_0];50; FAST-NEXT: ld.param.b64 %rd2, [t1_param_1];51; FAST-NEXT: add.f32x2 %rd3, %rd1, %rd2;52; FAST-NEXT: sub.f32x2 %rd4, %rd1, %rd2;53; FAST-NEXT: mul.f32x2 %rd5, %rd3, %rd4;54; FAST-NEXT: st.param.b64 [func_retval0], %rd5;55; FAST-NEXT: ret;56;57; DEFAULT-LABEL: t1(58; DEFAULT: {59; DEFAULT-NEXT: .reg .b64 %rd<6>;60; DEFAULT-EMPTY:61; DEFAULT-NEXT: // %bb.0:62; DEFAULT-NEXT: ld.param.b64 %rd1, [t1_param_0];63; DEFAULT-NEXT: ld.param.b64 %rd2, [t1_param_1];64; DEFAULT-NEXT: add.rn.f32x2 %rd3, %rd1, %rd2;65; DEFAULT-NEXT: sub.rn.f32x2 %rd4, %rd1, %rd2;66; DEFAULT-NEXT: mul.rn.f32x2 %rd5, %rd3, %rd4;67; DEFAULT-NEXT: st.param.b64 [func_retval0], %rd5;68; DEFAULT-NEXT: ret;69 %v1 = fadd <2 x float> %a, %b70 %v2 = fsub <2 x float> %a, %b71 %v3 = fmul <2 x float> %v1, %v272 ret <2 x float> %v373}74 75;; Make sure we generate the non ".rn" version when the "contract" flag is76;; present on the instructions77define <2 x float> @t2(<2 x float> %a, <2 x float> %b) {78; CHECK-LABEL: t2(79; CHECK: {80; CHECK-NEXT: .reg .b64 %rd<6>;81; CHECK-EMPTY:82; CHECK-NEXT: // %bb.0:83; CHECK-NEXT: ld.param.b64 %rd1, [t2_param_0];84; CHECK-NEXT: ld.param.b64 %rd2, [t2_param_1];85; CHECK-NEXT: add.f32x2 %rd3, %rd1, %rd2;86; CHECK-NEXT: sub.f32x2 %rd4, %rd1, %rd2;87; CHECK-NEXT: mul.f32x2 %rd5, %rd3, %rd4;88; CHECK-NEXT: st.param.b64 [func_retval0], %rd5;89; CHECK-NEXT: ret;90 %v1 = fadd contract <2 x float> %a, %b91 %v2 = fsub contract <2 x float> %a, %b92 %v3 = fmul contract <2 x float> %v1, %v293 ret <2 x float> %v394}95 96;; Make sure we always fold to fma when the "contract" flag is present97define <2 x float> @t3(<2 x float> %a, <2 x float> %b, <2 x float> %c) {98; CHECK-LABEL: t3(99; CHECK: {100; CHECK-NEXT: .reg .b64 %rd<5>;101; CHECK-EMPTY:102; CHECK-NEXT: // %bb.0:103; CHECK-NEXT: ld.param.b64 %rd1, [t3_param_0];104; CHECK-NEXT: ld.param.b64 %rd2, [t3_param_1];105; CHECK-NEXT: ld.param.b64 %rd3, [t3_param_2];106; CHECK-NEXT: fma.rn.f32x2 %rd4, %rd1, %rd2, %rd3;107; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;108; CHECK-NEXT: ret;109 %v0 = fmul contract <2 x float> %a, %b110 %v1 = fadd contract <2 x float> %v0, %c111 ret <2 x float> %v1112}113