[NVPTX] Add more FMA intriniscs/builtins

[jchlanda]

Summary

This patch adds builtins/intrinsics for the following variants of FMA:

• f16
  • rn
  • rn_ftz
  • rn_sat
  • rn_ftz_sat
  • rn_relu
  • rn_ftz_relu
• f16x2
  • rn
  • rn_ftz
  • rn_sat
  • rn_ftz_sat
  • rn_relu
  • rn_ftz_relu
• bf16
  • rn
  • rn_relu
• bf16x2
  • rn
  • rn_relu

They all require PTX 7.0, SM_80.

ptxas (Cuda compilation tools, release 11.0, V11.0.194) is happy with the generated assembly.

Depends on D<117887>

[tra]

They all require PTX 7.0, SM_80.

According to https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-fma only fma.relu and bf16* variants require ptx70/sm80:

PTX ISA Notes
Introduced in PTX ISA version 4.2.

fma.relu.{f16, f16x2} and fma{.relu}.{bf16, bf16x2} introduced in PTX ISA version 7.0.

Target ISA Notes
Requires sm_53 or higher.

fma.relu.{f16, f16x2} and fma{.relu}.{bf16, bf16x2} require sm_80 or higher.

[jchlanda]

Set correct SM and PTX version.

[jchlanda]

In D118977#3297465, @tra wrote:

They all require PTX 7.0, SM_80.

According to https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-fma only fma.relu and bf16* variants require ptx70/sm80:

PTX ISA Notes
Introduced in PTX ISA version 4.2.

fma.relu.{f16, f16x2} and fma{.relu}.{bf16, bf16x2} introduced in PTX ISA version 7.0.

Target ISA Notes
Requires sm_53 or higher.

fma.relu.{f16, f16x2} and fma{.relu}.{bf16, bf16x2} require sm_80 or higher.

My bad, sorry. Fixed now.

[tra]

In D118977#3299974, @jchlanda wrote:

Target ISA Notes
Requires sm_53 or higher.

I think we do need this constraint applied to the new builtins, too. Right now nothing stops using them on a GPU where they do not exist and that will likely crash the compiler when we fail to find a matching intrinsic.

[jchlanda]

Add sm/ptx version guard to f16{x2} builtins.

[jchlanda]

In D118977#3302158, @tra wrote:

In D118977#3299974, @jchlanda wrote:

Target ISA Notes
Requires sm_53 or higher.

I think we do need this constraint applied to the new builtins, too. Right now nothing stops using them on a GPU where they do not exist and that will likely crash the compiler when we fail to find a matching intrinsic.

You are right, it would lead to unexpected crashes, I wrongly assumed that since the lowest version guarded against is PTX 6.0/SM_60, anything lower than that should be OK. Fixed now.

[tra]

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

937

I think the default should be the most useful/common and the least surprising value.
I'd argue that in this case it would be []. This would give reader a reasonable idea about what's going on even without looking at FMA_TUPLE implementation.

[jchlanda]

Tidy up FMA_TUPLE class.

[jchlanda]

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

937

Agreed, I wrote that class before folding in f32 and f64, hasPTX70, hasSM80 made more sense then. Changed now.

[jchlanda]

PTX/sm version tidy up.

[jchlanda]

@tra I've fixed the test failure (math-intrins.ll) the rest seems to be unrelated timeouts, would you be able to merge those patches in, as I don't have the commit access please? The same goes for https://reviews.llvm.org/D117887 and https://reviews.llvm.org/D119157 from Nicolas. Thanks.

[tra]

In D118977#3329146, @jchlanda wrote:

@tra I've fixed the test failure (math-intrins.ll) the rest seems to be unrelated timeouts,

Thank you.

would you be able to merge those patches in, as I don't have the commit access please? The same goes for https://reviews.llvm.org/D117887 and https://reviews.llvm.org/D119157 from Nicolas. Thanks.

I'll be out of office till Tuesday. I will land the patches when I'm back.

clang/include/clang/Basic/BuiltinsNVPTX.def

 //===----------------------------------------------------------------------===//
 
 // The format of this database matches clang/Basic/Builtins.def.
 
 #if defined(BUILTIN) && !defined(TARGET_BUILTIN)
 #   define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS)
 #endif
 
+#pragma push_macro("SM_53")
 #pragma push_macro("SM_70")
 #pragma push_macro("SM_72")
 #pragma push_macro("SM_75")
 #pragma push_macro("SM_80")
 #pragma push_macro("SM_86")
 #define SM_86 "sm_86"
 #define SM_80 "sm_80|" SM_86
 #define SM_75 "sm_75|" SM_80
 #define SM_72 "sm_72|" SM_75
 #define SM_70 "sm_70|" SM_72
 
 #pragma push_macro("SM_60")
 #define SM_60 "sm_60|sm_61|sm_62|" SM_70
+#define SM_53 "sm_53|" SM_60
 
+#pragma push_macro("PTX42")
 #pragma push_macro("PTX60")
 #pragma push_macro("PTX61")
 #pragma push_macro("PTX63")
 #pragma push_macro("PTX64")
 #pragma push_macro("PTX65")
 #pragma push_macro("PTX70")
 #pragma push_macro("PTX71")
 #pragma push_macro("PTX72")
 #pragma push_macro("PTX73")
 #pragma push_macro("PTX74")
 #pragma push_macro("PTX75")
 #define PTX75 "ptx75"
 #define PTX74 "ptx74|" PTX75
 #define PTX73 "ptx73|" PTX74
 #define PTX72 "ptx72|" PTX73
 #define PTX71 "ptx71|" PTX72
 #define PTX70 "ptx70|" PTX71
 #define PTX65 "ptx65|" PTX70
 #define PTX64 "ptx64|" PTX65
 #define PTX63 "ptx63|" PTX64
 #define PTX61 "ptx61|" PTX63
 #define PTX60 "ptx60|" PTX61
+#define PTX42 "ptx42|" PTX60
 
 #pragma push_macro("AND")
 #define AND(a, b) "(" a "),(" b ")"
 
 // Special Registers
 
 BUILTIN(__nvvm_read_ptx_sreg_tid_x, "i", "nc")
 BUILTIN(__nvvm_read_ptx_sreg_tid_y, "i", "nc")

 BUILTIN(__nvvm_sin_approx_ftz_f, "ff", "")
 BUILTIN(__nvvm_sin_approx_f, "ff", "")
 
 BUILTIN(__nvvm_cos_approx_ftz_f, "ff", "")
 BUILTIN(__nvvm_cos_approx_f, "ff", "")
 
 // Fma
 
