https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/221859
>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 1/2] 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); +} >From 1c9b5c586f037061bf64ce3e38e6fb2202ed4dc8 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Tue, 8 Sep 2026 03:23:47 +0200 Subject: [PATCH 2/2] remove unneccessary comments --- clang/test/CIR/CodeGenCUDA/builtins-nvvm-fma-oob.cu | 7 ------- 1 file changed, 7 deletions(-) diff --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-fma-oob.cu b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-fma-oob.cu index 2385f6d851d39..de40e6b55d78e 100644 --- a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-fma-oob.cu +++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-fma-oob.cu @@ -17,13 +17,6 @@ // 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))); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
