llvm/test/CodeGen/NVPTX/bf16-sub.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | FileCheck %s
; RUN: %if ptxas-sm_90 && ptxas-isa-7.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | %ptxas-verify -arch=sm_90 %}

define bfloat @sub_rn_bf16(bfloat %a, bfloat %b) {
; CHECK-LABEL: sub_rn_bf16(
; CHECK:       {
; CHECK-NEXT:    .reg .b16 %rs<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b16 %rs1, [sub_rn_bf16_param_0];
; CHECK-NEXT:    ld.param.b16 %rs2, [sub_rn_bf16_param_1];
; CHECK-NEXT:    sub.rn.bf16 %rs3, %rs1, %rs2;
; CHECK-NEXT:    st.param.b16 [func_retval0], %rs3;
; CHECK-NEXT:    ret;
  %1 = fneg bfloat %b
  %res = call bfloat @llvm.nvvm.fadd.bf16(bfloat %a, bfloat %1, i32 1)
  ret bfloat %res
}

define <2 x bfloat> @sub_rn_bf16x2(<2 x bfloat> %a, <2 x bfloat> %b) {
; CHECK-LABEL: sub_rn_bf16x2(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [sub_rn_bf16x2_param_0];
; CHECK-NEXT:    ld.param.b32 %r2, [sub_rn_bf16x2_param_1];
; CHECK-NEXT:    sub.rn.bf16x2 %r3, %r1, %r2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
  %1 = fneg <2 x bfloat> %b
  %res = call <2 x bfloat> @llvm.nvvm.fadd.v2bf16(<2 x bfloat> %a, <2 x bfloat> %1, i32 1)
  ret <2 x bfloat> %res
}