+TARGET_BUILTIN(__nvvm_fma_rn_f16, "hhhh", "", AND(SM_53, PTX42))
+TARGET_BUILTIN(__nvvm_fma_rn_ftz_f16, "hhhh", "", AND(SM_53, PTX42))
+TARGET_BUILTIN(__nvvm_fma_rn_sat_f16, "hhhh", "", AND(SM_53, PTX42))
+TARGET_BUILTIN(__nvvm_fma_rn_ftz_sat_f16, "hhhh", "", AND(SM_53, PTX42))
+TARGET_BUILTIN(__nvvm_fma_rn_relu_f16, "hhhh", "", AND(SM_80, PTX70))
+TARGET_BUILTIN(__nvvm_fma_rn_ftz_relu_f16, "hhhh", "", AND(SM_80, PTX70))
+TARGET_BUILTIN(__nvvm_fma_rn_f16x2, "V2hV2hV2hV2h", "", AND(SM_53, PTX42))
+TARGET_BUILTIN(__nvvm_fma_rn_ftz_f16x2, "V2hV2hV2hV2h", "", AND(SM_53, PTX42))
+TARGET_BUILTIN(__nvvm_fma_rn_sat_f16x2, "V2hV2hV2hV2h", "", AND(SM_53, PTX42))
+TARGET_BUILTIN(__nvvm_fma_rn_ftz_sat_f16x2, "V2hV2hV2hV2h", "", AND(SM_53, PTX42))
+TARGET_BUILTIN(__nvvm_fma_rn_relu_f16x2, "V2hV2hV2hV2h", "", AND(SM_80, PTX70))
+TARGET_BUILTIN(__nvvm_fma_rn_ftz_relu_f16x2, "V2hV2hV2hV2h", "", AND(SM_80, PTX70))
+TARGET_BUILTIN(__nvvm_fma_rn_bf16, "UsUsUsUs", "", AND(SM_80, PTX70))
+TARGET_BUILTIN(__nvvm_fma_rn_relu_bf16, "UsUsUsUs", "", AND(SM_80, PTX70))
+TARGET_BUILTIN(__nvvm_fma_rn_bf16x2, "ZUiZUiZUiZUi", "", AND(SM_80, PTX70))
+TARGET_BUILTIN(__nvvm_fma_rn_relu_bf16x2, "ZUiZUiZUiZUi", "", AND(SM_80, PTX70))
 BUILTIN(__nvvm_fma_rn_ftz_f, "ffff", "")
 BUILTIN(__nvvm_fma_rn_f, "ffff", "")
 BUILTIN(__nvvm_fma_rz_ftz_f, "ffff", "")
 BUILTIN(__nvvm_fma_rz_f, "ffff", "")
 BUILTIN(__nvvm_fma_rm_ftz_f, "ffff", "")
 BUILTIN(__nvvm_fma_rm_f, "ffff", "")
 BUILTIN(__nvvm_fma_rp_ftz_f, "ffff", "")
 BUILTIN(__nvvm_fma_rp_f, "ffff", "")

 TARGET_BUILTIN(__nvvm_abs_bf16, "UsUs", "", AND(SM_80,PTX70))
 TARGET_BUILTIN(__nvvm_abs_bf16x2, "ZUiZUi", "", AND(SM_80,PTX70))
 TARGET_BUILTIN(__nvvm_neg_bf16, "UsUs", "", AND(SM_80,PTX70))
 TARGET_BUILTIN(__nvvm_neg_bf16x2, "ZUiZUi", "", AND(SM_80,PTX70))
 
 #undef BUILTIN
 #undef TARGET_BUILTIN
 #pragma pop_macro("AND")
+#pragma pop_macro("SM_53")
 #pragma pop_macro("SM_60")
 #pragma pop_macro("SM_70")
 #pragma pop_macro("SM_72")
 #pragma pop_macro("SM_75")
 #pragma pop_macro("SM_80")
 #pragma pop_macro("SM_86")
+#pragma pop_macro("PTX42")
 #pragma pop_macro("PTX60")
 #pragma pop_macro("PTX61")
 #pragma pop_macro("PTX63")
 #pragma pop_macro("PTX64")
 #pragma pop_macro("PTX65")
 #pragma pop_macro("PTX70")
 #pragma pop_macro("PTX71")
 #pragma pop_macro("PTX72")
 #pragma pop_macro("PTX73")
 #pragma pop_macro("PTX74")
 #pragma pop_macro("PTX75")

clang/test/CodeGen/builtins-nvptx-native-half-type.c

 // RUN:   -emit-llvm -o - -x cuda %s \
 // RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX72_SM86 %s
 
 // RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown \
 // RUN:   -target-cpu sm_86 -target-feature +ptx72 -fcuda-is-device \
 // RUN:   -fnative-half-type -S -emit-llvm -o - -x cuda %s \
 // RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX72_SM86 %s
 
+// RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu \
+// RUN:   sm_53 -target-feature +ptx42 -fcuda-is-device -fnative-half-type -S \
+// RUN:   -emit-llvm -o - -x cuda %s \
+// RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX42_SM53 %s
+
+// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown \
+// RUN:   -target-cpu sm_53 -target-feature +ptx42 -fcuda-is-device \
+// RUN:   -fnative-half-type -S -emit-llvm -o - -x cuda %s \
+// RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX42_SM53 %s
+
 #define __device__ __attribute__((device))
 
 // CHECK-LABEL: nvvm_min_max_sm80
 __device__ void nvvm_min_max_sm80() {
 #if __CUDA_ARCH__ >= 800
   // CHECK_PTX70_SM80: call half @llvm.nvvm.fmin.f16
   __nvvm_fmin_f16(0.1f16, 0.1f16);
   // CHECK_PTX70_SM80: call half @llvm.nvvm.fmin.ftz.f16
#if __CUDA_ARCH__ >= 800
   // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmax.nan.f16x2
   __nvvm_fmax_nan_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
   // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmax.ftz.nan.f16x2
   __nvvm_fmax_ftz_nan_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
 #endif
   // CHECK: ret void
 }
 
+// CHECK-LABEL: nvvm_fma_f16_f16x2_sm80
+__device__ void nvvm_fma_f16_f16x2_sm80() {
+#if __CUDA_ARCH__ >= 800
+  // CHECK_PTX70_SM80: call half @llvm.nvvm.fma.rn.relu.f16
+  __nvvm_fma_rn_relu_f16(0.1f16, 0.1f16, 0.1f16);
+  // CHECK_PTX70_SM80: call half @llvm.nvvm.fma.rn.ftz.relu.f16
+  __nvvm_fma_rn_ftz_relu_f16(0.1f16, 0.1f16, 0.1f16);
+
+  // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fma.rn.relu.f16x2
+  __nvvm_fma_rn_relu_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},
+                           {0.1f16, 0.7f16});
+  // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fma.rn.ftz.relu.f16x2
+  __nvvm_fma_rn_ftz_relu_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},
+                               {0.1f16, 0.7f16});
+#endif
+  // CHECK: ret void
+}
+
+// CHECK-LABEL: nvvm_fma_f16_f16x2_sm53
+__device__ void nvvm_fma_f16_f16x2_sm53() {
+#if __CUDA_ARCH__ >= 530
+  // CHECK_PTX42_SM53: call half @llvm.nvvm.fma.rn.f16
+  __nvvm_fma_rn_f16(0.1f16, 0.1f16, 0.1f16);
+  // CHECK_PTX42_SM53: call half @llvm.nvvm.fma.rn.ftz.f16
+  __nvvm_fma_rn_ftz_f16(0.1f16, 0.1f16, 0.1f16);
+  // CHECK_PTX42_SM53: call half @llvm.nvvm.fma.rn.sat.f16
+  __nvvm_fma_rn_sat_f16(0.1f16, 0.1f16, 0.1f16);
+  // CHECK_PTX42_SM53: call half @llvm.nvvm.fma.rn.ftz.sat.f16
+  __nvvm_fma_rn_ftz_sat_f16(0.1f16, 0.1f16, 0.1f16);
+
+  // CHECK_PTX42_SM53: call <2 x half> @llvm.nvvm.fma.rn.f16x2
+  __nvvm_fma_rn_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},
+                      {0.1f16, 0.7f16});
+  // CHECK_PTX42_SM53: call <2 x half> @llvm.nvvm.fma.rn.ftz.f16x2
+  __nvvm_fma_rn_ftz_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},
+                          {0.1f16, 0.7f16});
+  // CHECK_PTX42_SM53: call <2 x half> @llvm.nvvm.fma.rn.sat.f16x2
+  __nvvm_fma_rn_sat_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},
+                          {0.1f16, 0.7f16});
+  // CHECK_PTX42_SM53: call <2 x half> @llvm.nvvm.fma.rn.ftz.sat.f16x2
+  __nvvm_fma_rn_ftz_sat_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},
+                              {0.1f16, 0.7f16});
+#endif
+  // CHECK: ret void
+}
+
 // CHECK-LABEL: nvvm_min_max_sm86
 __device__ void nvvm_min_max_sm86() {
 #if __CUDA_ARCH__ >= 860
   // CHECK_PTX72_SM86: call half @llvm.nvvm.fmin.xorsign.abs.f16
   __nvvm_fmin_xorsign_abs_f16(0.1f16, 0.1f16);
   // CHECK_PTX72_SM86: call half @llvm.nvvm.fmin.ftz.xorsign.abs.f16
   __nvvm_fmin_ftz_xorsign_abs_f16(0.1f16, 0.1f16);
   // CHECK_PTX72_SM86: call half @llvm.nvvm.fmin.nan.xorsign.abs.f16

clang/test/CodeGen/builtins-nvptx.c

#if __CUDA_ARCH__ >= 800
   // CHECK_PTX70_SM80: call float @llvm.nvvm.fmax.nan.f
   __nvvm_fmax_nan_f(0.1f, (float)NAN32);
   // CHECK_PTX70_SM80: call float @llvm.nvvm.fmax.ftz.nan.f
   __nvvm_fmax_ftz_nan_f(0.1f, (float)NAN32);
 
 #endif
   // CHECK: ret void
 }
