https://github.com/ayokunle321 created
https://github.com/llvm/llvm-project/pull/221859
Adds codegen for the following NVVM half type builtins:
- `__nvvm_ex2_approx_{f16,f16x2}`
- `__nvvm_fmin{_ftz,}{_nan,}{_xorsign_abs,}_{f16,f16x2}`
- `__nvvm_fmax{_ftz,}{_nan,}{_xorsign_abs,}_{f16,f16x2}`
- `__nvvm_fma_rn{_ftz,}{_relu,_sat,}_{f16,f16x2}`
- `__nvvm_fma_rn_oob{_relu,}_{f16,bf16,f16x2,bf16x2}`
These are lowered to the corresponding llvm.nvvm.* intrinsics.
This diverges a bit from OG, as no half-type helper is needed here. OG routes
these through `MakeHalfType`, which coerces each argument to the intrinsic's
declared parameter type, appends the intrinsic's defaulted arguments, and
asserts that no argument is an integer constant expression. None of that is
needed for the builtins in this patch as each already agrees with its intrinsic
on operand and result types, and none has defaulted parameters. The existing
`emitUnaryNVVMIntrinsic` and `emitBuiltinWithOneOverloadedType<N>` helpers are
therefore sufficient.
More generally, nothing `MakeHalfType` does is half-specific, and its coercion
looks unexercises, i.e., the NVVM intrinsics whose builtins carry defaulted
arguments are resolved by `Intrinsic::getIntrinsicForClangBuiltin` and handled
by the generic intrinsic path before target dispatch, which does its own
argument coercion and appends the defaults. The one part of `MakeHalfType` that
is load-bearing is its trailing-argument parameter, used by the `_pzo`
conversion variants; those are
not added here.
Assisted by Claude.
>From 26b4f025b0493f738bbdd2d5f2fe540e4254cb6a Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <[email protected]>
Date: Tue, 8 Sep 2026 02:42:29 +0200
Subject: [PATCH] add half type builtins
---
clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp | 330 +++++----------
.../CIR/CodeGenCUDA/builtins-nvvm-fma-oob.cu | 94 +++++
.../CIR/CodeGenCUDA/builtins-nvvm-half.cu | 391 ++++++++++++++++++
3 files changed, 597 insertions(+), 218 deletions(-)
create mode 100644 clang/test/CIR/CodeGenCUDA/builtins-nvvm-fma-oob.cu
create mode 100644 clang/test/CIR/CodeGenCUDA/builtins-nvvm-half.cu
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
index 8b33d85d49fe1..9bf5660100c14 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
@@ -517,16 +517,6 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId,
const CallExpr *expr) {
getContext().BuiltinInfo.getName(builtinId));
return mlir::Value{};
// The following builtins require half type support
- case NVPTX::BI__nvvm_ex2_approx_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- case NVPTX::BI__nvvm_ex2_approx_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
case NVPTX::BI__nvvm_ff2f16x2_rn:
cgm.errorNYI(expr->getSourceRange(),
std::string("unimplemented NVPTX builtin call: ") +
@@ -548,265 +538,167 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned
builtinId, const CallExpr *expr) {
getContext().BuiltinInfo.getName(builtinId));
return mlir::Value{};
case NVPTX::BI__nvvm_fma_rn_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.f16")
+ .getValue();
case NVPTX::BI__nvvm_fma_rn_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fma_rn_ftz_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.ftz.f16")
+ .getValue();
case NVPTX::BI__nvvm_fma_rn_ftz_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.ftz.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fma_rn_ftz_relu_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<3>(expr,
"nvvm.fma.rn.ftz.relu.f16")
+ .getValue();
case NVPTX::BI__nvvm_fma_rn_ftz_relu_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<3>(expr,
+ "nvvm.fma.rn.ftz.relu.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fma_rn_ftz_sat_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.ftz.sat.f16")
+ .getValue();
case NVPTX::BI__nvvm_fma_rn_ftz_sat_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<3>(expr,
+ "nvvm.fma.rn.ftz.sat.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fma_rn_relu_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.relu.f16")
+ .getValue();
case NVPTX::BI__nvvm_fma_rn_relu_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.relu.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fma_rn_sat_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.sat.f16")
+ .getValue();
case NVPTX::BI__nvvm_fma_rn_sat_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.sat.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fma_rn_oob_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
case NVPTX::BI__nvvm_fma_rn_oob_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
case NVPTX::BI__nvvm_fma_rn_oob_bf16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
case NVPTX::BI__nvvm_fma_rn_oob_bf16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.oob")
+ .getValue();
case NVPTX::BI__nvvm_fma_rn_oob_relu_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
case NVPTX::BI__nvvm_fma_rn_oob_relu_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
case NVPTX::BI__nvvm_fma_rn_oob_relu_bf16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
case NVPTX::BI__nvvm_fma_rn_oob_relu_bf16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.oob.relu")
+ .getValue();
case NVPTX::BI__nvvm_fmax_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmax.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmax_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmax.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmax_ftz_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmax.ftz.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmax_ftz_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmax.ftz.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmax_ftz_nan_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmax.ftz.nan.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmax_ftz_nan_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmax.ftz.nan.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmax_ftz_nan_xorsign_abs_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(
+ expr, "nvvm.fmax.ftz.nan.xorsign.abs.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmax_ftz_nan_xorsign_abs_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(
+ expr, "nvvm.fmax.ftz.nan.xorsign.abs.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmax_ftz_xorsign_abs_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr,
+ "nvvm.fmax.ftz.xorsign.abs.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmax_ftz_xorsign_abs_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(
+ expr, "nvvm.fmax.ftz.xorsign.abs.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmax_nan_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmax.nan.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmax_nan_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmax.nan.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmax_nan_xorsign_abs_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr,
+ "nvvm.fmax.nan.xorsign.abs.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmax_nan_xorsign_abs_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(
+ expr, "nvvm.fmax.nan.xorsign.abs.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmax_xorsign_abs_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr,
+ "nvvm.fmax.xorsign.abs.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmax_xorsign_abs_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr,
+ "nvvm.fmax.xorsign.abs.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmin_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmin.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmin_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmin.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmin_ftz_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmin.ftz.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmin_ftz_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmin.ftz.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmin_ftz_nan_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmin.ftz.nan.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmin_ftz_nan_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmin.ftz.nan.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmin_ftz_nan_xorsign_abs_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(
+ expr, "nvvm.fmin.ftz.nan.xorsign.abs.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmin_ftz_nan_xorsign_abs_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(
+ expr, "nvvm.fmin.ftz.nan.xorsign.abs.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmin_ftz_xorsign_abs_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr,
+ "nvvm.fmin.ftz.xorsign.abs.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmin_ftz_xorsign_abs_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(
+ expr, "nvvm.fmin.ftz.xorsign.abs.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmin_nan_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmin.nan.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmin_nan_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr, "nvvm.fmin.nan.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmin_nan_xorsign_abs_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr,
+ "nvvm.fmin.nan.xorsign.abs.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmin_nan_xorsign_abs_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(
+ expr, "nvvm.fmin.nan.xorsign.abs.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fmin_xorsign_abs_f16:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr,
+ "nvvm.fmin.xorsign.abs.f16")
+ .getValue();
case NVPTX::BI__nvvm_fmin_xorsign_abs_f16x2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return emitBuiltinWithOneOverloadedType<2>(expr,
+ "nvvm.fmin.xorsign.abs.f16x2")
+ .getValue();
case NVPTX::BI__nvvm_fabs_f:
case NVPTX::BI__nvvm_abs_bf16:
case NVPTX::BI__nvvm_abs_bf16x2:
@@ -821,6 +713,8 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId,
const CallExpr *expr) {
return emitUnaryNVVMIntrinsic(*this, expr, "fabs");
case NVPTX::BI__nvvm_ex2_approx_d:
case NVPTX::BI__nvvm_ex2_approx_f:
+ case NVPTX::BI__nvvm_ex2_approx_f16:
+ case NVPTX::BI__nvvm_ex2_approx_f16x2:
return emitUnaryNVVMIntrinsic(*this, expr, "nvvm.ex2.approx");
case NVPTX::BI__nvvm_ex2_approx_ftz_f:
return emitUnaryNVVMIntrinsic(*this, expr, "nvvm.ex2.approx.ftz");
diff --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-fma-oob.cu
b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-fma-oob.cu
new file mode 100644
index 0000000000000..2385f6d851d39
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-fma-oob.cu
@@ -0,0 +1,94 @@
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_90a \
+// RUN: -target-feature +ptx87 -x cuda -fcuda-is-device -fclangir
-emit-cir \
+// RUN: %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_90a \
+// RUN: -target-feature +ptx87 -x cuda -fcuda-is-device -fclangir
-emit-llvm \
+// RUN: %s -o %t-cir.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_90a \
+// RUN: -target-feature +ptx87 -x cuda -fcuda-is-device -emit-llvm \
+// RUN: %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// FIXME: CIR doesn't propagate the 'contract' fast-math flag to LLVM IR
+// calls yet, so LLVM check lines use {{.*}} to tolerate the difference
+// between CIR (no flags) and classic codegen ('contract').
+
+// A scalar __fp16 cannot be a parameter or return type, so _Float16 is
+// used instead; both are `half`, so the calls below are unaffected. The
+// functions are extern "C" so the check labels are exact.
+//
+// nvvm.fma.rn.oob and nvvm.fma.rn.oob.relu are overloaded on the operand
+// type, so they cover both half and bfloat and take the unsuffixed name.
+
+#define __device__ __attribute__((device))
+
+typedef __fp16 f16x2 __attribute__((ext_vector_type(2)));
+typedef __bf16 bf16x2 __attribute__((ext_vector_type(2)));
+
+// CIR-LABEL: @test_fma_rn_oob_bf16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.oob" {{.*}} : (!cir.bf16,
!cir.bf16, !cir.bf16) -> !cir.bf16
+// LLVM-LABEL: @test_fma_rn_oob_bf16(
+// LLVM: call{{.*}} bfloat @llvm.nvvm.fma.rn.oob.bf16(bfloat %{{.*}}, bfloat
%{{.*}}, bfloat %{{.*}})
+extern "C" __device__ __bf16 test_fma_rn_oob_bf16(__bf16 a0, __bf16 a1, __bf16
a2) {
+ return __nvvm_fma_rn_oob_bf16(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_oob_bf16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.oob" {{.*}} : (!cir.vector<2 x
!cir.bf16>, !cir.vector<2 x !cir.bf16>, !cir.vector<2 x !cir.bf16>) ->
!cir.vector<2 x !cir.bf16>
+// LLVM-LABEL: @test_fma_rn_oob_bf16x2(
+// LLVM: call{{.*}} <2 x bfloat> @llvm.nvvm.fma.rn.oob.v2bf16(<2 x bfloat>
%{{.*}}, <2 x bfloat> %{{.*}}, <2 x bfloat> %{{.*}})
+extern "C" __device__ bf16x2 test_fma_rn_oob_bf16x2(bf16x2 a0, bf16x2 a1,
bf16x2 a2) {
+ return __nvvm_fma_rn_oob_bf16x2(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_oob_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.oob" {{.*}} : (!cir.f16,
!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fma_rn_oob_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fma.rn.oob.f16(half %{{.*}}, half %{{.*}},
half %{{.*}})
+extern "C" __device__ _Float16 test_fma_rn_oob_f16(_Float16 a0, _Float16 a1,
_Float16 a2) {
+ return __nvvm_fma_rn_oob_f16(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_oob_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.oob" {{.*}} : (!cir.vector<2 x
!cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) ->
!cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_fma_rn_oob_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.rn.oob.v2f16(<2 x half> %{{.*}},
<2 x half> %{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fma_rn_oob_f16x2(f16x2 a0, f16x2 a1, f16x2
a2) {
+ return __nvvm_fma_rn_oob_f16x2(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_oob_relu_bf16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.oob.relu" {{.*}} : (!cir.bf16,
!cir.bf16, !cir.bf16) -> !cir.bf16
+// LLVM-LABEL: @test_fma_rn_oob_relu_bf16(
+// LLVM: call{{.*}} bfloat @llvm.nvvm.fma.rn.oob.relu.bf16(bfloat %{{.*}},
bfloat %{{.*}}, bfloat %{{.*}})
+extern "C" __device__ __bf16 test_fma_rn_oob_relu_bf16(__bf16 a0, __bf16 a1,
__bf16 a2) {
+ return __nvvm_fma_rn_oob_relu_bf16(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_oob_relu_bf16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.oob.relu" {{.*}} : (!cir.vector<2
x !cir.bf16>, !cir.vector<2 x !cir.bf16>, !cir.vector<2 x !cir.bf16>) ->
!cir.vector<2 x !cir.bf16>
+// LLVM-LABEL: @test_fma_rn_oob_relu_bf16x2(
+// LLVM: call{{.*}} <2 x bfloat> @llvm.nvvm.fma.rn.oob.relu.v2bf16(<2 x
bfloat> %{{.*}}, <2 x bfloat> %{{.*}}, <2 x bfloat> %{{.*}})
+extern "C" __device__ bf16x2 test_fma_rn_oob_relu_bf16x2(bf16x2 a0, bf16x2 a1,
bf16x2 a2) {
+ return __nvvm_fma_rn_oob_relu_bf16x2(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_oob_relu_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.oob.relu" {{.*}} : (!cir.f16,
!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fma_rn_oob_relu_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fma.rn.oob.relu.f16(half %{{.*}}, half
%{{.*}}, half %{{.*}})
+extern "C" __device__ _Float16 test_fma_rn_oob_relu_f16(_Float16 a0, _Float16
a1, _Float16 a2) {
+ return __nvvm_fma_rn_oob_relu_f16(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_oob_relu_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.oob.relu" {{.*}} : (!cir.vector<2
x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) ->
!cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_fma_rn_oob_relu_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.rn.oob.relu.v2f16(<2 x half>
%{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fma_rn_oob_relu_f16x2(f16x2 a0, f16x2 a1,
f16x2 a2) {
+ return __nvvm_fma_rn_oob_relu_f16x2(a0, a1, a2);
+}
diff --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-half.cu
b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-half.cu
new file mode 100644
index 0000000000000..d8dd88616d830
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-half.cu
@@ -0,0 +1,391 @@
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_90a \
+// RUN: -target-feature +ptx87 -x cuda -fcuda-is-device -fclangir
-emit-cir \
+// RUN: %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_90a \
+// RUN: -target-feature +ptx87 -x cuda -fcuda-is-device -fclangir
-emit-llvm \
+// RUN: %s -o %t-cir.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_90a \
+// RUN: -target-feature +ptx87 -x cuda -fcuda-is-device -emit-llvm \
+// RUN: %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// FIXME: CIR doesn't propagate the 'contract' fast-math flag to LLVM IR
+// calls yet, so LLVM check lines use {{.*}} to tolerate the difference
+// between CIR (no flags) and classic codegen ('contract').
+
+#define __device__ __attribute__((device))
+
+typedef __fp16 f16x2 __attribute__((ext_vector_type(2)));
+typedef __bf16 bf16x2 __attribute__((ext_vector_type(2)));
+
+// CIR-LABEL: @test_ex2_approx_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.ex2.approx" {{.*}} : (!cir.f16) ->
!cir.f16
+// LLVM-LABEL: @test_ex2_approx_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.ex2.approx.f16(half %{{.*}})
+extern "C" __device__ _Float16 test_ex2_approx_f16(_Float16 a0) {
+ return __nvvm_ex2_approx_f16(a0);
+}
+
+// CIR-LABEL: @test_ex2_approx_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.ex2.approx" {{.*}} : (!cir.vector<2 x
!cir.f16>) -> !cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_ex2_approx_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.ex2.approx.v2f16(<2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_ex2_approx_f16x2(f16x2 a0) {
+ return __nvvm_ex2_approx_f16x2(a0);
+}
+
+// CIR-LABEL: @test_fma_rn_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.f16" {{.*}} : (!cir.f16,
!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fma_rn_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fma.rn.f16(half %{{.*}}, half %{{.*}},
half %{{.*}})
+extern "C" __device__ _Float16 test_fma_rn_f16(_Float16 a0, _Float16 a1,
_Float16 a2) {
+ return __nvvm_fma_rn_f16(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.f16x2" {{.*}} : (!cir.vector<2 x
!cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) ->
!cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_fma_rn_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.rn.f16x2(<2 x half> %{{.*}}, <2
x half> %{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fma_rn_f16x2(f16x2 a0, f16x2 a1, f16x2 a2) {
+ return __nvvm_fma_rn_f16x2(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_ftz_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.ftz.f16" {{.*}} : (!cir.f16,
!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fma_rn_ftz_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fma.rn.ftz.f16(half %{{.*}}, half %{{.*}},
half %{{.*}})
+extern "C" __device__ _Float16 test_fma_rn_ftz_f16(_Float16 a0, _Float16 a1,
_Float16 a2) {
+ return __nvvm_fma_rn_ftz_f16(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_ftz_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.ftz.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x
!cir.f16>) -> !cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_fma_rn_ftz_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.rn.ftz.f16x2(<2 x half> %{{.*}},
<2 x half> %{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fma_rn_ftz_f16x2(f16x2 a0, f16x2 a1, f16x2
a2) {
+ return __nvvm_fma_rn_ftz_f16x2(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_ftz_relu_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.ftz.relu.f16" {{.*}} : (!cir.f16,
!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fma_rn_ftz_relu_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fma.rn.ftz.relu.f16(half %{{.*}}, half
%{{.*}}, half %{{.*}})
+extern "C" __device__ _Float16 test_fma_rn_ftz_relu_f16(_Float16 a0, _Float16
a1, _Float16 a2) {
+ return __nvvm_fma_rn_ftz_relu_f16(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_ftz_relu_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.ftz.relu.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x
!cir.f16>) -> !cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_fma_rn_ftz_relu_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.rn.ftz.relu.f16x2(<2 x half>
%{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fma_rn_ftz_relu_f16x2(f16x2 a0, f16x2 a1,
f16x2 a2) {
+ return __nvvm_fma_rn_ftz_relu_f16x2(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_ftz_sat_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.ftz.sat.f16" {{.*}} : (!cir.f16,
!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fma_rn_ftz_sat_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fma.rn.ftz.sat.f16(half %{{.*}}, half
%{{.*}}, half %{{.*}})
+extern "C" __device__ _Float16 test_fma_rn_ftz_sat_f16(_Float16 a0, _Float16
a1, _Float16 a2) {
+ return __nvvm_fma_rn_ftz_sat_f16(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_ftz_sat_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.ftz.sat.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x
!cir.f16>) -> !cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_fma_rn_ftz_sat_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.rn.ftz.sat.f16x2(<2 x half>
%{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fma_rn_ftz_sat_f16x2(f16x2 a0, f16x2 a1,
f16x2 a2) {
+ return __nvvm_fma_rn_ftz_sat_f16x2(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_relu_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.relu.f16" {{.*}} : (!cir.f16,
!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fma_rn_relu_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fma.rn.relu.f16(half %{{.*}}, half
%{{.*}}, half %{{.*}})
+extern "C" __device__ _Float16 test_fma_rn_relu_f16(_Float16 a0, _Float16 a1,
_Float16 a2) {
+ return __nvvm_fma_rn_relu_f16(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_relu_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.relu.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x
!cir.f16>) -> !cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_fma_rn_relu_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.rn.relu.f16x2(<2 x half>
%{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fma_rn_relu_f16x2(f16x2 a0, f16x2 a1, f16x2
a2) {
+ return __nvvm_fma_rn_relu_f16x2(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_sat_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.sat.f16" {{.*}} : (!cir.f16,
!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fma_rn_sat_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fma.rn.sat.f16(half %{{.*}}, half %{{.*}},
half %{{.*}})
+extern "C" __device__ _Float16 test_fma_rn_sat_f16(_Float16 a0, _Float16 a1,
_Float16 a2) {
+ return __nvvm_fma_rn_sat_f16(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fma_rn_sat_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.sat.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x
!cir.f16>) -> !cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_fma_rn_sat_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.rn.sat.f16x2(<2 x half> %{{.*}},
<2 x half> %{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fma_rn_sat_f16x2(f16x2 a0, f16x2 a1, f16x2
a2) {
+ return __nvvm_fma_rn_sat_f16x2(a0, a1, a2);
+}
+
+// CIR-LABEL: @test_fmax_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.f16" {{.*}} : (!cir.f16, !cir.f16)
-> !cir.f16
+// LLVM-LABEL: @test_fmax_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmax.f16(half %{{.*}}, half %{{.*}})
+extern "C" __device__ _Float16 test_fmax_f16(_Float16 a0, _Float16 a1) {
+ return __nvvm_fmax_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.f16x2" {{.*}} : (!cir.vector<2 x
!cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_fmax_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmax.f16x2(<2 x half> %{{.*}}, <2 x
half> %{{.*}})
+extern "C" __device__ f16x2 test_fmax_f16x2(f16x2 a0, f16x2 a1) {
+ return __nvvm_fmax_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_ftz_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.ftz.f16" {{.*}} : (!cir.f16,
!cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fmax_ftz_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmax.ftz.f16(half %{{.*}}, half %{{.*}})
+extern "C" __device__ _Float16 test_fmax_ftz_f16(_Float16 a0, _Float16 a1) {
+ return __nvvm_fmax_ftz_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_ftz_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.ftz.f16x2" {{.*}} : (!cir.vector<2
x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_fmax_ftz_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmax.ftz.f16x2(<2 x half> %{{.*}},
<2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fmax_ftz_f16x2(f16x2 a0, f16x2 a1) {
+ return __nvvm_fmax_ftz_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_ftz_nan_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.ftz.nan.f16" {{.*}} : (!cir.f16,
!cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fmax_ftz_nan_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmax.ftz.nan.f16(half %{{.*}}, half
%{{.*}})
+extern "C" __device__ _Float16 test_fmax_ftz_nan_f16(_Float16 a0, _Float16 a1)
{
+ return __nvvm_fmax_ftz_nan_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_ftz_nan_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.ftz.nan.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x
!cir.f16>
+// LLVM-LABEL: @test_fmax_ftz_nan_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmax.ftz.nan.f16x2(<2 x half>
%{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fmax_ftz_nan_f16x2(f16x2 a0, f16x2 a1) {
+ return __nvvm_fmax_ftz_nan_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_ftz_nan_xorsign_abs_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.ftz.nan.xorsign.abs.f16" {{.*}} :
(!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fmax_ftz_nan_xorsign_abs_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmax.ftz.nan.xorsign.abs.f16(half %{{.*}},
half %{{.*}})
+extern "C" __device__ _Float16 test_fmax_ftz_nan_xorsign_abs_f16(_Float16 a0,
_Float16 a1) {
+ return __nvvm_fmax_ftz_nan_xorsign_abs_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_ftz_nan_xorsign_abs_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.ftz.nan.xorsign.abs.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x
!cir.f16>
+// LLVM-LABEL: @test_fmax_ftz_nan_xorsign_abs_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmax.ftz.nan.xorsign.abs.f16x2(<2 x
half> %{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fmax_ftz_nan_xorsign_abs_f16x2(f16x2 a0,
f16x2 a1) {
+ return __nvvm_fmax_ftz_nan_xorsign_abs_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_ftz_xorsign_abs_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.ftz.xorsign.abs.f16" {{.*}} :
(!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fmax_ftz_xorsign_abs_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmax.ftz.xorsign.abs.f16(half %{{.*}},
half %{{.*}})
+extern "C" __device__ _Float16 test_fmax_ftz_xorsign_abs_f16(_Float16 a0,
_Float16 a1) {
+ return __nvvm_fmax_ftz_xorsign_abs_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_ftz_xorsign_abs_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.ftz.xorsign.abs.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x
!cir.f16>
+// LLVM-LABEL: @test_fmax_ftz_xorsign_abs_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmax.ftz.xorsign.abs.f16x2(<2 x
half> %{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fmax_ftz_xorsign_abs_f16x2(f16x2 a0, f16x2
a1) {
+ return __nvvm_fmax_ftz_xorsign_abs_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_nan_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.nan.f16" {{.*}} : (!cir.f16,
!cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fmax_nan_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmax.nan.f16(half %{{.*}}, half %{{.*}})
+extern "C" __device__ _Float16 test_fmax_nan_f16(_Float16 a0, _Float16 a1) {
+ return __nvvm_fmax_nan_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_nan_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.nan.f16x2" {{.*}} : (!cir.vector<2
x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_fmax_nan_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmax.nan.f16x2(<2 x half> %{{.*}},
<2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fmax_nan_f16x2(f16x2 a0, f16x2 a1) {
+ return __nvvm_fmax_nan_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_nan_xorsign_abs_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.nan.xorsign.abs.f16" {{.*}} :
(!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fmax_nan_xorsign_abs_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmax.nan.xorsign.abs.f16(half %{{.*}},
half %{{.*}})
+extern "C" __device__ _Float16 test_fmax_nan_xorsign_abs_f16(_Float16 a0,
_Float16 a1) {
+ return __nvvm_fmax_nan_xorsign_abs_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_nan_xorsign_abs_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.nan.xorsign.abs.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x
!cir.f16>
+// LLVM-LABEL: @test_fmax_nan_xorsign_abs_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmax.nan.xorsign.abs.f16x2(<2 x
half> %{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fmax_nan_xorsign_abs_f16x2(f16x2 a0, f16x2
a1) {
+ return __nvvm_fmax_nan_xorsign_abs_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_xorsign_abs_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.xorsign.abs.f16" {{.*}} :
(!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fmax_xorsign_abs_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmax.xorsign.abs.f16(half %{{.*}}, half
%{{.*}})
+extern "C" __device__ _Float16 test_fmax_xorsign_abs_f16(_Float16 a0, _Float16
a1) {
+ return __nvvm_fmax_xorsign_abs_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmax_xorsign_abs_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmax.xorsign.abs.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x
!cir.f16>
+// LLVM-LABEL: @test_fmax_xorsign_abs_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmax.xorsign.abs.f16x2(<2 x half>
%{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fmax_xorsign_abs_f16x2(f16x2 a0, f16x2 a1) {
+ return __nvvm_fmax_xorsign_abs_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.f16" {{.*}} : (!cir.f16, !cir.f16)
-> !cir.f16
+// LLVM-LABEL: @test_fmin_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmin.f16(half %{{.*}}, half %{{.*}})
+extern "C" __device__ _Float16 test_fmin_f16(_Float16 a0, _Float16 a1) {
+ return __nvvm_fmin_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.f16x2" {{.*}} : (!cir.vector<2 x
!cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_fmin_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmin.f16x2(<2 x half> %{{.*}}, <2 x
half> %{{.*}})
+extern "C" __device__ f16x2 test_fmin_f16x2(f16x2 a0, f16x2 a1) {
+ return __nvvm_fmin_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_ftz_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.ftz.f16" {{.*}} : (!cir.f16,
!cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fmin_ftz_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmin.ftz.f16(half %{{.*}}, half %{{.*}})
+extern "C" __device__ _Float16 test_fmin_ftz_f16(_Float16 a0, _Float16 a1) {
+ return __nvvm_fmin_ftz_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_ftz_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.ftz.f16x2" {{.*}} : (!cir.vector<2
x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_fmin_ftz_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmin.ftz.f16x2(<2 x half> %{{.*}},
<2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fmin_ftz_f16x2(f16x2 a0, f16x2 a1) {
+ return __nvvm_fmin_ftz_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_ftz_nan_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.ftz.nan.f16" {{.*}} : (!cir.f16,
!cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fmin_ftz_nan_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmin.ftz.nan.f16(half %{{.*}}, half
%{{.*}})
+extern "C" __device__ _Float16 test_fmin_ftz_nan_f16(_Float16 a0, _Float16 a1)
{
+ return __nvvm_fmin_ftz_nan_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_ftz_nan_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.ftz.nan.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x
!cir.f16>
+// LLVM-LABEL: @test_fmin_ftz_nan_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmin.ftz.nan.f16x2(<2 x half>
%{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fmin_ftz_nan_f16x2(f16x2 a0, f16x2 a1) {
+ return __nvvm_fmin_ftz_nan_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_ftz_nan_xorsign_abs_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.ftz.nan.xorsign.abs.f16" {{.*}} :
(!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fmin_ftz_nan_xorsign_abs_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmin.ftz.nan.xorsign.abs.f16(half %{{.*}},
half %{{.*}})
+extern "C" __device__ _Float16 test_fmin_ftz_nan_xorsign_abs_f16(_Float16 a0,
_Float16 a1) {
+ return __nvvm_fmin_ftz_nan_xorsign_abs_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_ftz_nan_xorsign_abs_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.ftz.nan.xorsign.abs.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x
!cir.f16>
+// LLVM-LABEL: @test_fmin_ftz_nan_xorsign_abs_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmin.ftz.nan.xorsign.abs.f16x2(<2 x
half> %{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fmin_ftz_nan_xorsign_abs_f16x2(f16x2 a0,
f16x2 a1) {
+ return __nvvm_fmin_ftz_nan_xorsign_abs_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_ftz_xorsign_abs_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.ftz.xorsign.abs.f16" {{.*}} :
(!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fmin_ftz_xorsign_abs_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmin.ftz.xorsign.abs.f16(half %{{.*}},
half %{{.*}})
+extern "C" __device__ _Float16 test_fmin_ftz_xorsign_abs_f16(_Float16 a0,
_Float16 a1) {
+ return __nvvm_fmin_ftz_xorsign_abs_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_ftz_xorsign_abs_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.ftz.xorsign.abs.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x
!cir.f16>
+// LLVM-LABEL: @test_fmin_ftz_xorsign_abs_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmin.ftz.xorsign.abs.f16x2(<2 x
half> %{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fmin_ftz_xorsign_abs_f16x2(f16x2 a0, f16x2
a1) {
+ return __nvvm_fmin_ftz_xorsign_abs_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_nan_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.nan.f16" {{.*}} : (!cir.f16,
!cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fmin_nan_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmin.nan.f16(half %{{.*}}, half %{{.*}})
+extern "C" __device__ _Float16 test_fmin_nan_f16(_Float16 a0, _Float16 a1) {
+ return __nvvm_fmin_nan_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_nan_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.nan.f16x2" {{.*}} : (!cir.vector<2
x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x !cir.f16>
+// LLVM-LABEL: @test_fmin_nan_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmin.nan.f16x2(<2 x half> %{{.*}},
<2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fmin_nan_f16x2(f16x2 a0, f16x2 a1) {
+ return __nvvm_fmin_nan_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_nan_xorsign_abs_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.nan.xorsign.abs.f16" {{.*}} :
(!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fmin_nan_xorsign_abs_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmin.nan.xorsign.abs.f16(half %{{.*}},
half %{{.*}})
+extern "C" __device__ _Float16 test_fmin_nan_xorsign_abs_f16(_Float16 a0,
_Float16 a1) {
+ return __nvvm_fmin_nan_xorsign_abs_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_nan_xorsign_abs_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.nan.xorsign.abs.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x
!cir.f16>
+// LLVM-LABEL: @test_fmin_nan_xorsign_abs_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmin.nan.xorsign.abs.f16x2(<2 x
half> %{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fmin_nan_xorsign_abs_f16x2(f16x2 a0, f16x2
a1) {
+ return __nvvm_fmin_nan_xorsign_abs_f16x2(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_xorsign_abs_f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.xorsign.abs.f16" {{.*}} :
(!cir.f16, !cir.f16) -> !cir.f16
+// LLVM-LABEL: @test_fmin_xorsign_abs_f16(
+// LLVM: call{{.*}} half @llvm.nvvm.fmin.xorsign.abs.f16(half %{{.*}}, half
%{{.*}})
+extern "C" __device__ _Float16 test_fmin_xorsign_abs_f16(_Float16 a0, _Float16
a1) {
+ return __nvvm_fmin_xorsign_abs_f16(a0, a1);
+}
+
+// CIR-LABEL: @test_fmin_xorsign_abs_f16x2
+// CIR: cir.call_llvm_intrinsic "nvvm.fmin.xorsign.abs.f16x2" {{.*}} :
(!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x
!cir.f16>
+// LLVM-LABEL: @test_fmin_xorsign_abs_f16x2(
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fmin.xorsign.abs.f16x2(<2 x half>
%{{.*}}, <2 x half> %{{.*}})
+extern "C" __device__ f16x2 test_fmin_xorsign_abs_f16x2(f16x2 a0, f16x2 a1) {
+ return __nvvm_fmin_xorsign_abs_f16x2(a0, a1);
+}
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
