[NVPTX] abs, neg, min, max intrinsics for half types

[jchlanda]

Summary

This patch extends abs and neg intrinsics to support bf16 and bf16x2 data formats.
Extends min and max to support .NaN modifier and f16, f16x2, bf16 and bf16x2 data formats.

[jchlanda]

I need to update this to the upstream, as it was based on intel's fork. Abandoning for now.

clang/include/clang/Basic/BuiltinsNVPTX.def

 BUILTIN(__builtin_ptx_implicit_offset, "Ui*", "nc")
 
 // MISC
 
 BUILTIN(__nvvm_prmt, "UiUiUiUi", "")
 
 // Min Max
 
-BUILTIN(__nvvm_fmax_ftz_f, "fff",  "")
-BUILTIN(__nvvm_fmax_f, "fff",  "")
-BUILTIN(__nvvm_fmin_ftz_f, "fff",  "")
 BUILTIN(__nvvm_fmin_f, "fff",  "")
+BUILTIN(__nvvm_fmin_ftz_f, "fff",  "")
+TARGET_BUILTIN(__nvvm_fmin_nan_f, "fff",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmin_ftz_nan_f, "fff",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmin_f16, "hhh",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmin_ftz_f16, "hhh",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmin_nan_f16, "hhh",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmin_ftz_nan_f16, "hhh",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmin_f16x2, "V2hV2hV2h",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmin_ftz_f16x2, "V2hV2hV2h",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmin_nan_f16x2, "V2hV2hV2h",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmin_ftz_nan_f16x2, "V2hV2hV2h",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmin_bf16, "UsUsUs", "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmin_nan_bf16, "UsUsUs", "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmin_bf16x2, "ZUiZUiZUi", "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmin_nan_bf16x2, "ZUiZUiZUi", "", AND(SM_80,PTX70))
+
+BUILTIN(__nvvm_fmax_f, "fff",  "")
+BUILTIN(__nvvm_fmax_ftz_f, "fff",  "")
+TARGET_BUILTIN(__nvvm_fmax_nan_f, "fff",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmax_ftz_nan_f, "fff",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmax_f16, "hhh",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmax_ftz_f16, "hhh",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmax_nan_f16, "hhh",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmax_ftz_nan_f16, "hhh",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmax_f16x2, "V2hV2hV2h",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmax_ftz_f16x2, "V2hV2hV2h",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmax_nan_f16x2, "V2hV2hV2h",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmax_ftz_nan_f16x2, "V2hV2hV2h",  "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmax_bf16, "UsUsUs", "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmax_nan_bf16, "UsUsUs", "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmax_bf16x2, "ZUiZUiZUi", "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_fmax_nan_bf16x2, "ZUiZUiZUi", "", AND(SM_80,PTX70))
 
-BUILTIN(__nvvm_fmax_d, "ddd", "")
 BUILTIN(__nvvm_fmin_d, "ddd", "")
+BUILTIN(__nvvm_fmax_d, "ddd", "")
 
 // Multiplication
 
 BUILTIN(__nvvm_mulhi_i, "iii", "")
 BUILTIN(__nvvm_mulhi_ui, "UiUiUi", "")
 BUILTIN(__nvvm_mulhi_ll, "LLiLLiLLi", "")
 BUILTIN(__nvvm_mulhi_ull, "ULLiULLiULLi", "")
 

 TARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_8, "vv*3vC*1", "", AND(SM_80,PTX70))
 TARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_16, "vv*3vC*1", "", AND(SM_80,PTX70))
 TARGET_BUILTIN(__nvvm_cp_async_cg_shared_global_16, "vv*3vC*1", "", AND(SM_80,PTX70))
 
 TARGET_BUILTIN(__nvvm_cp_async_commit_group, "v", "", AND(SM_80,PTX70))
 TARGET_BUILTIN(__nvvm_cp_async_wait_group, "vIi", "", AND(SM_80,PTX70))
 TARGET_BUILTIN(__nvvm_cp_async_wait_all, "v", "", AND(SM_80,PTX70))
 
+// bf16, bf16x2 abs, neg
+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_60")
 #pragma pop_macro("SM_70")
 #pragma pop_macro("SM_72")
 #pragma pop_macro("SM_75")
 #pragma pop_macro("SM_80")

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