+
+// CHECK-LABEL: nvvm_fma_bf16_bf16x2_sm80
+__device__ void nvvm_fma_bf16_bf16x2_sm80() {
+#if __CUDA_ARCH__ >= 800
+  // CHECK_PTX70_SM80: call i16 @llvm.nvvm.fma.rn.bf16
+  __nvvm_fma_rn_bf16(0x1234, 0x7FBF, 0x1234);
+  // CHECK_PTX70_SM80: call i16 @llvm.nvvm.fma.rn.relu.bf16
+  __nvvm_fma_rn_relu_bf16(0x1234, 0x7FBF, 0x1234);
+  // CHECK_PTX70_SM80: call i32 @llvm.nvvm.fma.rn.bf16x2
+  __nvvm_fma_rn_bf16x2(0x7FBFFFFF, 0xFFFFFFFF, 0x7FBFFFFF);
+  // CHECK_PTX70_SM80: call i32 @llvm.nvvm.fma.rn.relu.bf16x2
+  __nvvm_fma_rn_relu_bf16x2(0x7FBFFFFF, 0xFFFFFFFF, 0x7FBFFFFF);
+#endif
+  // CHECK: ret void
+}
+
 // CHECK-LABEL: nvvm_min_max_sm86
 __device__ void nvvm_min_max_sm86() {
 #if __CUDA_ARCH__ >= 860
 
   // CHECK_PTX72_SM86: call i16 @llvm.nvvm.fmin.xorsign.abs.bf16
   __nvvm_fmin_xorsign_abs_bf16(BF16, BF16_2);
   // CHECK_PTX72_SM86: call i16 @llvm.nvvm.fmin.nan.xorsign.abs.bf16
   __nvvm_fmin_nan_xorsign_abs_bf16(BF16, NANBF16);

llvm/include/llvm/IR/IntrinsicsNVVM.td

def int_nvvm_cos_approx_ftz_f : GCCBuiltin<"__nvvm_cos_approx_ftz_f">,
       DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
   def int_nvvm_cos_approx_f : GCCBuiltin<"__nvvm_cos_approx_f">,
       DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
 
 //
 // Fma
 //
 
-  def int_nvvm_fma_rn_ftz_f : GCCBuiltin<"__nvvm_fma_rn_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
+  foreach variant = ["_rn_f16", "_rn_ftz_f16", "_rn_sat_f16",
+    "_rn_ftz_sat_f16", "_rn_relu_f16", "_rn_ftz_relu_f16"] in {
+    def int_nvvm_fma # variant : GCCBuiltin<!strconcat("__nvvm_fma", variant)>,
+        DefaultAttrsIntrinsic<[llvm_half_ty],
+          [llvm_half_ty, llvm_half_ty, llvm_half_ty],
           [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_fma_rn_f : GCCBuiltin<"__nvvm_fma_rn_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_fma_rz_ftz_f : GCCBuiltin<"__nvvm_fma_rz_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_fma_rz_f : GCCBuiltin<"__nvvm_fma_rz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_fma_rm_ftz_f : GCCBuiltin<"__nvvm_fma_rm_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_fma_rm_f : GCCBuiltin<"__nvvm_fma_rm_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_fma_rp_ftz_f : GCCBuiltin<"__nvvm_fma_rp_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
-        [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_fma_rp_f : GCCBuiltin<"__nvvm_fma_rp_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
+  }
+
+  foreach variant = ["_rn_f16x2", "_rn_ftz_f16x2", "_rn_sat_f16x2",
+    "_rn_ftz_sat_f16x2", "_rn_relu_f16x2", "_rn_ftz_relu_f16x2"] in {
+    def int_nvvm_fma # variant : GCCBuiltin<!strconcat("__nvvm_fma", variant)>,
+      DefaultAttrsIntrinsic<[llvm_v2f16_ty],
+        [llvm_v2f16_ty, llvm_v2f16_ty, llvm_v2f16_ty],
         [IntrNoMem, IntrSpeculatable]>;
+  }
 
-  def int_nvvm_fma_rn_d : GCCBuiltin<"__nvvm_fma_rn_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty],
-        [llvm_double_ty, llvm_double_ty, llvm_double_ty],
+  foreach variant = ["_rn_bf16", "_rn_relu_bf16"] in {
+    def int_nvvm_fma # variant : GCCBuiltin<!strconcat("__nvvm_fma", variant)>,
+      DefaultAttrsIntrinsic<[llvm_i16_ty],
+        [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
         [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_fma_rz_d : GCCBuiltin<"__nvvm_fma_rz_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty],
-        [llvm_double_ty, llvm_double_ty, llvm_double_ty],
+  }
+
+  foreach variant = ["_rn_bf16x2", "_rn_relu_bf16x2"] in {
+    def int_nvvm_fma # variant : GCCBuiltin<!strconcat("__nvvm_fma", variant)>,
+      DefaultAttrsIntrinsic<[llvm_i32_ty],
+        [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
         [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_fma_rm_d : GCCBuiltin<"__nvvm_fma_rm_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty],
-        [llvm_double_ty, llvm_double_ty, llvm_double_ty],
+  }
+
+  foreach variant = ["_rn_ftz_f", "_rn_f", "_rz_ftz_f", "_rz_f", "_rm_ftz_f",
+    "_rm_f", "_rp_ftz_f", "_rp_f"] in {
+    def int_nvvm_fma # variant : GCCBuiltin<!strconcat("__nvvm_fma", variant)>,
+      DefaultAttrsIntrinsic<[llvm_float_ty],
+        [llvm_float_ty, llvm_float_ty, llvm_float_ty],
         [IntrNoMem, IntrSpeculatable]>;
-  def int_nvvm_fma_rp_d : GCCBuiltin<"__nvvm_fma_rp_d">,
+  }
+
+  foreach variant = ["_rn_d", "_rz_d", "_rm_d", "_rp_d"] in {
+    def int_nvvm_fma # variant : GCCBuiltin<!strconcat("__nvvm_fma", variant)>,
       DefaultAttrsIntrinsic<[llvm_double_ty],
         [llvm_double_ty, llvm_double_ty, llvm_double_ty],
         [IntrNoMem, IntrSpeculatable]>;
+  }
 
 //
 // Rcp
 //
 
   def int_nvvm_rcp_rn_ftz_f : GCCBuiltin<"__nvvm_rcp_rn_ftz_f">,
       DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
   def int_nvvm_rcp_rn_f : GCCBuiltin<"__nvvm_rcp_rn_f">,

llvm/lib/Target/NVPTX/NVPTXInstrInfo.td

 def do_SQRTF32_RN : Predicate<"usePrecSqrtF32()">;
 
 def hasHWROT32 : Predicate<"Subtarget->hasHWROT32()">;
 def noHWROT32 : Predicate<"!Subtarget->hasHWROT32()">;
 
 def True : Predicate<"true">;
 
 def hasPTX31 : Predicate<"Subtarget->getPTXVersion() >= 31">;
+def hasPTX42 : Predicate<"Subtarget->getPTXVersion() >= 42">;
 def hasPTX60 : Predicate<"Subtarget->getPTXVersion() >= 60">;
 def hasPTX61 : Predicate<"Subtarget->getPTXVersion() >= 61">;
 def hasPTX63 : Predicate<"Subtarget->getPTXVersion() >= 63">;
 def hasPTX64 : Predicate<"Subtarget->getPTXVersion() >= 64">;
 def hasPTX65 : Predicate<"Subtarget->getPTXVersion() >= 65">;
 def hasPTX70 : Predicate<"Subtarget->getPTXVersion() >= 70">;
 def hasPTX71 : Predicate<"Subtarget->getPTXVersion() >= 71">;
 def hasPTX72 : Predicate<"Subtarget->getPTXVersion() >= 72">;
 
 def hasSM30 : Predicate<"Subtarget->getSmVersion() >= 30">;
+def hasSM53 : Predicate<"Subtarget->getSmVersion() >= 53">;
 def hasSM70 : Predicate<"Subtarget->getSmVersion() >= 70">;
 def hasSM72 : Predicate<"Subtarget->getSmVersion() >= 72">;
 def hasSM75 : Predicate<"Subtarget->getSmVersion() >= 75">;
 def hasSM80 : Predicate<"Subtarget->getSmVersion() >= 80">;
 def hasSM86 : Predicate<"Subtarget->getSmVersion() >= 86">;
 
 // non-sync shfl instructions are not available on sm_70+ in PTX6.4+
 def hasSHFL : Predicate<"!(Subtarget->getSmVersion() >= 70" 

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

list<Predicate> Preds = []>
             : NVPTXInst<(outs t_regclass:$dst),
               (ins s0_regclass:$src0, s1_regclass:$src1),
             OpcStr,
         [(set t_regclass:$dst, (IntOP s0_regclass:$src0, s1_regclass:$src1))]>,
         Requires<Preds>;
 
 class F_MATH_3<string OpcStr, NVPTXRegClass t_regclass,
   NVPTXRegClass s0_regclass, NVPTXRegClass s1_regclass,
-  NVPTXRegClass s2_regclass, Intrinsic IntOP>
+  NVPTXRegClass s2_regclass, Intrinsic IntOP, list<Predicate> Preds = []>
             : NVPTXInst<(outs t_regclass:$dst),
               (ins s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2),
             OpcStr,
         [(set t_regclass:$dst,
-          (IntOP s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2))]>;
+          (IntOP s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2))]>,
+          Requires<Preds>;
 
 //
 // MISC
 //
 
 def INT_NVVM_PRMT : F_MATH_3<"prmt.b32 \t$dst, $src0, $src1, $src2;", Int32Regs,
   Int32Regs, Int32Regs, Int32Regs, int_nvvm_prmt>;
 

 def INT_NVVM_FMIN_D : F_MATH_2<"min.f64 \t$dst, $src0, $src1;", Float64Regs,
   Float64Regs, Float64Regs, int_nvvm_fmin_d>;
 def INT_NVVM_FMAX_D : F_MATH_2<"max.f64 \t$dst, $src0, $src1;", Float64Regs,
   Float64Regs, Float64Regs, int_nvvm_fmax_d>;
 
 //
 // Min Max f16, f16x2, bf16, bf16x2
 //
+
 class MIN_MAX_TUPLE<string V, Intrinsic I, NVPTXRegClass RC,
                     list<Predicate> Preds = [hasPTX70, hasSM80]> {
   string Variant = V;
   Intrinsic Intr = I;
   NVPTXRegClass RegClass = RC;
   list<Predicate> Predicates = Preds;
 }
 
def INT_NVVM_COS_APPROX_FTZ_F : F_MATH_1<"cos.approx.ftz.f32 \t$dst, $src0;",
   Float32Regs, Float32Regs, int_nvvm_cos_approx_ftz_f>;
 def INT_NVVM_COS_APPROX_F : F_MATH_1<"cos.approx.f32 \t$dst, $src0;",
   Float32Regs, Float32Regs, int_nvvm_cos_approx_f>;
 
 //
 // Fma
 //
 
-def INT_NVVM_FMA_RN_FTZ_F
-  : F_MATH_3<"fma.rn.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
-    Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rn_ftz_f>;
-def INT_NVVM_FMA_RN_F : F_MATH_3<"fma.rn.f32 \t$dst, $src0, $src1, $src2;",
-  Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rn_f>;
-def INT_NVVM_FMA_RZ_FTZ_F
-  : F_MATH_3<"fma.rz.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
-    Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rz_ftz_f>;
-def INT_NVVM_FMA_RZ_F : F_MATH_3<"fma.rz.f32 \t$dst, $src0, $src1, $src2;",
-  Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rz_f>;
-def INT_NVVM_FMA_RM_FTZ_F
-  : F_MATH_3<"fma.rm.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
-    Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rm_ftz_f>;
-def INT_NVVM_FMA_RM_F : F_MATH_3<"fma.rm.f32 \t$dst, $src0, $src1, $src2;",
-  Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rm_f>;
-def INT_NVVM_FMA_RP_FTZ_F
-  : F_MATH_3<"fma.rp.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
-    Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rp_ftz_f>;
-def INT_NVVM_FMA_RP_F : F_MATH_3<"fma.rp.f32 \t$dst, $src0, $src1, $src2;",
-  Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rp_f>;
-
-def INT_NVVM_FMA_RN_D : F_MATH_3<"fma.rn.f64 \t$dst, $src0, $src1, $src2;",
-  Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rn_d>;
-def INT_NVVM_FMA_RZ_D : F_MATH_3<"fma.rz.f64 \t$dst, $src0, $src1, $src2;",
-  Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rz_d>;
-def INT_NVVM_FMA_RM_D : F_MATH_3<"fma.rm.f64 \t$dst, $src0, $src1, $src2;",
-  Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rm_d>;
-def INT_NVVM_FMA_RP_D : F_MATH_3<"fma.rp.f64 \t$dst, $src0, $src1, $src2;",
-  Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rp_d>;
+class FMA_TUPLE<string V, Intrinsic I, NVPTXRegClass RC,
+                list<Predicate> Preds = [hasPTX70, hasSM80]> {
+  string Variant = V;
+  Intrinsic Intr = I;
+  NVPTXRegClass RegClass = RC;
+  list<Predicate> Predicates = Preds;
+}
+
+multiclass FMA_INST {
+  foreach P = [
+    FMA_TUPLE<"_rn_f64", int_nvvm_fma_rn_d, Float64Regs, []>,
+    FMA_TUPLE<"_rz_f64", int_nvvm_fma_rz_d, Float64Regs, []>,
+    FMA_TUPLE<"_rm_f64", int_nvvm_fma_rm_d, Float64Regs, []>,
+    FMA_TUPLE<"_rp_f64", int_nvvm_fma_rp_d, Float64Regs, []>,
+
+    FMA_TUPLE<"_rn_ftz_f32", int_nvvm_fma_rn_ftz_f, Float32Regs, []>,
+    FMA_TUPLE<"_rn_f32", int_nvvm_fma_rn_f, Float32Regs, []>,
+    FMA_TUPLE<"_rz_ftz_f32", int_nvvm_fma_rz_ftz_f, Float32Regs, []>,
+    FMA_TUPLE<"_rz_f32", int_nvvm_fma_rz_f, Float32Regs, []>,
+    FMA_TUPLE<"_rm_f32", int_nvvm_fma_rm_f, Float32Regs, []>,
+    FMA_TUPLE<"_rm_ftz_f32", int_nvvm_fma_rm_ftz_f, Float32Regs, []>,
+    FMA_TUPLE<"_rp_f32", int_nvvm_fma_rp_f, Float32Regs, []>,
+    FMA_TUPLE<"_rp_ftz_f32", int_nvvm_fma_rp_ftz_f, Float32Regs, []>,
+
+    FMA_TUPLE<"_rn_f16", int_nvvm_fma_rn_f16, Float16Regs, [hasPTX42, hasSM53]>,
+    FMA_TUPLE<"_rn_ftz_f16", int_nvvm_fma_rn_ftz_f16, Float16Regs,
+      [hasPTX42, hasSM53]>,
+    FMA_TUPLE<"_rn_sat_f16", int_nvvm_fma_rn_sat_f16, Float16Regs,
+      [hasPTX42, hasSM53]>,
+    FMA_TUPLE<"_rn_ftz_sat_f16", int_nvvm_fma_rn_ftz_sat_f16, Float16Regs,
+      [hasPTX42, hasSM53]>,
+    FMA_TUPLE<"_rn_relu_f16", int_nvvm_fma_rn_relu_f16, Float16Regs>,
+    FMA_TUPLE<"_rn_ftz_relu_f16", int_nvvm_fma_rn_ftz_relu_f16, Float16Regs>,
+
+    FMA_TUPLE<"_rn_f16x2", int_nvvm_fma_rn_f16x2, Float16x2Regs,
+      [hasPTX42, hasSM53]>,
+    FMA_TUPLE<"_rn_ftz_f16x2", int_nvvm_fma_rn_ftz_f16x2, Float16x2Regs,
+      [hasPTX42, hasSM53]>,
+    FMA_TUPLE<"_rn_sat_f16x2", int_nvvm_fma_rn_sat_f16x2, Float16x2Regs,
+      [hasPTX42, hasSM53]>,
+    FMA_TUPLE<"_rn_ftz_sat_f16x2", int_nvvm_fma_rn_ftz_sat_f16x2,
+      Float16x2Regs>,
+    FMA_TUPLE<"_rn_relu_f16x2", int_nvvm_fma_rn_relu_f16x2, Float16x2Regs>,
+    FMA_TUPLE<"_rn_ftz_relu_f16x2", int_nvvm_fma_rn_ftz_relu_f16x2,
+      Float16x2Regs>,
+
+    FMA_TUPLE<"_rn_bf16", int_nvvm_fma_rn_bf16, Int16Regs>,
+    FMA_TUPLE<"_rn_relu_bf16", int_nvvm_fma_rn_relu_bf16, Int16Regs>,
+
+    FMA_TUPLE<"_rn_bf16x2", int_nvvm_fma_rn_bf16x2, Int32Regs>,
+    FMA_TUPLE<"_rn_relu_bf16x2", int_nvvm_fma_rn_relu_bf16x2, Int32Regs>
+  ] in {
+    def P.Variant :
+      F_MATH_3<!strconcat("fma",
+        !subst("_", ".", P.Variant), " \t$dst, $src0, $src1, $src2;"),
+        P.RegClass, P.RegClass, P.RegClass, P.RegClass, P.Intr, P.Predicates>;
+  }
+}
+
+defm INT_NVVM_FMA : FMA_INST;
 
 //
 // Rcp
 //
 
 def INT_NVVM_RCP_RN_FTZ_F : F_MATH_1<"rcp.rn.ftz.f32 \t$dst, $src0;",
   Float32Regs, Float32Regs, int_nvvm_rcp_rn_ftz_f>;
 def INT_NVVM_RCP_RN_F : F_MATH_1<"rcp.rn.f32 \t$dst, $src0;",

llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp

const SimplifyAction Action = [II]() -> SimplifyAction {
     case Intrinsic::nvvm_floor_ftz_f:
       return {Intrinsic::floor, FTZ_MustBeOn};
     case Intrinsic::nvvm_fma_rn_d:
       return {Intrinsic::fma, FTZ_Any};
     case Intrinsic::nvvm_fma_rn_f:
       return {Intrinsic::fma, FTZ_MustBeOff};
     case Intrinsic::nvvm_fma_rn_ftz_f:
       return {Intrinsic::fma, FTZ_MustBeOn};
+    case Intrinsic::nvvm_fma_rn_f16:
+      return {Intrinsic::fma, FTZ_MustBeOff, true};
+    case Intrinsic::nvvm_fma_rn_ftz_f16:
+      return {Intrinsic::fma, FTZ_MustBeOn, true};
+    case Intrinsic::nvvm_fma_rn_f16x2:
+      return {Intrinsic::fma, FTZ_MustBeOff, true};
+    case Intrinsic::nvvm_fma_rn_ftz_f16x2:
+      return {Intrinsic::fma, FTZ_MustBeOn, true};
     case Intrinsic::nvvm_fmax_d:
       return {Intrinsic::maxnum, FTZ_Any};
     case Intrinsic::nvvm_fmax_f:
       return {Intrinsic::maxnum, FTZ_MustBeOff};
     case Intrinsic::nvvm_fmax_ftz_f:
       return {Intrinsic::maxnum, FTZ_MustBeOn};
     case Intrinsic::nvvm_fmax_nan_f:
       return {Intrinsic::maximum, FTZ_MustBeOff};

llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70-instcombine.ll

 declare <2 x half> @llvm.nvvm.fmax.ftz.f16x2(<2 x half>, <2 x half>)
 declare float @llvm.nvvm.fmax.nan.f(float, float)
 declare float @llvm.nvvm.fmax.ftz.nan.f(float, float)
 declare half @llvm.nvvm.fmax.nan.f16(half, half)
 declare half @llvm.nvvm.fmax.ftz.nan.f16(half, half)
 declare <2 x half> @llvm.nvvm.fmax.nan.f16x2(<2 x half>, <2 x half>)
 declare <2 x half> @llvm.nvvm.fmax.ftz.nan.f16x2(<2 x half>, <2 x half>)
 
+; f16 and f16x2 fma are available since ptx 4.2 and sm_53.
+declare half @llvm.nvvm.fma.rn.f16(half, half, half)
+declare half @llvm.nvvm.fma.rn.ftz.f16(half, half, half)
+declare <2 x half> @llvm.nvvm.fma.rn.f16x2(<2 x half>, <2 x half>, <2 x half>)
+declare <2 x half> @llvm.nvvm.fma.rn.ftz.f16x2(<2 x half>, <2 x half>, <2 x half>)
+
 ; CHECK-LABEL: fmin_f16
 define half @fmin_f16(half %0, half %1) {
   ; CHECK-NOT: @llvm.nvvm.fmin.f16
   ; CHECK: @llvm.minnum.f16
   %res = call half @llvm.nvvm.fmin.f16(half %0, half %1)
   ret half %res
 }
 

 ; CHECK-LABEL: fmax_ftz_nan_f16x2_no_attr
 define <2 x half> @fmax_ftz_nan_f16x2_no_attr(<2 x half> %0, <2 x half> %1) {
   ; CHECK-NOT: @llvm.maximum.v2f16
   ; CHECK: @llvm.nvvm.fmax.ftz.nan.f16x2
   %res = call <2 x half> @llvm.nvvm.fmax.ftz.nan.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
+; CHECK-LABEL: fma_rn_f16
+define half @fma_rn_f16(half %0, half %1, half %2) {
+  ; CHECK-NOT: @llvm.nvvm.fma.rn.f16
+  ; CHECK: @llvm.fma.f16
+  %res = call half @llvm.nvvm.fma.rn.f16(half %0, half %1, half %2)
+  ret half %res
+}
+
+; CHECK-LABEL: fma_rn_ftz_f16_no_attr
+define half @fma_rn_ftz_f16_no_attr(half %0, half %1, half %2) {
+  ; CHECK-NOT: @llvm.fma.f16
+  ; CHECK: @llvm.nvvm.fma.rn.ftz.f16
+  %res = call half @llvm.nvvm.fma.rn.ftz.f16(half %0, half %1, half %2)
+  ret half %res
+}
+
+; CHECK-LABEL: fma_rn_ftz_f16
+define half @fma_rn_ftz_f16(half %0, half %1, half %2) #0 {
+  ; CHECK-NOT: @llvm.nvvm.fma.rn.ftz.f16
+  ; CHECK: @llvm.fma.f16
+  %res = call half @llvm.nvvm.fma.rn.ftz.f16(half %0, half %1, half %2)
+  ret half %res
+}
+
+; CHECK-LABEL: fma_rn_f16x2
+define <2 x half> @fma_rn_f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2) {
+  ; CHECK-NOT: @llvm.nvvm.fma.rn.f16x2
+  ; CHECK: @llvm.fma.v2f16
+  %res = call <2 x half> @llvm.nvvm.fma.rn.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2)
+  ret <2 x half> %res
+}
+
+; CHECK-LABEL: fma_rn_ftz_f16x2
+define <2 x half> @fma_rn_ftz_f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2) #0 {
+  ; CHECK-NOT: @llvm.nvvm.fma.rn.ftz.f16x2
+  ; CHECK: @llvm.fma.v2f16
+  %res = call <2 x half> @llvm.nvvm.fma.rn.ftz.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2)
+  ret <2 x half> %res
+}
+
+; CHECK-LABEL: fma_rn_ftz_f16x2_no_attr
+define <2 x half> @fma_rn_ftz_f16x2_no_attr(<2 x half> %0, <2 x half> %1, <2 x half> %2) {
+  ; CHECK-NOT: @llvm.fma.v2f16
+  ; CHECK: @llvm.nvvm.fma.rn.ftz.f16x2
+  %res = call <2 x half> @llvm.nvvm.fma.rn.ftz.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2)
+  ret <2 x half> %res
+}
+
 attributes #0 = { "denormal-fp-math"="preserve-sign" }
 attributes #1 = { "denormal-fp-math-f32"="preserve-sign" }

llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70.ll

 declare <2 x half> @llvm.nvvm.fmax.ftz.f16x2(<2 x half>, <2 x half>)
 declare <2 x half> @llvm.nvvm.fmax.nan.f16x2(<2 x half>, <2 x half>)
 declare <2 x half> @llvm.nvvm.fmax.ftz.nan.f16x2(<2 x half>, <2 x half>)
 declare i16 @llvm.nvvm.fmax.bf16(i16, i16)
 declare i16 @llvm.nvvm.fmax.nan.bf16(i16, i16)
 declare i32 @llvm.nvvm.fmax.bf16x2(i32, i32)
 declare i32 @llvm.nvvm.fmax.nan.bf16x2(i32, i32)
 
+declare half @llvm.nvvm.fma.rn.relu.f16(half, half, half)
+declare half @llvm.nvvm.fma.rn.ftz.relu.f16(half, half, half)
+declare <2 x half> @llvm.nvvm.fma.rn.relu.f16x2(<2 x half>, <2 x half>, <2 x half>)
+declare <2 x half> @llvm.nvvm.fma.rn.ftz.relu.f16x2(<2 x half>, <2 x half>, <2 x half>)
+declare i16 @llvm.nvvm.fma.rn.bf16(i16, i16, i16)
+declare i16 @llvm.nvvm.fma.rn.relu.bf16(i16, i16, i16)
+declare i32 @llvm.nvvm.fma.rn.bf16x2(i32, i32, i32)
+declare i32 @llvm.nvvm.fma.rn.relu.bf16x2(i32, i32, i32)
+
 ; CHECK-LABEL: abs_bf16
 define i16 @abs_bf16(i16 %0) {
+  ; CHECK-NOT: call
   ; CHECK: abs.bf16
   %res = call i16 @llvm.nvvm.abs.bf16(i16 %0);
   ret i16 %res
 }
 
 ; CHECK-LABEL: abs_bf16x2
 define i32 @abs_bf16x2(i32 %0) {
+  ; CHECK-NOT: call
   ; CHECK: abs.bf16x2
   %res = call i32 @llvm.nvvm.abs.bf16x2(i32 %0);
   ret i32 %res
 }
 
 ; CHECK-LABEL: neg_bf16
 define i16 @neg_bf16(i16 %0) {
+  ; CHECK-NOT: call
   ; CHECK: neg.bf16
   %res = call i16 @llvm.nvvm.neg.bf16(i16 %0);
   ret i16 %res
 }
 
 ; CHECK-LABEL: neg_bf16x2
 define i32 @neg_bf16x2(i32 %0) {
+  ; CHECK-NOT: call
   ; CHECK: neg.bf16x2
   %res = call i32 @llvm.nvvm.neg.bf16x2(i32 %0);
   ret i32 %res
 }
 
 ; CHECK-LABEL: fmin_nan_f
 define float @fmin_nan_f(float %0, float %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.NaN.f32
   %res = call float @llvm.nvvm.fmin.nan.f(float %0, float %1);
   ret float %res
 }
 
 ; CHECK-LABEL: fmin_ftz_nan_f
 define float @fmin_ftz_nan_f(float %0, float %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.ftz.NaN.f32
   %res = call float @llvm.nvvm.fmin.ftz.nan.f(float %0, float %1);
   ret float %res
 }
 
 ; CHECK-LABEL: fmin_f16
 define half @fmin_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.f16
   %res = call half @llvm.nvvm.fmin.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmin_ftz_f16
 define half @fmin_ftz_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.ftz.f16
   %res = call half @llvm.nvvm.fmin.ftz.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmin_nan_f16
 define half @fmin_nan_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.NaN.f16
   %res = call half @llvm.nvvm.fmin.nan.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmin_ftz_nan_f16
 define half @fmin_ftz_nan_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.ftz.NaN.f16
   %res = call half @llvm.nvvm.fmin.ftz.nan.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmin_f16x2
 define <2 x half> @fmin_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.f16x2
   %res = call <2 x half> @llvm.nvvm.fmin.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmin_ftz_f16x2
 define <2 x half> @fmin_ftz_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.ftz.f16x2
   %res = call <2 x half> @llvm.nvvm.fmin.ftz.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmin_nan_f16x2
 define <2 x half> @fmin_nan_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.NaN.f16x2
   %res = call <2 x half> @llvm.nvvm.fmin.nan.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmin_ftz_nan_f16x2
 define <2 x half> @fmin_ftz_nan_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.ftz.NaN.f16x2
   %res = call <2 x half> @llvm.nvvm.fmin.ftz.nan.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmin_bf16
 define i16 @fmin_bf16(i16 %0, i16 %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.bf16
   %res = call i16 @llvm.nvvm.fmin.bf16(i16 %0, i16 %1)
   ret i16 %res
 }
 
 ; CHECK-LABEL: fmin_nan_bf16
 define i16 @fmin_nan_bf16(i16 %0, i16 %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.NaN.bf16
   %res = call i16 @llvm.nvvm.fmin.nan.bf16(i16 %0, i16 %1)
   ret i16 %res
 }
 
 ; CHECK-LABEL: fmin_bf16x2
 define i32 @fmin_bf16x2(i32 %0, i32 %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.bf16x2
   %res = call i32 @llvm.nvvm.fmin.bf16x2(i32 %0, i32 %1)
   ret i32 %res
 }
 
 ; CHECK-LABEL: fmin_nan_bf16x2
 define i32 @fmin_nan_bf16x2(i32 %0, i32 %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.NaN.bf16x2
   %res = call i32 @llvm.nvvm.fmin.nan.bf16x2(i32 %0, i32 %1)
   ret i32 %res
 }
 
 ; CHECK-LABEL: fmax_nan_f
 define float @fmax_nan_f(float %0, float %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.NaN.f32
   %res = call float @llvm.nvvm.fmax.nan.f(float %0, float %1);
   ret float %res
 }
 
 ; CHECK-LABEL: fmax_ftz_nan_f
 define float @fmax_ftz_nan_f(float %0, float %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.ftz.NaN.f32
   %res = call float @llvm.nvvm.fmax.ftz.nan.f(float %0, float %1);
   ret float %res
 }
 
 ; CHECK-LABEL: fmax_f16
 define half @fmax_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.f16
   %res = call half @llvm.nvvm.fmax.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmax_ftz_f16
 define half @fmax_ftz_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.ftz.f16
   %res = call half @llvm.nvvm.fmax.ftz.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmax_nan_f16
 define half @fmax_nan_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.NaN.f16
   %res = call half @llvm.nvvm.fmax.nan.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmax_ftz_nan_f16
 define half @fmax_ftz_nan_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.ftz.NaN.f16
   %res = call half @llvm.nvvm.fmax.ftz.nan.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmax_f16x2
 define <2 x half> @fmax_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.f16x2
   %res = call <2 x half> @llvm.nvvm.fmax.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmax_ftz_f16x2
 define <2 x half> @fmax_ftz_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.ftz.f16x2
   %res = call <2 x half> @llvm.nvvm.fmax.ftz.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmax_nan_f16x2
 define <2 x half> @fmax_nan_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.NaN.f16x2
   %res = call <2 x half> @llvm.nvvm.fmax.nan.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmax_ftz_nan_f16x2
 define <2 x half> @fmax_ftz_nan_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.ftz.NaN.f16x2
   %res = call <2 x half> @llvm.nvvm.fmax.ftz.nan.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmax_bf16
 define i16 @fmax_bf16(i16 %0, i16 %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.bf16
   %res = call i16 @llvm.nvvm.fmax.bf16(i16 %0, i16 %1)
   ret i16 %res
 }
 
 ; CHECK-LABEL: fmax_nan_bf16
 define i16 @fmax_nan_bf16(i16 %0, i16 %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.NaN.bf16
   %res = call i16 @llvm.nvvm.fmax.nan.bf16(i16 %0, i16 %1)
   ret i16 %res
 }
 
 ; CHECK-LABEL: fmax_bf16x2
 define i32 @fmax_bf16x2(i32 %0, i32 %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.bf16x2
   %res = call i32 @llvm.nvvm.fmax.bf16x2(i32 %0, i32 %1)
   ret i32 %res
 }
 
 ; CHECK-LABEL: fmax_nan_bf16x2
 define i32 @fmax_nan_bf16x2(i32 %0, i32 %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.NaN.bf16x2
   %res = call i32 @llvm.nvvm.fmax.nan.bf16x2(i32 %0, i32 %1)
   ret i32 %res
 }
