[NVPTX] add f32x2 version of fp-contract test

Author: Prince781

llvm/test/CodeGen/NVPTX/fp-contract-f32x2.ll

@@ -0,0 +1,112 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -fp-contract=fast | FileCheck %s --check-prefixes=CHECK,FAST
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 | FileCheck %s --check-prefixes=CHECK,DEFAULT
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -fp-contract=fast | %ptxas-verify %}
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 | %ptxas-verify %}
+
+target triple = "nvptx64-unknown-cuda"
+
+;; FAST-LABEL: @t0
+;; DEFAULT-LABEL: @t0
+define <2 x float> @t0(<2 x float> %a, <2 x float> %b, <2 x float> %c) {
+; FAST-LABEL: t0(
+; FAST:       {
+; FAST-NEXT:    .reg .b64 %rd<5>;
+; FAST-EMPTY:
+; FAST-NEXT:  // %bb.0:
+; FAST-NEXT:    ld.param.b64 %rd1, [t0_param_2];
+; FAST-NEXT:    ld.param.b64 %rd2, [t0_param_1];
+; FAST-NEXT:    ld.param.b64 %rd3, [t0_param_0];
+; FAST-NEXT:    fma.rn.f32x2 %rd4, %rd3, %rd2, %rd1;
+; FAST-NEXT:    st.param.b64 [func_retval0], %rd4;
+; FAST-NEXT:    ret;
+;
+; DEFAULT-LABEL: t0(
+; DEFAULT:       {
+; DEFAULT-NEXT:    .reg .b64 %rd<6>;
+; DEFAULT-EMPTY:
+; DEFAULT-NEXT:  // %bb.0:
+; DEFAULT-NEXT:    ld.param.b64 %rd1, [t0_param_2];
+; DEFAULT-NEXT:    ld.param.b64 %rd2, [t0_param_1];
+; DEFAULT-NEXT:    ld.param.b64 %rd3, [t0_param_0];
+; DEFAULT-NEXT:    mul.rn.f32x2 %rd4, %rd3, %rd2;
+; DEFAULT-NEXT:    add.rn.f32x2 %rd5, %rd4, %rd1;
+; DEFAULT-NEXT:    st.param.b64 [func_retval0], %rd5;
+; DEFAULT-NEXT:    ret;
+  %v0 = fmul <2 x float> %a, %b
+  %v1 = fadd <2 x float> %v0, %c
+  ret <2 x float> %v1
+}
+
+;; We cannot form an fma here, but make sure we explicitly emit add.rn.f32x2
+;; to prevent ptxas from fusing this with anything else.
+define <2 x float> @t1(<2 x float> %a, <2 x float> %b) {
+; FAST-LABEL: t1(
+; FAST:       {
+; FAST-NEXT:    .reg .b64 %rd<6>;
+; FAST-EMPTY:
+; FAST-NEXT:  // %bb.0:
+; FAST-NEXT:    ld.param.b64 %rd1, [t1_param_1];
+; FAST-NEXT:    ld.param.b64 %rd2, [t1_param_0];
+; FAST-NEXT:    add.f32x2 %rd3, %rd2, %rd1;
+; FAST-NEXT:    sub.f32x2 %rd4, %rd2, %rd1;
+; FAST-NEXT:    mul.f32x2 %rd5, %rd3, %rd4;
+; FAST-NEXT:    st.param.b64 [func_retval0], %rd5;
+; FAST-NEXT:    ret;
+;
+; DEFAULT-LABEL: t1(
+; DEFAULT:       {
+; DEFAULT-NEXT:    .reg .b64 %rd<6>;
+; DEFAULT-EMPTY:
+; DEFAULT-NEXT:  // %bb.0:
+; DEFAULT-NEXT:    ld.param.b64 %rd1, [t1_param_1];
+; DEFAULT-NEXT:    ld.param.b64 %rd2, [t1_param_0];
+; DEFAULT-NEXT:    add.rn.f32x2 %rd3, %rd2, %rd1;
+; DEFAULT-NEXT:    sub.rn.f32x2 %rd4, %rd2, %rd1;
+; DEFAULT-NEXT:    mul.rn.f32x2 %rd5, %rd3, %rd4;
+; DEFAULT-NEXT:    st.param.b64 [func_retval0], %rd5;
+; DEFAULT-NEXT:    ret;
+  %v1 = fadd <2 x float> %a, %b
+  %v2 = fsub <2 x float> %a, %b
+  %v3 = fmul <2 x float> %v1, %v2
+  ret <2 x float> %v3
+}
+
+;; Make sure we generate the non ".rn" version when the "contract" flag is
+;; present on the instructions
+define <2 x float> @t2(<2 x float> %a, <2 x float> %b) {
+; CHECK-LABEL: t2(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b64 %rd<6>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b64 %rd1, [t2_param_1];
+; CHECK-NEXT:    ld.param.b64 %rd2, [t2_param_0];
+; CHECK-NEXT:    add.f32x2 %rd3, %rd2, %rd1;
+; CHECK-NEXT:    sub.f32x2 %rd4, %rd2, %rd1;
+; CHECK-NEXT:    mul.f32x2 %rd5, %rd3, %rd4;
+; CHECK-NEXT:    st.param.b64 [func_retval0], %rd5;
+; CHECK-NEXT:    ret;
+  %v1 = fadd contract <2 x float> %a, %b
+  %v2 = fsub contract <2 x float> %a, %b
+  %v3 = fmul contract <2 x float> %v1, %v2
+  ret <2 x float> %v3
+}
+
+;; Make sure we always fold to fma when the "contract" flag is present
+define <2 x float> @t3(<2 x float> %a, <2 x float> %b, <2 x float> %c) {
+; CHECK-LABEL: t3(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b64 %rd1, [t3_param_2];
+; CHECK-NEXT:    ld.param.b64 %rd2, [t3_param_1];
+; CHECK-NEXT:    ld.param.b64 %rd3, [t3_param_0];
+; CHECK-NEXT:    fma.rn.f32x2 %rd4, %rd3, %rd2, %rd1;
+; CHECK-NEXT:    st.param.b64 [func_retval0], %rd4;
+; CHECK-NEXT:    ret;
+  %v0 = fmul contract <2 x float> %a, %b
+  %v1 = fadd contract <2 x float> %v0, %c
+  ret <2 x float> %v1
+}