+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu \
+// RUN:   sm_80 -target-feature +ptx70 -fcuda-is-device -fnative-half-type -S \
+// RUN:   -emit-llvm -o - -x cuda %s \
+// RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX70_SM80 %s
+
+// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown \
+// RUN:   -target-cpu sm_80 -target-feature +ptx70 -fcuda-is-device \
+// RUN:   -fnative-half-type -S -emit-llvm -o - -x cuda %s \
+// RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX70_SM80 %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
+  __nvvm_fmin_ftz_f16(0.1f16, 0.1f16);
+  // CHECK_PTX70_SM80: call half @llvm.nvvm.fmin.nan.f16
+  __nvvm_fmin_nan_f16(0.1f16, 0.1f16);
+  // CHECK_PTX70_SM80: call half @llvm.nvvm.fmin.ftz.nan.f16
+  __nvvm_fmin_ftz_nan_f16(0.1f16, 0.1f16);
+  // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmin.f16x2
+  __nvvm_fmin_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
+  // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmin.ftz.f16x2
+  __nvvm_fmin_ftz_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
+  // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmin.nan.f16x2
+  __nvvm_fmin_nan_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
+  // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmin.ftz.nan.f16x2
+  __nvvm_fmin_ftz_nan_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
+
+  // CHECK_PTX70_SM80: call half @llvm.nvvm.fmax.f16
+  __nvvm_fmax_f16(0.1f16, 0.1f16);
+  // CHECK_PTX70_SM80: call half @llvm.nvvm.fmax.ftz.f16
+  __nvvm_fmax_ftz_f16(0.1f16, 0.1f16);
+  // CHECK_PTX70_SM80: call half @llvm.nvvm.fmax.nan.f16
+  __nvvm_fmax_nan_f16(0.1f16, 0.1f16);
+  // CHECK_PTX70_SM80: call half @llvm.nvvm.fmax.ftz.nan.f16
+  __nvvm_fmax_ftz_nan_f16(0.1f16, 0.1f16);
+  // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmax.f16x2
+  __nvvm_fmax_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
+  // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmax.ftz.f16x2
+  __nvvm_fmax_ftz_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
+  // 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
+}

clang/test/CodeGen/builtins-nvptx.c

__nvvm_cp_async_wait_group(0);
     // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.wait.group(i32 8)
   __nvvm_cp_async_wait_group(8);
     // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.wait.group(i32 16)
   __nvvm_cp_async_wait_group(16);
   // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.wait.all
   __nvvm_cp_async_wait_all();
   #endif
   // CHECK: ret void
-}
+}
+
+// CHECK-LABEL: nvvm_abs_neg_bf16_bf16x2_sm80
+__device__ void nvvm_abs_neg_bf16_bf16x2_sm80() {
+#if __CUDA_ARCH__ >= 800
+
+  // CHECK_PTX70_SM80: call i16 @llvm.nvvm.abs.bf16(i16 -1)
+  __nvvm_abs_bf16(0xFFFF);
+  // CHECK_PTX70_SM80: call i32 @llvm.nvvm.abs.bf16x2(i32 -1)
+  __nvvm_abs_bf16x2(0xFFFFFFFF);
+
+  // CHECK_PTX70_SM80: call i16 @llvm.nvvm.neg.bf16(i16 -1)
+  __nvvm_neg_bf16(0xFFFF);
+  // CHECK_PTX70_SM80: call i32 @llvm.nvvm.neg.bf16x2(i32 -1)
+  __nvvm_neg_bf16x2(0xFFFFFFFF);
+#endif
+  // CHECK: ret void
+}
+
+// CHECK-LABEL: nvvm_min_max_sm80
+__device__ void nvvm_min_max_sm80() {
+#if __CUDA_ARCH__ >= 800
+
+  // CHECK_PTX70_SM80: call float @llvm.nvvm.fmin.nan.f
+  __nvvm_fmin_nan_f(0.1f, (float)0x7FBFFFFF);
+  // CHECK_PTX70_SM80: call float @llvm.nvvm.fmin.ftz.nan.f
+  __nvvm_fmin_ftz_nan_f(0.1f, (float)0x7FBFFFFF);
+
+  // CHECK_PTX70_SM80: call i16 @llvm.nvvm.fmin.bf16
+  __nvvm_fmin_bf16(0x1234, 0x7FBF);
+  // CHECK_PTX70_SM80: call i16 @llvm.nvvm.fmin.nan.bf16
+  __nvvm_fmin_nan_bf16(0x1234, 0x7FBF);
+  // CHECK_PTX70_SM80: call i32 @llvm.nvvm.fmin.bf16x2
+  __nvvm_fmin_bf16x2(0x7FBFFFFF, 0xFFFFFFFF);
+  // CHECK_PTX70_SM80: call i32 @llvm.nvvm.fmin.nan.bf16x2
+  __nvvm_fmin_nan_bf16x2(0x7FBFFFFF, 0xFFFFFFFF);
+  // CHECK_PTX70_SM80: call float @llvm.nvvm.fmax.nan.f
+  __nvvm_fmax_nan_f(0.1f, (float)0x7FBFFFFF);
+  // CHECK_PTX70_SM80: call float @llvm.nvvm.fmax.ftz.nan.f
+  __nvvm_fmax_ftz_nan_f(0.1f, (float)0x7FBFFFFF);
+
+  // CHECK_PTX70_SM80: call float @llvm.nvvm.fmax.nan.f
+  __nvvm_fmax_nan_f(0.1f, (float)0x7FBFFFFF);
+  // CHECK_PTX70_SM80: call float @llvm.nvvm.fmax.ftz.nan.f
+  __nvvm_fmax_ftz_nan_f(0.1f, (float)0x7FBFFFFF);
+  // CHECK_PTX70_SM80: call i16 @llvm.nvvm.fmax.bf16
+  __nvvm_fmax_bf16(0x1234, 0x7FBF);
+  // CHECK_PTX70_SM80: call i16 @llvm.nvvm.fmax.nan.bf16
+  __nvvm_fmax_nan_bf16(0x1234, 0x7FBF);
+  // CHECK_PTX70_SM80: call i32 @llvm.nvvm.fmax.bf16x2
+  __nvvm_fmax_bf16x2(0x7FBFFFFF, 0xFFFFFFFF);
+  // CHECK_PTX70_SM80: call i32 @llvm.nvvm.fmax.nan.bf16x2
+  __nvvm_fmax_nan_bf16x2(0x7FBFFFFF, 0xFFFFFFFF);
+  // CHECK_PTX70_SM80: call float @llvm.nvvm.fmax.nan.f
+  __nvvm_fmax_nan_f(0.1f, (float)0x7FBFFFFF);
+  // CHECK_PTX70_SM80: call float @llvm.nvvm.fmax.ftz.nan.f
+  __nvvm_fmax_ftz_nan_f(0.1f, (float)0x7FBFFFFF);
+#endif
+  // CHECK: ret void
+}