+
+; CHECK-LABEL: fma_rn_relu_f16
+define half @fma_rn_relu_f16(half %0, half %1, half %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.relu.f16
+  %res = call half @llvm.nvvm.fma.rn.relu.f16(half %0, half %1, half %2)
+  ret half %res
+}
+
+; CHECK-LABEL: fma_rn_ftz_relu_f16
+define half @fma_rn_ftz_relu_f16(half %0, half %1, half %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.ftz.relu.f16
+  %res = call half @llvm.nvvm.fma.rn.ftz.relu.f16(half %0, half %1, half %2)
+  ret half %res
+}
+
+; CHECK-LABEL: fma_rn_relu_f16x2
+define <2 x half> @fma_rn_relu_f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.relu.f16x2
+  %res = call <2 x half> @llvm.nvvm.fma.rn.relu.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2)
+  ret <2 x half> %res
+}
+
+; CHECK-LABEL: fma_rn_ftz_relu_f16x2
+define <2 x half> @fma_rn_ftz_relu_f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.ftz.relu.f16x2
+  %res = call <2 x half> @llvm.nvvm.fma.rn.ftz.relu.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2)
+  ret <2 x half> %res
+}
+
+; CHECK-LABEL: fma_rn_bf16
+define i16 @fma_rn_bf16(i16 %0, i16 %1, i16 %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.bf16
+  %res = call i16 @llvm.nvvm.fma.rn.bf16(i16 %0, i16 %1, i16 %2)
+  ret i16 %res
+}
+
+; CHECK-LABEL: fma_rn_relu_bf16
+define i16 @fma_rn_relu_bf16(i16 %0, i16 %1, i16 %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.relu.bf16
+  %res = call i16 @llvm.nvvm.fma.rn.relu.bf16(i16 %0, i16 %1, i16 %2)
+  ret i16 %res
+}
+
+; CHECK-LABEL: fma_rn_bf16x2
+define i32 @fma_rn_bf16x2(i32 %0, i32 %1, i32 %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.bf16x2
+  %res = call i32 @llvm.nvvm.fma.rn.bf16x2(i32 %0, i32 %1, i32 %2)
+  ret i32 %res
+}
+
+; CHECK-LABEL: fma_rn_relu_bf16x2
+define i32 @fma_rn_relu_bf16x2(i32 %0, i32 %1, i32 %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.relu.bf16x2
+  %res = call i32 @llvm.nvvm.fma.rn.relu.bf16x2(i32 %0, i32 %1, i32 %2)
+  ret i32 %res
+}

