[SYCL] Add support for -foffload-fp32-prec-div/sqrt options.

clang/include/clang/Basic/DiagnosticCommonKinds.td

@@ -379,6 +379,11 @@ def err_ppc_impossible_musttail: Error<
 def err_aix_musttail_unsupported: Error<
   "'musttail' attribute is not supported on AIX">;
 
+def warn_acuracy_conflicts_with_explicit_target_prec_option : Warning<
+  "floating point accuracy control '%0' conflicts with explicit target "
+  "precision option '%1'">,
+  InGroup<DiagGroup<"accuracy-conflicts-with-explicit-target-prec-option">>;
+
 // Source manager
 def err_cannot_open_file : Error<"cannot open file '%0': %1">, DefaultFatal;
 def err_file_modified : Error<

clang/include/clang/Basic/FPOptions.def

@@ -30,4 +30,6 @@ OPTION(BFloat16ExcessPrecision, LangOptions::ExcessPrecisionKind, 2, Float16Exce
 OPTION(FPAccuracy, LangOptions::FPAccuracyKind, 3, BFloat16ExcessPrecision)
 OPTION(MathErrno, bool, 1, FPAccuracy)
 OPTION(ComplexRange, LangOptions::ComplexRangeKind, 2, MathErrno)
+OPTION(OffloadFp32PrecDiv, bool, 1, ComplexRange)
+OPTION(OffloadFp32PrecSqrt, bool, 1, OffloadFp32PrecDiv)
 #undef OPTION

clang/include/clang/Basic/LangOptions.def

@@ -372,6 +372,8 @@ BENIGN_ENUM_LANGOPT(FPEvalMethod, FPEvalMethodKind, 2, FEM_UnsetOnCommandLine, "
 ENUM_LANGOPT(Float16ExcessPrecision, ExcessPrecisionKind, 2, FPP_Standard, "Intermediate truncation behavior for Float16 arithmetic")
 ENUM_LANGOPT(BFloat16ExcessPrecision, ExcessPrecisionKind, 2, FPP_Standard, "Intermediate truncation behavior for BFloat16 arithmetic")
 BENIGN_ENUM_LANGOPT(FPAccuracy, FPAccuracyKind, 3, FPA_Default, "Accuracy for floating point operations and library functions")
+LANGOPT(OffloadFp32PrecDiv, 1, 1, "Return correctly rounded results of fdiv")
+LANGOPT(OffloadFp32PrecSqrt, 1, 1, "Return correctly rounded results of sqrt")
 LANGOPT(NoBitFieldTypeAlign , 1, 0, "bit-field type alignment")
 LANGOPT(HexagonQdsp6Compat , 1, 0, "hexagon-qdsp6 backward compatibility")
 LANGOPT(ObjCAutoRefCount , 1, 0, "Objective-C automated reference counting")

clang/include/clang/Driver/Options.td

@@ -1157,6 +1157,22 @@ defm cx_fortran_rules: BoolOptionWithoutMarshalling<"f", "cx-fortran-rules",
   NegFlag<SetFalse, [], [ClangOption, CC1Option], "Range reduction is disabled "
   "for complex arithmetic operations">>;
 
+ defm offload_fp32_prec_div: BoolOption<"f", "offload-fp32-prec-div",
+  LangOpts<"OffloadFp32PrecDiv">, DefaultTrue,
+  PosFlag<SetTrue, [], [ClangOption, CC1Option], "fdiv operations in offload device "
+  "code are required to return correctly rounded results.">,
+  NegFlag<SetFalse, [], [ClangOption, CC1Option], "fdiv operations in offload device "
+  "code are not required to return correctly rounded results.">>,
+  Group<f_Group>;
+
+ defm offload_fp32_prec_sqrt: BoolOption<"f", "offload-fp32-prec-sqrt",
+ LangOpts<"OffloadFp32PrecSqrt">, DefaultTrue,
+  PosFlag<SetTrue, [], [ClangOption, CC1Option], "sqrt operations in offload device "
+  "code are required to return correctly rounded results.">,
+  NegFlag<SetFalse, [], [ClangOption, CC1Option], "sqrt operations in offload device "
+  "code are not required to return correctly rounded results.">>,
+  Group<f_Group>;
+
 // OpenCL-only Options
 def cl_opt_disable : Flag<["-"], "cl-opt-disable">, Group<opencl_Group>,
   Visibility<[ClangOption, CC1Option]>,

clang/lib/CodeGen/CGBuiltin.cpp

@@ -24099,6 +24099,7 @@ llvm::CallInst *CodeGenFunction::MaybeEmitFPBuiltinofFD(
             .Case("sincos", llvm::Intrinsic::fpbuiltin_sincos)
             .Case("exp10", llvm::Intrinsic::fpbuiltin_exp10)
             .Case("rsqrt", llvm::Intrinsic::fpbuiltin_rsqrt)
+            .Case("sqrt", llvm::Intrinsic::fpbuiltin_sqrt)
             .Default(0);
   } else {
     // The function has a clang builtin. Create an attribute for it
@@ -24200,7 +24201,8 @@ llvm::CallInst *CodeGenFunction::MaybeEmitFPBuiltinofFD(
   // a TU fp-accuracy requested.
   const LangOptions &LangOpts = getLangOpts();
   if (hasFuncNameRequestedFPAccuracy(Name, LangOpts) ||
-      !LangOpts.FPAccuracyVal.empty()) {
+      !LangOpts.FPAccuracyVal.empty() || !LangOpts.OffloadFp32PrecDiv ||
+      !LangOpts.OffloadFp32PrecSqrt) {
     llvm::Function *Func =
         CGM.getIntrinsic(FPAccuracyIntrinsicID, IRArgs[0]->getType());
     return CreateBuiltinCallWithAttr(*this, Name, Func, ArrayRef(IRArgs),

clang/lib/CodeGen/CGCall.cpp

@@ -1889,15 +1889,23 @@ void CodeGenModule::getDefaultFunctionFPAccuracyAttributes(
           Int32Ty, convertFPAccuracyToAspect(FuncMapIt->second)));
     }
   }
-  if (FuncAttrs.attrs().size() == 0)
+  if (FuncAttrs.attrs().size() == 0) {
+    StringRef FPAccuracyVal;
     if (!getLangOpts().FPAccuracyVal.empty()) {
-      StringRef FPAccuracyVal = llvm::fp::getAccuracyForFPBuiltin(
+      FPAccuracyVal = llvm::fp::getAccuracyForFPBuiltin(
           ID, FuncType, convertFPAccuracy(getLangOpts().FPAccuracyVal));
       assert(!FPAccuracyVal.empty() && "A valid accuracy value is expected");
       FuncAttrs.addAttribute("fpbuiltin-max-error", FPAccuracyVal);
       MD = llvm::ConstantAsMetadata::get(llvm::ConstantInt::get(
           Int32Ty, convertFPAccuracyToAspect(getLangOpts().FPAccuracyVal)));
     }
+    if (Name == "sqrt" && !getLangOpts().OffloadFp32PrecSqrt)
+      FPAccuracyVal = "3.0";
+    if (Name == "fdiv" && !getLangOpts().OffloadFp32PrecDiv)
+      FPAccuracyVal = "2.5";
+    if (!FPAccuracyVal.empty())
+      FuncAttrs.addAttribute("fpbuiltin-max-error", FPAccuracyVal);
+  }
 }
 
 /// Add denormal-fp-math and denormal-fp-math-f32 as appropriate for the
@@ -5790,10 +5798,16 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
   // Emit the actual call/invoke instruction.
   llvm::CallBase *CI;
   if (!InvokeDest) {
-    if (!getLangOpts().FPAccuracyFuncMap.empty() ||
-        !getLangOpts().FPAccuracyVal.empty()) {
-      const auto *FD = dyn_cast_if_present<FunctionDecl>(TargetDecl);
-      if (FD && FD->getNameInfo().getName().isIdentifier()) {
+    const auto *FD = dyn_cast_if_present<FunctionDecl>(TargetDecl);
+    if (FD && FD->getNameInfo().getName().isIdentifier()) {
+      StringRef FuncName = FD->getName();
+      const bool IsFloat32Type = FD->getReturnType()->isFloat32Type();
+      if (!getLangOpts().FPAccuracyFuncMap.empty() ||
+          !getLangOpts().FPAccuracyVal.empty() ||
+          (FuncName == "sqrt" && !getLangOpts().OffloadFp32PrecSqrt &&
+           IsFloat32Type) ||
+          (FuncName == "fdiv" && !getLangOpts().OffloadFp32PrecDiv &&
+           IsFloat32Type)) {
         CI = MaybeEmitFPBuiltinofFD(IRFuncTy, IRCallArgs, CalleePtr,
                                     FD->getName(), FD->getBuiltinID());
         if (CI)

clang/lib/Driver/ToolChains/Clang.cpp

@@ -1721,6 +1721,39 @@ static void CollectARMPACBTIOptions(const ToolChain &TC, const ArgList &Args,
   }
 }
 
+static void EmitAccuracyDiag(const Driver &D, const JobAction &JA,
+                             StringRef AccuracValStr, StringRef TargetPrecStr) {
+  if (JA.isDeviceOffloading(Action::OFK_SYCL)) {
+    D.Diag(clang::diag::warn_acuracy_conflicts_with_explicit_target_prec_option)
+        << AccuracValStr << TargetPrecStr;
+  }
+}
+
+void Clang::AddSPIRTargetArgs(const ArgList &Args, ArgStringList &CmdArgs,
+                              const JobAction &JA, const Driver &D) const {
+  if (JA.isDeviceOffloading(Action::OFK_SYCL)) {
+    if (Arg *A = Args.getLastArg(options::OPT_ffp_model_EQ)) {
+      if (!strcmp(A->getValue(), "fast")) {
+        CmdArgs.push_back("-fno-offload-fp32-prec-div");
+        CmdArgs.push_back("-fno-offload-fp32-prec-sqrt");
+      }
+    } else {
+      if (Arg *A = Args.getLastArg(options::OPT_ffp_accuracy_EQ)) {
+        if (Args.getLastArg(options::OPT_fno_offload_fp32_prec_div))
+          EmitAccuracyDiag(D, JA, A->getValue(), "-fno-offload-fp32-prec-div");
+        if (Args.getLastArg(options::OPT_fno_offload_fp32_prec_sqrt))
+          EmitAccuracyDiag(D, JA, A->getValue(), "-fno-offload-fp32-prec-sqrt");
+      }
+      if (!Args.hasFlag(options::OPT_foffload_fp32_prec_div,
+                        options::OPT_fno_offload_fp32_prec_div, true))
+        CmdArgs.push_back("-fno-offload-fp32-prec-div");
+      if (!Args.hasFlag(options::OPT_foffload_fp32_prec_sqrt,
+                        options::OPT_fno_offload_fp32_prec_sqrt, true))
+        CmdArgs.push_back("-fno-offload-fp32-prec-sqrt");
+    }
+  }
+}
+
 void Clang::AddARMTargetArgs(const llvm::Triple &Triple, const ArgList &Args,
                              ArgStringList &CmdArgs, bool KernelOrKext) const {
   RenderARMABI(getToolChain().getDriver(), Triple, Args, CmdArgs);
@@ -1771,8 +1804,10 @@ void Clang::AddARMTargetArgs(const llvm::Triple &Triple, const ArgList &Args,
 
 void Clang::RenderTargetOptions(const llvm::Triple &EffectiveTriple,
                                 const ArgList &Args, bool KernelOrKext,
-                                ArgStringList &CmdArgs) const {
+                                ArgStringList &CmdArgs,
+                                const JobAction &JA) const {
   const ToolChain &TC = getToolChain();
+  const Driver &D = TC.getDriver();
 
   // Add the target features
   getTargetFeatures(TC.getDriver(), EffectiveTriple, Args, CmdArgs, false);
@@ -1782,6 +1817,12 @@ void Clang::RenderTargetOptions(const llvm::Triple &EffectiveTriple,
   default:
     break;
 
+  case llvm::Triple::spir:
+  case llvm::Triple::spir64:
+  case llvm::Triple::spirv32:
+  case llvm::Triple::spirv64:
+    AddSPIRTargetArgs(Args, CmdArgs, JA, D);
+    break;
   case llvm::Triple::arm:
   case llvm::Triple::armeb:
   case llvm::Triple::thumb:
@@ -6806,7 +6847,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     CmdArgs.push_back(Args.MakeArgString(CPU));
   }
 
-  RenderTargetOptions(Triple, Args, KernelOrKext, CmdArgs);
+  RenderTargetOptions(Triple, Args, KernelOrKext, CmdArgs, JA);
 
   // Add clang-cl arguments.
   types::ID InputType = Input.getType();

clang/lib/Driver/ToolChains/Clang.h

@@ -47,14 +47,18 @@ class LLVM_LIBRARY_VISIBILITY Clang : public Tool {
 
   void RenderTargetOptions(const llvm::Triple &EffectiveTriple,
                            const llvm::opt::ArgList &Args, bool KernelOrKext,
-                           llvm::opt::ArgStringList &CmdArgs) const;
+                           llvm::opt::ArgStringList &CmdArgs,
+                           const JobAction &JA) const;
 
   void AddAArch64TargetArgs(const llvm::opt::ArgList &Args,
                             llvm::opt::ArgStringList &CmdArgs) const;
   void AddARMTargetArgs(const llvm::Triple &Triple,
                         const llvm::opt::ArgList &Args,
                         llvm::opt::ArgStringList &CmdArgs,
                         bool KernelOrKext) const;
+  void AddSPIRTargetArgs(const llvm::opt::ArgList &Args,
+                         llvm::opt::ArgStringList &CmdArgs, const JobAction &JA,
+                         const Driver &D) const;
   void AddARM64TargetArgs(const llvm::opt::ArgList &Args,
                           llvm::opt::ArgStringList &CmdArgs) const;
   void AddLoongArchTargetArgs(const llvm::opt::ArgList &Args,

clang/test/CodeGenSYCL/offload-fp32-div-sqrt.cpp

@@ -0,0 +1,156 @@
+// DEFINE: %{common_opts_spirv32} = -internal-isystem %S/Inputs \
+// DEFINE: -fsycl-is-device -emit-llvm -triple spirv32-unknown-unknown
+
+// DEFINE: %{common_opts_spirv64} = -internal-isystem %S/Inputs \
+// DEFINE: -fsycl-is-device -emit-llvm -triple spirv32-unknown-unknown
+
+// DEFINE: %{common_opts_spir} = -internal-isystem %S/Inputs \
+// DEFINE: -fsycl-is-device -emit-llvm -triple spirv32-unknown-unknown
+
+// DEFINE: %{common_opts_spir64} = -internal-isystem %S/Inputs \
+// DEFINE: -fsycl-is-device -emit-llvm -triple spirv32-unknown-unknown
+
+// RUN: %clang_cc1 %{common_opts_spirv32} %s -o - \
+// RUN: | FileCheck --check-prefix PREC-SQRT %s
+
+// RUN: %clang_cc1 %{common_opts_spirv32} -foffload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix PREC-SQRT %s
+
+// RUN: %clang_cc1 %{common_opts_spirv32} -fno-offload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-SQRT %s
+
+// RUN: %clang_cc1 %{common_opts_spirv32} -foffload-fp32-prec-div %s -o - \
+// RUN: | FileCheck --check-prefix PREC-DIV %s
+
+// RUN: %clang_cc1 %{common_opts_spirv32} -fno-offload-fp32-prec-div %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-DIV %s
+
+// RUN: %clang_cc1 %{common_opts_spirv32} -ffast-math \
+// RUN:-fno-offload-fp32-prec-div -fno-offload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-SQRT-FAST %s
+
+// RUN: %clang_cc1 %{common_opts_spirv32} -ffast-math \
+// RUN: -fno-offload-fp32-prec-div -fno-offload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-DIV-FAST %s
+
+//
+
+// RUN: %clang_cc1 %{common_opts_spirv64} %s -o - \
+// RUN: | FileCheck --check-prefix PREC-SQRT %s
+
+// RUN: %clang_cc1 %{common_opts_spirv64} -foffload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix PREC-SQRT %s
+
+// RUN: %clang_cc1 %{common_opts_spirv64} -fno-offload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-SQRT %s
+
+// RUN: %clang_cc1 %{common_opts_spirv64} -foffload-fp32-prec-div %s -o - \
+// RUN: | FileCheck --check-prefix PREC-DIV %s
+
+// RUN: %clang_cc1 %{common_opts_spirv64} -fno-offload-fp32-prec-div %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-DIV %s
+
+// RUN: %clang_cc1 %{common_opts_spirv64} -ffast-math -fno-offload-fp32-prec-div \
+// RUN: -fno-offload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-SQRT-FAST %s
+
+// RUN: %clang_cc1 %{common_opts_spirv64} -ffast-math -fno-offload-fp32-prec-div \
+// RUN: -fno-offload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-DIV-FAST %s
+
+//
+
+// RUN: %clang_cc1 %{common_opts_spir} %s -o - \
+// RUN: | FileCheck --check-prefix PREC-SQRT %s
+
+// RUN: %clang_cc1 %{common_opts_spir} -foffload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix PREC-SQRT %s
+
+// RUN: %clang_cc1 %{common_opts_spir} -fno-offload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-SQRT %s
+
+// RUN: %clang_cc1 %{common_opts_spir} -foffload-fp32-prec-div %s -o - \
+// RUN: | FileCheck --check-prefix PREC-DIV %s
+
+// RUN: %clang_cc1 %{common_opts_spir} -fno-offload-fp32-prec-div %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-DIV %s
+
+// RUN: %clang_cc1 %{common_opts_spir} -ffast-math -fno-offload-fp32-prec-div \
+// RUN: -fno-offload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-SQRT-FAST %s
+
+// RUN: %clang_cc1 %{common_opts_spir} -ffast-math -fno-offload-fp32-prec-div \
+// RUN: -fno-offload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-DIV-FAST %s
+
+//
+
+// RUN: %clang_cc1 %{common_opts_spir64} %s -o - \
+// RUN: | FileCheck --check-prefix PREC-SQRT %s
+
+// RUN: %clang_cc1 %{common_opts_spir64} -foffload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix PREC-SQRT %s
+
+// RUN: %clang_cc1 %{common_opts_spir64} -fno-offload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-SQRT %s
+
+// RUN: %clang_cc1 %{common_opts_spir64} -foffload-fp32-prec-div %s -o - \
+// RUN: | FileCheck --check-prefix PREC-DIV %s
+
+// RUN: %clang_cc1 %{common_opts_spir64} -fno-offload-fp32-prec-div %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-DIV %s
+
+// RUN: %clang_cc1 %{common_opts_spir64} -ffast-math -fno-offload-fp32-prec-div \
+// RUN: -fno-offload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-SQRT-FAST %s
+
+// RUN: %clang_cc1 %{common_opts_spir64} -ffast-math -fno-offload-fp32-prec-div \
+// RUN: -fno-offload-fp32-prec-sqrt %s -o - \
+// RUN: | FileCheck --check-prefix ROUNDED-DIV-FAST %s
+
+#include "sycl.hpp"
+
+extern "C" SYCL_EXTERNAL float sqrt(float);
+extern "C" SYCL_EXTERNAL float fdiv(float, float);
+
+using namespace sycl;
+
+int main() {
+  const unsigned array_size = 4;
+  range<1> numOfItems{array_size};
+  float Value1 = .5f;
+  float Value2 = .9f;
+  queue deviceQueue;
+
+  deviceQueue.submit([&](handler& cgh) {
+    cgh.parallel_for<class KernelSqrt>(numOfItems,
+    [=](id<1> wiID) {
+      // PREC-SQRT: call spir_func float @sqrt(float noundef {{.*}})
+      // ROUNDED-SQRT: call float @llvm.fpbuiltin.sqrt.f32(float {{.*}}) #[[ATTR_SQRT:[0-9]+]]
+      // ROUNDED-DIV:  call spir_func float @sqrt(float noundef {{.*}})
+      (void)sqrt(Value1);
+    });
+  });
+
+  deviceQueue.submit([&](handler& cgh) {
+    cgh.parallel_for<class KernelFdiv>(numOfItems,
+    [=](id<1> wiID) {
+      // PREC-SQRT: call spir_func float @fdiv(float noundef {{.*}}, float noundef {{.*}})
+      // ROUNDED-SQRT: call spir_func float @fdiv(float noundef {{.*}}, float noundef {{.*}})
+
+      // ROUNDED-SQRT-FAST: call reassoc nnan ninf nsz arcp afn float @llvm.fpbuiltin.sqrt.f32(float {{.*}}) #[[ATTR_SQRT:[0-9]+]]
+
+      // PREC-DIV: call spir_func float @fdiv(float noundef {{.*}}, float noundef {{.*}})
+      // ROUNDED-DIV: call float @llvm.fpbuiltin.fdiv.f32(float {{.*}}, float {{.*}}) #[[ATTR_DIV:[0-9]+]]
+      // ROUNDED-DIV-FAST: call reassoc nnan ninf nsz arcp afn float @llvm.fpbuiltin.fdiv.f32(float {{.*}}, float {{.*}}) #[[ATTR_DIV:[0-9]+]]
+      (void)fdiv(Value1, Value1);
+    });
+  });
+
+return 0;
+}
+
+// ROUNDED-SQRT: attributes #[[ATTR_SQRT]] = {{.*}}"fpbuiltin-max-error"="3.0"
+// ROUNDED-SQRT-FAST: attributes #[[ATTR_SQRT]] = {{.*}}"fpbuiltin-max-error"="3.0"
+// ROUNDED-DIV: attributes #[[ATTR_DIV]] = {{.*}}"fpbuiltin-max-error"="2.5"
+// ROUNDED-DIV-FAST: attributes #[[ATTR_DIV]] = {{.*}}"fpbuiltin-max-error"="2.5"