llvm/include/llvm/IR/IntrinsicsNVVM.td

let TargetPrefix = "nvvm" in {
   def int_nvvm_prmt : GCCBuiltin<"__nvvm_prmt">,
       DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
         [IntrNoMem, Commutative]>;
 
 //
 // Min Max
 //
 
-  def int_nvvm_fmin_f : GCCBuiltin<"__nvvm_fmin_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
+  foreach operation = ["min", "max"] in {
+    def int_nvvm_f # operation # _d :
+      GCCBuiltin<!strconcat("__nvvm_f", operation, "_d")>,
+      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
         [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_fmin_ftz_f : GCCBuiltin<"__nvvm_fmin_ftz_f">,
+
+    foreach capability = ["_f", "_ftz_f", "_nan_f", "_ftz_nan_f"] in {
+      def int_nvvm_f # operation # capability :
+        GCCBuiltin<!strconcat("__nvvm_f", operation, capability)>,
         DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
           [IntrNoMem, IntrSpeculatable, Commutative]>;
+    }
 
-  def int_nvvm_fmax_f : GCCBuiltin<"__nvvm_fmax_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty]
-        , [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_fmax_ftz_f : GCCBuiltin<"__nvvm_fmax_ftz_f">,
-      DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
+    foreach capability = ["_f16", "_ftz_f16", "_nan_f16", "_ftz_nan_f16"] in {
+      def int_nvvm_f # operation # capability :
+        GCCBuiltin<!strconcat("__nvvm_f", operation, capability)>,
+        DefaultAttrsIntrinsic<[llvm_half_ty], [llvm_half_ty, llvm_half_ty],
           [IntrNoMem, IntrSpeculatable, Commutative]>;
+    }
 
-  def int_nvvm_fmin_d : GCCBuiltin<"__nvvm_fmin_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
+    foreach capability = ["_f16x2", "_ftz_f16x2", "_nan_f16x2",
+      "_ftz_nan_f16x2"] in {
+      def int_nvvm_f # operation # capability :
+        GCCBuiltin<!strconcat("__nvvm_f", operation, capability)>,
+        DefaultAttrsIntrinsic<[llvm_v2f16_ty], [llvm_v2f16_ty, llvm_v2f16_ty],
           [IntrNoMem, IntrSpeculatable, Commutative]>;
-  def int_nvvm_fmax_d : GCCBuiltin<"__nvvm_fmax_d">,
-      DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
+    }
+
+    foreach capability = ["_bf16", "_nan_bf16"] in {
+      def int_nvvm_f # operation # capability :
+        GCCBuiltin<!strconcat("__nvvm_f", operation, capability)>,
+        DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty],
           [IntrNoMem, IntrSpeculatable, Commutative]>;
+    }
+
+    foreach capability = ["_bf16x2", "_nan_bf16x2"] in {
+      def int_nvvm_f # operation # capability :
+        GCCBuiltin<!strconcat("__nvvm_f", operation, capability)>,
+        DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
+          [IntrNoMem, IntrSpeculatable, Commutative]>;
+    }
+  }
 
 //
 // Multiplication
 //
 
   def int_nvvm_mulhi_i : GCCBuiltin<"__nvvm_mulhi_i">,
       DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
         [IntrNoMem, IntrSpeculatable, Commutative]>;
def int_nvvm_ull2d_rp : GCCBuiltin<"__nvvm_ull2d_rp">,
       DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem, IntrSpeculatable]>;
 
   def int_nvvm_f2h_rn_ftz : GCCBuiltin<"__nvvm_f2h_rn_ftz">,
       DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
   def int_nvvm_f2h_rn : GCCBuiltin<"__nvvm_f2h_rn">,
       DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
 
 //
+// bf16, bf16x2 abs, neg
+//
+  foreach unary = ["abs", "neg"] in {
+    def int_nvvm_ # unary # _bf16 :
+      GCCBuiltin<!strconcat("__nvvm_", unary, "_bf16")>,
+      DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty], [IntrNoMem]>;
+    def int_nvvm_ # unary # _bf16x2 :
+      GCCBuiltin<!strconcat("__nvvm_", unary, "_bf16x2")>,
+      DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem]>;
+  }
+
+//
 // Bitcast
 //
 
   def int_nvvm_bitcast_f2i : GCCBuiltin<"__nvvm_bitcast_f2i">,
       DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>;
   def int_nvvm_bitcast_i2f : GCCBuiltin<"__nvvm_bitcast_i2f">,
       DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem, IntrSpeculatable]>;
 