llvm/test/CodeGen/NVPTX/math-intrins-sm86-ptx72.ll

 declare i32 @llvm.nvvm.fmax.nan.xorsign.abs.bf16x2(i32, i32)
 declare float @llvm.nvvm.fmax.xorsign.abs.f(float, float)
 declare float @llvm.nvvm.fmax.ftz.xorsign.abs.f(float, float)
 declare float @llvm.nvvm.fmax.nan.xorsign.abs.f(float, float)
 declare float @llvm.nvvm.fmax.ftz.nan.xorsign.abs.f(float, float)
 
 ; CHECK-LABEL: fmin_xorsign_abs_f16
 define half @fmin_xorsign_abs_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.xorsign.abs.f16
   %res = call half @llvm.nvvm.fmin.xorsign.abs.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmin_ftz_xorsign_abs_f16
 define half @fmin_ftz_xorsign_abs_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.ftz.xorsign.abs.f16
   %res = call half @llvm.nvvm.fmin.ftz.xorsign.abs.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmin_nan_xorsign_abs_f16
 define half @fmin_nan_xorsign_abs_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.NaN.xorsign.abs.f16
   %res = call half @llvm.nvvm.fmin.nan.xorsign.abs.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmin_ftz_nan_xorsign_abs_f16
 define half @fmin_ftz_nan_xorsign_abs_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.ftz.NaN.xorsign.abs.f16
   %res = call half @llvm.nvvm.fmin.ftz.nan.xorsign.abs.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmin_xorsign_abs_f16x2
 define <2 x half> @fmin_xorsign_abs_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.xorsign.abs.f16x2
   %res = call <2 x half> @llvm.nvvm.fmin.xorsign.abs.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmin_ftz_xorsign_abs_f16x2
 define <2 x half> @fmin_ftz_xorsign_abs_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.ftz.xorsign.abs.f16x2
   %res = call <2 x half> @llvm.nvvm.fmin.ftz.xorsign.abs.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmin_nan_xorsign_abs_f16x2
 define <2 x half> @fmin_nan_xorsign_abs_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.NaN.xorsign.abs.f16x2
   %res = call <2 x half> @llvm.nvvm.fmin.nan.xorsign.abs.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmin_ftz_nan_xorsign_abs_f16x2
 define <2 x half> @fmin_ftz_nan_xorsign_abs_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.ftz.NaN.xorsign.abs.f16x2
   %res = call <2 x half> @llvm.nvvm.fmin.ftz.nan.xorsign.abs.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmin_xorsign_abs_bf16
 define i16 @fmin_xorsign_abs_bf16(i16 %0, i16 %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.xorsign.abs.bf16
   %res = call i16 @llvm.nvvm.fmin.xorsign.abs.bf16(i16 %0, i16 %1)
   ret i16 %res
 }
 
 ; CHECK-LABEL: fmin_nan_xorsign_abs_bf16
 define i16 @fmin_nan_xorsign_abs_bf16(i16 %0, i16 %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.NaN.xorsign.abs.bf16
   %res = call i16 @llvm.nvvm.fmin.nan.xorsign.abs.bf16(i16 %0, i16 %1)
   ret i16 %res
 }
 
 ; CHECK-LABEL: fmin_xorsign_abs_bf16x2
 define i32 @fmin_xorsign_abs_bf16x2(i32 %0, i32 %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.xorsign.abs.bf16x2
   %res = call i32 @llvm.nvvm.fmin.xorsign.abs.bf16x2(i32 %0, i32 %1)
   ret i32 %res
 }
 
 ; CHECK-LABEL: fmin_nan_xorsign_abs_bf16x2
 define i32 @fmin_nan_xorsign_abs_bf16x2(i32 %0, i32 %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.NaN.xorsign.abs.bf16x2
   %res = call i32 @llvm.nvvm.fmin.nan.xorsign.abs.bf16x2(i32 %0, i32 %1)
   ret i32 %res
 }
 
 ; CHECK-LABEL: fmin_xorsign_abs_f
 define float @fmin_xorsign_abs_f(float %0, float %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.xorsign.abs.f
   %res = call float @llvm.nvvm.fmin.xorsign.abs.f(float %0, float %1)
   ret float %res
 }
 
 ; CHECK-LABEL: fmin_ftz_xorsign_abs_f
 define float @fmin_ftz_xorsign_abs_f(float %0, float %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.ftz.xorsign.abs.f
   %res = call float @llvm.nvvm.fmin.ftz.xorsign.abs.f(float %0, float %1)
   ret float %res
 }
 
 ; CHECK-LABEL: fmin_nan_xorsign_abs_f
 define float @fmin_nan_xorsign_abs_f(float %0, float %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.NaN.xorsign.abs.f
   %res = call float @llvm.nvvm.fmin.nan.xorsign.abs.f(float %0, float %1)
   ret float %res
 }
 
 ; CHECK-LABEL: fmin_ftz_nan_xorsign_abs_f
 define float @fmin_ftz_nan_xorsign_abs_f(float %0, float %1) {
+  ; CHECK-NOT: call
   ; CHECK: min.ftz.NaN.xorsign.abs.f
   %res = call float @llvm.nvvm.fmin.ftz.nan.xorsign.abs.f(float %0, float %1)
   ret float %res
 }
 
 ; CHECK-LABEL: fmax_xorsign_abs_f16
 define half @fmax_xorsign_abs_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.xorsign.abs.f16
   %res = call half @llvm.nvvm.fmax.xorsign.abs.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmax_ftz_xorsign_abs_f16
 define half @fmax_ftz_xorsign_abs_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.ftz.xorsign.abs.f16
   %res = call half @llvm.nvvm.fmax.ftz.xorsign.abs.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmax_nan_xorsign_abs_f16
 define half @fmax_nan_xorsign_abs_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.NaN.xorsign.abs.f16
   %res = call half @llvm.nvvm.fmax.nan.xorsign.abs.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmax_ftz_nan_xorsign_abs_f16
 define half @fmax_ftz_nan_xorsign_abs_f16(half %0, half %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.ftz.NaN.xorsign.abs.f16
   %res = call half @llvm.nvvm.fmax.ftz.nan.xorsign.abs.f16(half %0, half %1)
   ret half %res
 }
 
 ; CHECK-LABEL: fmax_xorsign_abs_f16x2
 define <2 x half> @fmax_xorsign_abs_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.xorsign.abs.f16x2
   %res = call <2 x half> @llvm.nvvm.fmax.xorsign.abs.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmax_ftz_xorsign_abs_f16x2
 define <2 x half> @fmax_ftz_xorsign_abs_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.ftz.xorsign.abs.f16x2
   %res = call <2 x half> @llvm.nvvm.fmax.ftz.xorsign.abs.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmax_nan_xorsign_abs_f16x2
 define <2 x half> @fmax_nan_xorsign_abs_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.NaN.xorsign.abs.f16x2
   %res = call <2 x half> @llvm.nvvm.fmax.nan.xorsign.abs.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmax_ftz_nan_xorsign_abs_f16x2
 define <2 x half> @fmax_ftz_nan_xorsign_abs_f16x2(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.ftz.NaN.xorsign.abs.f16x2
   %res = call <2 x half> @llvm.nvvm.fmax.ftz.nan.xorsign.abs.f16x2(<2 x half> %0, <2 x half> %1)
   ret <2 x half> %res
 }
 
 ; CHECK-LABEL: fmax_xorsign_abs_bf16
 define i16 @fmax_xorsign_abs_bf16(i16 %0, i16 %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.xorsign.abs.bf16
   %res = call i16 @llvm.nvvm.fmax.xorsign.abs.bf16(i16 %0, i16 %1)
   ret i16 %res
 }
 
 ; CHECK-LABEL: fmax_nan_xorsign_abs_bf16
 define i16 @fmax_nan_xorsign_abs_bf16(i16 %0, i16 %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.NaN.xorsign.abs.bf16
   %res = call i16 @llvm.nvvm.fmax.nan.xorsign.abs.bf16(i16 %0, i16 %1)
   ret i16 %res
 }
 
 ; CHECK-LABEL: fmax_xorsign_abs_bf16x2
 define i32 @fmax_xorsign_abs_bf16x2(i32 %0, i32 %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.xorsign.abs.bf16x2
   %res = call i32 @llvm.nvvm.fmax.xorsign.abs.bf16x2(i32 %0, i32 %1)
   ret i32 %res
 }
 
 ; CHECK-LABEL: fmax_nan_xorsign_abs_bf16x2
 define i32 @fmax_nan_xorsign_abs_bf16x2(i32 %0, i32 %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.NaN.xorsign.abs.bf16x2
   %res = call i32 @llvm.nvvm.fmax.nan.xorsign.abs.bf16x2(i32 %0, i32 %1)
   ret i32 %res
 }
 
 ; CHECK-LABEL: fmax_xorsign_abs_f
 define float @fmax_xorsign_abs_f(float %0, float %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.xorsign.abs.f
   %res = call float @llvm.nvvm.fmax.xorsign.abs.f(float %0, float %1)
   ret float %res
 }
 
 ; CHECK-LABEL: fmax_ftz_xorsign_abs_f
 define float @fmax_ftz_xorsign_abs_f(float %0, float %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.ftz.xorsign.abs.f
   %res = call float @llvm.nvvm.fmax.ftz.xorsign.abs.f(float %0, float %1)
   ret float %res
 }
 
 ; CHECK-LABEL: fmax_nan_xorsign_abs_f
 define float @fmax_nan_xorsign_abs_f(float %0, float %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.NaN.xorsign.abs.f
   %res = call float @llvm.nvvm.fmax.nan.xorsign.abs.f(float %0, float %1)
   ret float %res
 }
 
 ; CHECK-LABEL: fmax_ftz_nan_xorsign_abs_f
 define float @fmax_ftz_nan_xorsign_abs_f(float %0, float %1) {
+  ; CHECK-NOT: call
   ; CHECK: max.ftz.NaN.xorsign.abs.f
   %res = call float @llvm.nvvm.fmax.ftz.nan.xorsign.abs.f(float %0, float %1)
   ret float %res
 }

