jchlanda created this revision.
jchlanda added reviewers: csigg, tra, bkramer.
Herald added subscribers: asavonic, hiraditya, jholewinski.
jchlanda requested review of this revision.
Herald added subscribers: llvm-commits, cfe-commits, jdoerfert.
Herald added projects: clang, LLVM.
Adds support for the following builtins:
- abs, neg:
- bf16,
- bf16x2
- min, max
- f16, f16 ftz, f16 nan, f16 nan ftz
- f16x2, f16x2 ftz, f16x2 nan, f16x2 nan ftz
- bf16, bf16 nan
- bf16x2, bf16x2 nan
- f32 ftz, f32 nan, f32 nan ftz
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D117887
Files:
clang/include/clang/Basic/BuiltinsNVPTX.def
clang/test/CodeGen/builtins-nvptx-native-half-type.c
clang/test/CodeGen/builtins-nvptx.c
llvm/include/llvm/IR/IntrinsicsNVVM.td
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp
llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70-instcombine.ll
llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70.ll
Index: llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70.ll
===================================================================
--- /dev/null
+++ llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70.ll
@@ -0,0 +1,260 @@
+; 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
+}
Index: llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70-instcombine.ll
===================================================================
--- /dev/null
+++ llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70-instcombine.ll
@@ -0,0 +1,268 @@
+; 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 float @llvm.nvvm.fmin.nan.f(float, float)
+declare float @llvm.nvvm.fmin.ftz.nan.f(float, float)
+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.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 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>)
+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>)
+
+; 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: fmin_nan_f
+define float @fmin_nan_f(float %0, float %1) {
+ ; CHECK-NOT: @llvm.nvvm.fmin.nan.f
+ ; CHECK: @llvm.minimum.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) #1 {
+ ; CHECK-NOT: @llvm.nvvm.fmin.ftz.nan.f
+ ; CHECK: @llvm.minimum.f32
+ %res = call float @llvm.nvvm.fmin.ftz.nan.f(float %0, float %1)
+ ret float %res
+}
+
+; CHECK-LABEL: fmin_ftz_nan_f_no_attr
+define float @fmin_ftz_nan_f_no_attr(float %0, float %1) {
+ ; CHECK: @llvm.nvvm.fmin.ftz.nan.f
+ ; CHECK-NOT: @llvm.minimum.f32
+ %res = call float @llvm.nvvm.fmin.ftz.nan.f(float %0, float %1)
+ ret float %res
+}
+
+; CHECK-LABEL: fmin_nan_f16
+define half @fmin_nan_f16(half %0, half %1) {
+ ; CHECK-NOT: @llvm.nvvm.fmin.nan.f16
+ ; CHECK: @llvm.minimum.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) #0 {
+ ; CHECK-NOT: @llvm.nvvm.fmin.ftz.nan.f16
+ ; CHECK: @llvm.minimum.f16
+ %res = call half @llvm.nvvm.fmin.ftz.nan.f16(half %0, half %1)
+ ret half %res
+}
+
+; CHECK-LABEL: fmin_ftz_nan_f16_no_attr
+define half @fmin_ftz_nan_f16_no_attr(half %0, half %1) {
+ ; CHECK: @llvm.nvvm.fmin.ftz.nan.f16
+ ; CHECK-NOT: @llvm.minimum.f16
+ %res = call half @llvm.nvvm.fmin.ftz.nan.f16(half %0, half %1)
+ ret half %res
+}
+
+; CHECK-LABEL: fmin_nan_f16x2
+define <2 x half> @fmin_nan_f16x2(<2 x half> %0, <2 x half> %1) {
+ ; CHECK-NOT: @llvm.nvvm.fmin.nan.f16x2
+ ; CHECK: @llvm.minimum.v2f16
+ %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) #0 {
+ ; CHECK-NOT: @llvm.nvvm.fmin.ftz.nan.f16x2
+ ; CHECK: @llvm.minimum.v2f16
+ %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_ftz_nan_f16x2_no_attr
+define <2 x half> @fmin_ftz_nan_f16x2_no_attr(<2 x half> %0, <2 x half> %1) {
+ ; CHECK-NOT: @llvm.minimum.v2f16
+ ; CHECK: @llvm.nvvm.fmin.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: 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
+}
+
+; CHECK-LABEL: fmax_nan_f
+define float @fmax_nan_f(float %0, float %1) {
+ ; CHECK-NOT: @llvm.nvvm.fmax.nan.f
+ ; CHECK: @llvm.maximum.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) #1 {
+ ; CHECK-NOT: @llvm.nvvm.fmax.ftz.nan.f
+ ; CHECK: @llvm.maximum.f32
+ %res = call float @llvm.nvvm.fmax.ftz.nan.f(float %0, float %1)
+ ret float %res
+}
+
+; CHECK-LABEL: fmax_ftz_nan_f_no_attr
+define float @fmax_ftz_nan_f_no_attr(float %0, float %1) {
+ ; CHECK: @llvm.nvvm.fmax.ftz.nan.f
+ ; CHECK-NOT: @llvm.maximum.f32
+ %res = call float @llvm.nvvm.fmax.ftz.nan.f(float %0, float %1)
+ ret float %res
+}
+
+; CHECK-LABEL: fmax_nan_f16
+define half @fmax_nan_f16(half %0, half %1) {
+ ; CHECK-NOT: @llvm.nvvm.fmax.nan.f16
+ ; CHECK: @llvm.maximum.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) #0 {
+ ; CHECK-NOT: @llvm.nvvm.fmax.ftz.nan.f16
+ ; CHECK: @llvm.maximum.f16
+ %res = call half @llvm.nvvm.fmax.ftz.nan.f16(half %0, half %1)
+ ret half %res
+}
+
+; CHECK-LABEL: fmax_ftz_nan_f16_no_attr
+define half @fmax_ftz_nan_f16_no_attr(half %0, half %1) {
+ ; CHECK: @llvm.nvvm.fmax.ftz.nan.f16
+ ; CHECK-NOT: @llvm.maximum.f16
+ %res = call half @llvm.nvvm.fmax.ftz.nan.f16(half %0, half %1)
+ ret half %res
+}
+
+; CHECK-LABEL: fmax_nan_f16x2
+define <2 x half> @fmax_nan_f16x2(<2 x half> %0, <2 x half> %1) {
+ ; CHECK-NOT: @llvm.nvvm.fmax.nan.f16x2
+ ; CHECK: @llvm.maximum.v2f16
+ %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) #0 {
+ ; CHECK-NOT: @llvm.nvvm.fmax.ftz.nan.f16x2
+ ; CHECK: @llvm.maximum.v2f16
+ %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_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
+}
+
+attributes #0 = { "denormal-fp-math"="preserve-sign" }
+attributes #1 = { "denormal-fp-math-f32"="preserve-sign" }
Index: llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp
===================================================================
--- llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp
+++ llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp
@@ -145,11 +145,15 @@
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.
@@ -197,12 +201,52 @@
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};
+ case Intrinsic::nvvm_fmax_ftz_nan_f:
+ return {Intrinsic::maximum, 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_fmax_nan_f16:
+ return {Intrinsic::maximum, FTZ_MustBeOff, true};
+ case Intrinsic::nvvm_fmax_ftz_nan_f16:
+ return {Intrinsic::maximum, FTZ_MustBeOn, true};
+ case Intrinsic::nvvm_fmax_nan_f16x2:
+ return {Intrinsic::maximum, FTZ_MustBeOff, true};
+ case Intrinsic::nvvm_fmax_ftz_nan_f16x2:
+ return {Intrinsic::maximum, 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_nan_f:
+ return {Intrinsic::minimum, FTZ_MustBeOff};
+ case Intrinsic::nvvm_fmin_ftz_nan_f:
+ return {Intrinsic::minimum, 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_fmin_nan_f16:
+ return {Intrinsic::minimum, FTZ_MustBeOff, true};
+ case Intrinsic::nvvm_fmin_ftz_nan_f16:
+ return {Intrinsic::minimum, FTZ_MustBeOn, true};
+ case Intrinsic::nvvm_fmin_nan_f16x2:
+ return {Intrinsic::minimum, FTZ_MustBeOff, true};
+ case Intrinsic::nvvm_fmin_ftz_nan_f16x2:
+ return {Intrinsic::minimum, FTZ_MustBeOn, true};
case Intrinsic::nvvm_round_d:
return {Intrinsic::round, FTZ_Any};
case Intrinsic::nvvm_round_f:
@@ -316,9 +360,10 @@
// 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;
Index: llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
===================================================================
--- llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -549,19 +549,22 @@
// 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,
@@ -587,17 +590,73 @@
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
@@ -719,6 +778,19 @@
def INT_NVVM_FABS_D : F_MATH_1<"abs.f64 \t$dst, $src0;", Float64Regs,
Float64Regs, int_nvvm_fabs_d>;
+//
+// Abs, Neg bf16, bf16x2
+//
+
+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
//
Index: llvm/include/llvm/IR/IntrinsicsNVVM.td
===================================================================
--- llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -564,26 +564,48 @@
// Min Max
//
- def int_nvvm_fmin_f : GCCBuiltin<"__nvvm_fmin_f">,
- DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
- [IntrNoMem, IntrSpeculatable, Commutative]>;
- def int_nvvm_fmin_ftz_f : GCCBuiltin<"__nvvm_fmin_ftz_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_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],
- [IntrNoMem, IntrSpeculatable, Commutative]>;
+ 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_fmin_d : GCCBuiltin<"__nvvm_fmin_d">,
- DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
- [IntrNoMem, IntrSpeculatable, Commutative]>;
- def int_nvvm_fmax_d : GCCBuiltin<"__nvvm_fmax_d">,
- DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
- [IntrNoMem, IntrSpeculatable, Commutative]>;
+ 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]>;
+ }
+
+ 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]>;
+ }
+
+ 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
@@ -740,6 +762,19 @@
def int_nvvm_fabs_d : GCCBuiltin<"__nvvm_fabs_d">,
DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>;
+//
+// Abs, Neg bf16, bf16x2
+//
+
+ 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]>;
+ }
+
//
// Round
//
Index: clang/test/CodeGen/builtins-nvptx.c
===================================================================
--- clang/test/CodeGen/builtins-nvptx.c
+++ clang/test/CodeGen/builtins-nvptx.c
@@ -791,3 +791,62 @@
#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
+}
Index: clang/test/CodeGen/builtins-nvptx-native-half-type.c
===================================================================
--- /dev/null
+++ clang/test/CodeGen/builtins-nvptx-native-half-type.c
@@ -0,0 +1,53 @@
+// 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
+}
Index: clang/include/clang/Basic/BuiltinsNVPTX.def
===================================================================
--- clang/include/clang/Basic/BuiltinsNVPTX.def
+++ clang/include/clang/Basic/BuiltinsNVPTX.def
@@ -107,13 +107,41 @@
// Min Max
-BUILTIN(__nvvm_fmax_ftz_f, "fff", "")
-BUILTIN(__nvvm_fmax_f, "fff", "")
-BUILTIN(__nvvm_fmin_ftz_f, "fff", "")
+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_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))
+BUILTIN(__nvvm_fmin_d, "ddd", "")
+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_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))
BUILTIN(__nvvm_fmax_d, "ddd", "")
-BUILTIN(__nvvm_fmin_d, "ddd", "")
// Multiplication
@@ -827,6 +855,13 @@
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")
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