llvm/lib/Target/NVPTX/NVPTXInstrInfo.td

NVPTXInst<(outs Float32Regs:$dst),
                [(set Float32Regs:$dst, (OpNode Float32Regs:$a, Float32Regs:$b))]>;
    def f32ri :
      NVPTXInst<(outs Float32Regs:$dst),
                (ins Float32Regs:$a, f32imm:$b),
                !strconcat(OpcStr, ".f32 \t$dst, $a, $b;"),
                [(set Float32Regs:$dst, (OpNode Float32Regs:$a, fpimm:$b))]>;
 }
 
+multiclass F3_F16_F16X2<string OpcStr, SDNode OpNode> {
+   def f16rr :
+     NVPTXInst<(outs Float16Regs:$dst),
+               (ins Float16Regs:$a, Float16Regs:$b),
+               !strconcat(OpcStr, ".f16 \t$dst, $a, $b;"),
+               [(set Float16Regs:$dst,
+                  (OpNode Float16Regs:$a, Float16Regs:$b))]>,
+               Requires<[hasPTX70, hasSM80, useFP16Math]>;
+   def f16rr_ftz :
+     NVPTXInst<(outs Float16Regs:$dst),
+               (ins Float16Regs:$a, Float16Regs:$b),
+               !strconcat(OpcStr, ".ftz.f16 \t$dst, $a, $b;"),
+               [(set Float16Regs:$dst,
+                  (OpNode Float16Regs:$a, Float16Regs:$b))]>,
+               Requires<[hasPTX70, hasSM80, useFP16Math, doF32FTZ]>;
+
+   def f16x2rr :
+     NVPTXInst<(outs Float16x2Regs:$dst),
+               (ins Float16x2Regs:$a, Float16x2Regs:$b),
+               !strconcat(OpcStr, ".f16x2 \t$dst, $a, $b;"),
+               [(set Float16x2Regs:$dst,
+                  (OpNode Float16x2Regs:$a, Float16x2Regs:$b))]>,
+               Requires<[hasPTX70, hasSM80, useFP16Math]>;
+   def f16x2rr_ftz :
+     NVPTXInst<(outs Float16x2Regs:$dst),
+               (ins Float16x2Regs:$a, Float16x2Regs:$b),
+               !strconcat(OpcStr, ".ftz.f16x2 \t$dst, $a, $b;"),
+               [(set Float16x2Regs:$dst,
+                  (OpNode Float16x2Regs:$a, Float16x2Regs:$b))]>,
+               Requires<[hasPTX70, hasSM80, useFP16Math, doF32FTZ]>;
+}
+
 // Template for instructions which take three FP args.  The
 // instructions are named "<OpcStr>.f<Width>" (e.g. "add.f64").
 //
 // Also defines ftz (flush subnormal inputs and results to sign-preserving
 // zero) variants for fp32/fp16 functions.
 //
 // This multiclass should be used for nodes that can be folded to make fma ops.
 // In this case, we use the ".rn" variant when FMA is disabled, as this behaves
NVPTXInst<(outs Float16Regs:$dst), (ins f16imm:$a),
             "mov.b16 \t$dst, $a;", []>;
 
 defm FADD : F3_fma_component<"add", fadd>;
 defm FSUB : F3_fma_component<"sub", fsub>;
 defm FMUL : F3_fma_component<"mul", fmul>;
 
 defm FMIN : F3<"min", fminnum>;
 defm FMAX : F3<"max", fmaxnum>;