llvm/test/CodeGen/NVPTX/math-intrins.ll

 declare <2 x half> @llvm.minnum.v2f16(<2 x half>, <2 x half>) #0
 declare half @llvm.maxnum.f16(half, half) #0
 declare float @llvm.maxnum.f32(float, float) #0
 declare double @llvm.maxnum.f64(double, double) #0
 declare <2 x half> @llvm.maxnum.v2f16(<2 x half>, <2 x half>) #0
 declare float @llvm.fma.f32(float, float, float) #0
 declare double @llvm.fma.f64(double, double, double) #0
 
+declare half @llvm.nvvm.fma.rn.f16(half, half, half)
+declare half @llvm.nvvm.fma.rn.ftz.f16(half, half, half)
+declare half @llvm.nvvm.fma.rn.sat.f16(half, half, half)
+declare half @llvm.nvvm.fma.rn.ftz.sat.f16(half, half, half)
+declare <2 x half> @llvm.nvvm.fma.rn.f16x2(<2 x half>, <2 x half>, <2 x half>)
+declare <2 x half> @llvm.nvvm.fma.rn.ftz.f16x2(<2 x half>, <2 x half>, <2 x half>)
+declare <2 x half> @llvm.nvvm.fma.rn.sat.f16x2(<2 x half>, <2 x half>, <2 x half>)
+declare <2 x half> @llvm.nvvm.fma.rn.ftz.sat.f16x2(<2 x half>, <2 x half>, <2 x half>)
+
 ; ---- ceil ----
 
 ; CHECK-LABEL: ceil_float
 define float @ceil_float(float %a) {
   ; CHECK: cvt.rpi.f32.f32
   %b = call float @llvm.ceil.f32(float %a)
   ret float %b
 }

 
 ; CHECK-LABEL: @fma_double
 define double @fma_double(double %a, double %b, double %c) {
   ; CHECK: fma.rn.f64
   %x = call double @llvm.fma.f64(double %a, double %b, double %c)
   ret double %x
 }
 
+; CHECK-LABEL: fma_rn_f16
+define half @fma_rn_f16(half %0, half %1, half %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.f16
+  %res = call half @llvm.nvvm.fma.rn.f16(half %0, half %1, half %2)
+  ret half %res
+}
+
+; CHECK-LABEL: fma_rn_ftz_f16
+define half @fma_rn_ftz_f16(half %0, half %1, half %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.ftz.f16
+  %res = call half @llvm.nvvm.fma.rn.ftz.f16(half %0, half %1, half %2)
+  ret half %res
+}
+
+; CHECK-LABEL: fma_rn_sat_f16
+define half @fma_rn_sat_f16(half %0, half %1, half %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.sat.f16
+  %res = call half @llvm.nvvm.fma.rn.sat.f16(half %0, half %1, half %2)
+  ret half %res
+}
+
+; CHECK-LABEL: fma_rn_ftz_sat_f16
+define half @fma_rn_ftz_sat_f16(half %0, half %1, half %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.ftz.sat.f16
+  %res = call half @llvm.nvvm.fma.rn.ftz.sat.f16(half %0, half %1, half %2)
+  ret half %res
+}
+
+; CHECK-LABEL: fma_rn_f16x2
+define <2 x half> @fma_rn_f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.f16x2
+  %res = call <2 x half> @llvm.nvvm.fma.rn.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2)
+  ret <2 x half> %res
+}
+
+; CHECK-LABEL: fma_rn_ftz_f16x2
+define <2 x half> @fma_rn_ftz_f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.ftz.f16x2
+  %res = call <2 x half> @llvm.nvvm.fma.rn.ftz.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2)
+  ret <2 x half> %res
+}
+
+; CHECK-LABEL: fma_rn_sat_f16x2
+define <2 x half> @fma_rn_sat_f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.sat.f16x2
+  %res = call <2 x half> @llvm.nvvm.fma.rn.sat.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2)
+  ret <2 x half> %res
+}
+
+; CHECK-LABEL: fma_rn_ftz_sat_f16x2
+define <2 x half> @fma_rn_ftz_sat_f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2) {
+  ; CHECK-NOT: call
+  ; CHECK: fma.rn.ftz.sat.f16x2
+  %res = call <2 x half> @llvm.nvvm.fma.rn.ftz.sat.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2)
+  ret <2 x half> %res
+}
+
 attributes #0 = { nounwind readnone }
 attributes #1 = { "denormal-fp-math-f32" = "preserve-sign" }