-
+defm FMIN_F16_F16X2 : F3_F16_F16X2<"min", fminnum>;
+defm FMAX_F16_F16X2 : F3_F16_F16X2<"max", fmaxnum>;
 defm FABS  : F2<"abs", fabs>;
 defm FNEG  : F2<"neg", fneg>;
 defm FSQRT : F2<"sqrt.rn", fsqrt>;
 
 //
 // F64 division
 //
 def FDIV641r :

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

 def : Pat<(int_nvvm_fmin_d
             (int_nvvm_fmax_d Float64Regs:$a, immDouble0), immDouble1),
           (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;
 
 
 // We need a full string for OpcStr here because we need to deal with case like
 // INT_PTX_RECIP.
 class F_MATH_1<string OpcStr, NVPTXRegClass target_regclass,
-  NVPTXRegClass src_regclass, Intrinsic IntOP>
+  NVPTXRegClass src_regclass, Intrinsic IntOP, list<Predicate> Preds = []>
             : NVPTXInst<(outs target_regclass:$dst), (ins src_regclass:$src0),
             OpcStr,
-        [(set target_regclass:$dst, (IntOP src_regclass:$src0))]>;
+        [(set target_regclass:$dst, (IntOP src_regclass:$src0))]>,
+        Requires<Preds>;
 
 // We need a full string for OpcStr here because we need to deal with the case
 // like INT_PTX_NATIVE_POWR_F.
 class F_MATH_2<string OpcStr, NVPTXRegClass t_regclass,
-  NVPTXRegClass s0_regclass, NVPTXRegClass s1_regclass, Intrinsic IntOP>
+  NVPTXRegClass s0_regclass, NVPTXRegClass s1_regclass, Intrinsic IntOP,
+  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))]>;
+        [(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>
             : NVPTXInst<(outs t_regclass:$dst),
               (ins s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2),
             OpcStr,
         [(set t_regclass:$dst,

 // Min Max
 //
 
 def INT_NVVM_FMIN_F : F_MATH_2<"min.f32 \t$dst, $src0, $src1;", Float32Regs,
   Float32Regs, Float32Regs, int_nvvm_fmin_f>;
 def INT_NVVM_FMIN_FTZ_F : F_MATH_2<"min.ftz.f32 \t$dst, $src0, $src1;",
   Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmin_ftz_f>;
 
+def INT_NVVM_FMIN_NAN_F : F_MATH_2<"min.NaN.f32 \t$dst, $src0, $src1;",
+  Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmin_nan_f,
+  [hasPTX70, hasSM80]>;
+def INT_NVVM_FMIN_FTZ_NAN_F : F_MATH_2<"min.ftz.NaN.f32 \t$dst, $src0, $src1;",
+  Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmin_ftz_nan_f,
+  [hasPTX70, hasSM80]>;
+
 def INT_NVVM_FMAX_F : F_MATH_2<"max.f32 \t$dst, $src0, $src1;", Float32Regs,
   Float32Regs, Float32Regs, int_nvvm_fmax_f>;
 def INT_NVVM_FMAX_FTZ_F : F_MATH_2<"max.ftz.f32 \t$dst, $src0, $src1;",
   Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmax_ftz_f>;
 
+def INT_NVVM_FMAX_NAN_F : F_MATH_2<"max.NaN.f32 \t$dst, $src0, $src1;",
+  Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmax_nan_f,
+  [hasPTX70, hasSM80]>;
+def INT_NVVM_FMAX_FTZ_NAN_F : F_MATH_2<"max.ftz.NaN.f32 \t$dst, $src0, $src1;",
+  Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmax_ftz_nan_f,
+  [hasPTX70, hasSM80]>;
+
 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 C, Intrinsic I, NVPTXRegClass RC> {
+  string Capacity = C;
+  Intrinsic Intr = I;
+  NVPTXRegClass RegClass = RC;
+}
+
+multiclass MIN_MAX<string IntName> {
+  foreach P = [
+    MIN_MAX_TUPLE<"_f16", !if(!eq(IntName, "min"), int_nvvm_fmin_f16,
+      int_nvvm_fmax_f16), Float16Regs>,
+    MIN_MAX_TUPLE<"_ftz_f16", !if(!eq(IntName, "min"), int_nvvm_fmin_ftz_f16,
+      int_nvvm_fmax_ftz_f16), Float16Regs>,
+    MIN_MAX_TUPLE<"_NaN_f16", !if(!eq(IntName, "min"), int_nvvm_fmin_nan_f16,
+      int_nvvm_fmax_nan_f16), Float16Regs>,
+    MIN_MAX_TUPLE<"_ftz_NaN_f16", !if(!eq(IntName, "min"),
+      int_nvvm_fmin_ftz_nan_f16, int_nvvm_fmax_ftz_nan_f16), Float16Regs>,
+    MIN_MAX_TUPLE<"_f16x2", !if(!eq(IntName, "min"), int_nvvm_fmin_f16x2,
+      int_nvvm_fmax_f16x2), Float16x2Regs>,
+    MIN_MAX_TUPLE<"_ftz_f16x2", !if(!eq(IntName, "min"),
+      int_nvvm_fmin_ftz_f16x2, int_nvvm_fmax_ftz_f16x2), Float16x2Regs>,
+    MIN_MAX_TUPLE<"_NaN_f16x2", !if(!eq(IntName, "min"),
+      int_nvvm_fmin_nan_f16x2, int_nvvm_fmax_nan_f16x2), Float16x2Regs>,
+    MIN_MAX_TUPLE<"_ftz_NaN_f16x2", !if(!eq(IntName, "min"),
+      int_nvvm_fmin_ftz_nan_f16x2, int_nvvm_fmax_ftz_nan_f16x2), Float16x2Regs>,
+    MIN_MAX_TUPLE<"_bf16", !if(!eq(IntName, "min"),
+      int_nvvm_fmin_bf16, int_nvvm_fmax_bf16), Int16Regs>,
+    MIN_MAX_TUPLE<"_NaN_bf16", !if(!eq(IntName, "min"), int_nvvm_fmin_nan_bf16,
+      int_nvvm_fmax_nan_bf16), Int16Regs>,
+    MIN_MAX_TUPLE<"_bf16x2", !if(!eq(IntName, "min"), int_nvvm_fmin_bf16x2,
+      int_nvvm_fmax_bf16x2), Int32Regs>,
+    MIN_MAX_TUPLE<"_NaN_bf16x2", !if(!eq(IntName, "min"),
+      int_nvvm_fmin_nan_bf16x2, int_nvvm_fmax_nan_bf16x2), Int32Regs>] in {
+
+        def P.Capacity : F_MATH_2<!strconcat(
+          IntName, !subst("_", ".", P.Capacity), " \t$dst, $src0, $src1;"),
+          P.RegClass, P.RegClass, P.RegClass, P.Intr, [hasPTX70, hasSM80]>;
+  }
+}
+
+defm INT_NVVM_FMIN : MIN_MAX<"min">;
+defm INT_NVVM_FMAN : MIN_MAX<"max">;
 
 //
 // Multiplication
 //
 
 def INT_NVVM_MULHI_I : F_MATH_2<"mul.hi.s32 \t$dst, $src0, $src1;", Int32Regs,
   Int32Regs, Int32Regs, int_nvvm_mulhi_i>;
 def INT_NVVM_MULHI_UI : F_MATH_2<"mul.hi.u32 \t$dst, $src0, $src1;", Int32Regs,
def INT_NVVM_FABS_FTZ_F : F_MATH_1<"abs.ftz.f32 \t$dst, $src0;", Float32Regs,
   Float32Regs, int_nvvm_fabs_ftz_f>;
 def INT_NVVM_FABS_F : F_MATH_1<"abs.f32 \t$dst, $src0;", Float32Regs,
   Float32Regs, int_nvvm_fabs_f>;
 
 def INT_NVVM_FABS_D : F_MATH_1<"abs.f64 \t$dst, $src0;", Float64Regs,
   Float64Regs, int_nvvm_fabs_d>;
 
 //
+// bf16, bf16x2 abs, neg
+//
+def INT_NVVM_ABS_BF16 : F_MATH_1<"abs.bf16 \t$dst, $dst;", Int16Regs,
+  Int16Regs, int_nvvm_abs_bf16, [hasPTX70, hasSM80]>;
+def INT_NVVM_ABS_BF16X2 : F_MATH_1<"abs.bf16x2 \t$dst, $dst;", Int32Regs,
+  Int32Regs, int_nvvm_abs_bf16x2, [hasPTX70, hasSM80]>;
+def INT_NVVM_NEG_BF16 : F_MATH_1<"neg.bf16 \t$dst, $dst;", Int16Regs,
+  Int16Regs, int_nvvm_neg_bf16, [hasPTX70, hasSM80]>;
+def INT_NVVM_NEG_BF16X2 : F_MATH_1<"neg.bf16x2 \t$dst, $dst;", Int32Regs,
+  Int32Regs, int_nvvm_neg_bf16x2, [hasPTX70, hasSM80]>;
+
+//
 // Round
 //
 
 def : Pat<(int_nvvm_round_ftz_f Float32Regs:$a),
           (CVT_f32_f32 Float32Regs:$a, CvtRNI_FTZ)>;
 def : Pat<(int_nvvm_round_f Float32Regs:$a),
           (CVT_f32_f32 Float32Regs:$a, CvtRNI)>;
 def : Pat<(int_nvvm_round_d Float64Regs:$a),

llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp

static Instruction *simplifyNvvmIntrinsic(IntrinsicInst *II, InstCombiner &IC) {
   struct SimplifyAction {
     // Invariant: At most one of these Optionals has a value.
     Optional<Intrinsic::ID> IID;
     Optional<Instruction::CastOps> CastOp;
     Optional<Instruction::BinaryOps> BinaryOp;
     Optional<SpecialCase> Special;
 
     FtzRequirementTy FtzRequirement = FTZ_Any;
+    // Denormal handling is guarded by different attributes depending on the
+    // type (denormal-fp-math vs denormal-fp-math-f32), take note of halfs.
+    bool IsHalfTy = false;
 
     SimplifyAction() = default;
 
-    SimplifyAction(Intrinsic::ID IID, FtzRequirementTy FtzReq)
-        : IID(IID), FtzRequirement(FtzReq) {}
+    SimplifyAction(Intrinsic::ID IID, FtzRequirementTy FtzReq,
+                   bool IsHalfTy = false)
+        : IID(IID), FtzRequirement(FtzReq), IsHalfTy(IsHalfTy) {}
 
     // Cast operations don't have anything to do with FTZ, so we skip that
     // argument.
     SimplifyAction(Instruction::CastOps CastOp) : CastOp(CastOp) {}
 
     SimplifyAction(Instruction::BinaryOps BinaryOp, FtzRequirementTy FtzReq)
         : BinaryOp(BinaryOp), FtzRequirement(FtzReq) {}
 
const SimplifyAction Action = [II]() -> SimplifyAction {
     case Intrinsic::nvvm_fma_rn_ftz_f:
       return {Intrinsic::fma, FTZ_MustBeOn};
     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_f16:
+      return {Intrinsic::maxnum, FTZ_MustBeOff, true};
+    case Intrinsic::nvvm_fmax_ftz_f16:
+      return {Intrinsic::maxnum, FTZ_MustBeOn, true};
+    case Intrinsic::nvvm_fmax_f16x2:
+      return {Intrinsic::maxnum, FTZ_MustBeOff, true};
+    case Intrinsic::nvvm_fmax_ftz_f16x2:
+      return {Intrinsic::maxnum, FTZ_MustBeOn, true};
     case Intrinsic::nvvm_fmin_d:
       return {Intrinsic::minnum, FTZ_Any};
     case Intrinsic::nvvm_fmin_f:
       return {Intrinsic::minnum, FTZ_MustBeOff};
     case Intrinsic::nvvm_fmin_ftz_f:
       return {Intrinsic::minnum, FTZ_MustBeOn};
+    case Intrinsic::nvvm_fmin_f16:
+      return {Intrinsic::minnum, FTZ_MustBeOff, true};
+    case Intrinsic::nvvm_fmin_ftz_f16:
+      return {Intrinsic::minnum, FTZ_MustBeOn, true};
+    case Intrinsic::nvvm_fmin_f16x2:
+      return {Intrinsic::minnum, FTZ_MustBeOff, true};
+    case Intrinsic::nvvm_fmin_ftz_f16x2:
+      return {Intrinsic::minnum, FTZ_MustBeOn, true};
     case Intrinsic::nvvm_round_d:
       return {Intrinsic::round, FTZ_Any};
     case Intrinsic::nvvm_round_f:
       return {Intrinsic::round, FTZ_MustBeOff};
     case Intrinsic::nvvm_round_ftz_f:
       return {Intrinsic::round, FTZ_MustBeOn};
     case Intrinsic::nvvm_sqrt_rn_d:
       return {Intrinsic::sqrt, FTZ_Any};
const SimplifyAction Action = [II]() -> SimplifyAction {
     }
   }();
 
   // If Action.FtzRequirementTy is not satisfied by the module's ftz state, we
   // can bail out now.  (Notice that in the case that IID is not an NVVM
   // intrinsic, we don't have to look up any module metadata, as
   // FtzRequirementTy will be FTZ_Any.)
   if (Action.FtzRequirement != FTZ_Any) {
-    StringRef Attr = II->getFunction()
-                         ->getFnAttribute("denormal-fp-math-f32")
-                         .getValueAsString();
+    const char *AttrName =
+        Action.IsHalfTy ? "denormal-fp-math" : "denormal-fp-math-f32";
+    StringRef Attr =
+        II->getFunction()->getFnAttribute(AttrName).getValueAsString();
     DenormalMode Mode = parseDenormalFPAttribute(Attr);
     bool FtzEnabled = Mode.Output != DenormalMode::IEEE;
 
     if (FtzEnabled != (Action.FtzRequirement == FTZ_MustBeOn))
       return nullptr;
   }
 
   // Simplify to target-generic intrinsic.

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

+; RUN: opt < %s -instcombine -S -mtriple=nvptx-nvidia-cuda -march=nvptx64 \
+; RUN:    -mcpu=sm_80 -mattr=+ptx70 | \
+; RUN: FileCheck %s
+
+declare half @llvm.nvvm.fmin.f16(half, half)
+declare half @llvm.nvvm.fmin.ftz.f16(half, half)
+declare <2 x half> @llvm.nvvm.fmin.f16x2(<2 x half>, <2 x half>)
+declare <2 x half> @llvm.nvvm.fmin.ftz.f16x2(<2 x half>, <2 x half>)
+
+declare half @llvm.nvvm.fmax.f16(half, half)
+declare half @llvm.nvvm.fmax.ftz.f16(half, half)
+declare <2 x half> @llvm.nvvm.fmax.f16x2(<2 x half>, <2 x half>)
+declare <2 x half> @llvm.nvvm.fmax.ftz.f16x2(<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: fmin_ftz_f16
+define half @fmin_ftz_f16(half %0, half %1) #0 {
+  ; CHECK-NOT: @llvm.nvvm.fmin.ftz.f16
+  ; CHECK: @llvm.minnum.f16
+  %res = call half @llvm.nvvm.fmin.ftz.f16(half %0, half %1)
+  ret half %res
+}
+
+; CHECK-LABEL: fmin_ftz_f16_no_attr
+define half @fmin_ftz_f16_no_attr(half %0, half %1) {
+  ; CHECK-NOT: @llvm.minnum.f16
+  ; CHECK: @llvm.nvvm.fmin.ftz.f16
+  %res = call half @llvm.nvvm.fmin.ftz.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: @llvm.nvvm.fmin.f16x2
+  ; CHECK: @llvm.minnum.v2f16
+  %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) #0 {
+  ; CHECK-NOT: @llvm.nvvm.fmin.ftz.f16x2
+  ; CHECK: @llvm.minnum.v2f16
+  %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_ftz_f16x2_no_attr
+define <2 x half> @fmin_ftz_f16x2_no_attr(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: @llvm.minnum.v2f16
+  ; CHECK: @llvm.nvvm.fmin.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: fmax_f16
+define half @fmax_f16(half %0, half %1) {
+  ; CHECK-NOT: @llvm.nvvm.fmax.f16
+  ; CHECK: @llvm.maxnum.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) #0 {
+  ; CHECK-NOT: @llvm.nvvm.fmax.ftz.f16
+  ; CHECK: @llvm.maxnum.f16
+  %res = call half @llvm.nvvm.fmax.ftz.f16(half %0, half %1)
+  ret half %res
+}
+
+; CHECK-LABEL: fmax_ftz_f16_no_attr
+define half @fmax_ftz_f16_no_attr(half %0, half %1) {
+  ; CHECK-NOT: @llvm.maxnum.f16
+  ; CHECK: @llvm.nvvm.fmax.ftz.f16
+  %res = call half @llvm.nvvm.fmax.ftz.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: @llvm.nvvm.fmax.f16x2
+  ; CHECK: @llvm.maxnum.v2f16
+  %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) #0 {
+  ; CHECK-NOT: @llvm.nvvm.fmax.ftz.f16x2
+  ; CHECK: @llvm.maxnum.v2f16
+  %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_ftz_f16x2_no_attr
+define <2 x half> @fmax_ftz_f16x2_no_attr(<2 x half> %0, <2 x half> %1) {
+  ; CHECK-NOT: @llvm.maxnum.v2f16
+  ; CHECK: @llvm.nvvm.fmax.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
+}
+
+attributes #0 = { "denormal-fp-math"="preserve-sign" }

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

+; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | FileCheck %s
+
+declare i16 @llvm.nvvm.abs.bf16(i16)
+declare i32 @llvm.nvvm.abs.bf16x2(i32)
+declare i16 @llvm.nvvm.neg.bf16(i16)
+declare i32 @llvm.nvvm.neg.bf16x2(i32)
+
+declare float @llvm.nvvm.fmin.nan.f(float, float)
+declare float @llvm.nvvm.fmin.ftz.nan.f(float, float)
+declare half @llvm.nvvm.fmin.f16(half, half)
+declare half @llvm.nvvm.fmin.ftz.f16(half, half)
+declare half @llvm.nvvm.fmin.nan.f16(half, half)
+declare half @llvm.nvvm.fmin.ftz.nan.f16(half, half)
+declare <2 x half> @llvm.nvvm.fmin.f16x2(<2 x half>, <2 x half>)
+declare <2 x half> @llvm.nvvm.fmin.ftz.f16x2(<2 x half>, <2 x half>)
+declare <2 x half> @llvm.nvvm.fmin.nan.f16x2(<2 x half>, <2 x half>)
+declare <2 x half> @llvm.nvvm.fmin.ftz.nan.f16x2(<2 x half>, <2 x half>)
+declare i16 @llvm.nvvm.fmin.bf16(i16, i16)
+declare i16 @llvm.nvvm.fmin.nan.bf16(i16, i16)
+declare i32 @llvm.nvvm.fmin.bf16x2(i32, i32)
+declare i32 @llvm.nvvm.fmin.nan.bf16x2(i32, i32)
+
+declare float @llvm.nvvm.fmax.nan.f(float, float)
+declare float @llvm.nvvm.fmax.ftz.nan.f(float, float)
+declare half @llvm.nvvm.fmax.f16(half, half)
+declare half @llvm.nvvm.fmax.ftz.f16(half, half)
+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.f16x2(<2 x half>, <2 x half>)
+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)
+
+; CHECK-LABEL: abs_bf16
+define i16 @abs_bf16(i16 %0) {
+  ; 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: 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: max.NaN.bf16x2
+  %res = call i32 @llvm.nvvm.fmax.nan.bf16x2(i32 %0, i32 %1)
+  ret i32 %res
+}