[SYCL] Pass foffload-fp32-prec-[div/sqrt] options to device's BE

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -2061,9 +2061,18 @@ void SYCLToolChain::AddImpliedTargetArgs(const llvm::Triple &Triple,
     if (Args.hasFlag(options::OPT_ftarget_export_symbols,
                      options::OPT_fno_target_export_symbols, false))
       BeArgs.push_back("-library-compilation");
-  } else if (IsJIT)
+    // -foffload-fp32-prec-[sqrt/div]
+    if (Args.hasArg(options::OPT_foffload_fp32_prec_div) ||
+        Args.hasArg(options::OPT_foffload_fp32_prec_sqrt))
+      BeArgs.push_back("-ze-fp32-correctly-rounded-divide-sqrt");
+  } else if (IsJIT) {
     // -ftarget-compile-fast JIT
     Args.AddLastArg(BeArgs, options::OPT_ftarget_compile_fast);
+    // -foffload-fp32-prec-div JIT
+    Args.AddLastArg(BeArgs, options::OPT_foffload_fp32_prec_div);
+    // -foffload-fp32-prec-sqrt JIT
+    Args.AddLastArg(BeArgs, options::OPT_foffload_fp32_prec_sqrt);
+  }
   if (IsGen) {
     for (auto [DeviceName, BackendArgStr] : PerDeviceArgs) {
       CmdArgs.push_back("-device_options");

clang/test/Driver/sycl-foffload-fp32-prec-div-old-model.cpp

@@ -0,0 +1,13 @@
+// Test SYCL -foffload-fp32-prec-div
+
+// RUN: %clang -### -fsycl --no-offload-new-driver \
+// RUN:    -fsycl-targets=spir64_gen -foffload-fp32-prec-div %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=AOT %s
+
+// RUN: %clang -### -fsycl --no-offload-new-driver \
+// RUN:   -foffload-fp32-prec-div %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=JIT %s
+
+// AOT: "-ze-fp32-correctly-rounded-divide-sqrt"
+
+// JIT: clang-offload-wrapper{{.*}} "-compile-opts={{.*}}-foffload-fp32-prec-div"

clang/test/Driver/sycl-foffload-fp32-prec-div.cpp

@@ -0,0 +1,13 @@
+// Test SYCL -foffload-fp32-prec-div
+
+// RUN: %clang -### -fsycl --offload-new-driver \
+// RUN:    -fsycl-targets=spir64_gen -foffload-fp32-prec-div %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=AOT %s
+
+// RUN: %clang -### -fsycl --offload-new-driver \
+// RUN:   -foffload-fp32-prec-div %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=JIT %s
+
+// AOT: clang-offload-packager{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch={{.*}},kind=sycl,compile-opts=-options -ze-fp32-correctly-rounded-divide-sqrt{{.*}}"
+
+// JIT: clang-offload-packager{{.*}} "--image=file={{.*}}.bc,triple=spir64-unknown-unknown,arch={{.*}}compile-opts={{.*}}-foffload-fp32-prec-div"

clang/test/Driver/sycl-foffload-fp32-prec-sqrt-old-model.cpp

@@ -0,0 +1,13 @@
+// Test SYCL -foffload-fp32-prec-sqrt
+
+// RUN: %clang -### -fsycl --no-offload-new-driver \
+// RUN:    -fsycl-targets=spir64_gen -foffload-fp32-prec-sqrt %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=AOT %s
+
+// RUN: %clang -### -fsycl --no-offload-new-driver \
+// RUN:   -foffload-fp32-prec-sqrt %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=JIT %s
+
+// AOT: "-ze-fp32-correctly-rounded-divide-sqrt"
+
+// JIT: clang-offload-wrapper{{.*}} "-compile-opts={{.*}}-foffload-fp32-prec-sqrt"

clang/test/Driver/sycl-foffload-fp32-prec-sqrt.cpp

@@ -0,0 +1,13 @@
+// Test SYCL -foffload-fp32-prec-sqrt
+
+// RUN: %clang -### -fsycl --offload-new-driver \
+// RUN:    -fsycl-targets=spir64_gen -foffload-fp32-prec-sqrt %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=AOT %s
+
+// RUN: %clang -### -fsycl --offload-new-driver \
+// RUN:   -foffload-fp32-prec-sqrt %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=JIT %s
+
+// AOT: clang-offload-packager{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch={{.*}},kind=sycl,compile-opts=-options -ze-fp32-correctly-rounded-divide-sqrt{{.*}}"
+
+// JIT: clang-offload-packager{{.*}} "--image=file={{.*}}.bc,triple=spir64-unknown-unknown,arch={{.*}}compile-opts={{.*}}-foffload-fp32-prec-sqrt"

sycl/source/detail/program_manager/program_manager.cpp

@@ -446,6 +446,16 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts,
       CompileOpts = NewCompileOpts;
       OptPos = CompileOpts.find(TargetRegisterAllocMode);
     }
+    constexpr std::string_view ReplaceOpts[] = {"-foffload-fp32-prec-div",
+                                                "-foffload-fp32-prec-sqrt"};
+    for (const std::string_view Opt : ReplaceOpts) {
+      if (auto Pos = CompileOpts.find(Opt); Pos != std::string::npos) {
+        const char *BackendOption = nullptr;
+        PlatformImpl->getBackendOption(std::string(Opt).c_str(),
+                                       &BackendOption);
+        CompileOpts.replace(Pos, Opt.length(), BackendOption);
+      }
+    }
   }
 }
 

sycl/test-e2e/KernelAndProgram/fp32-precise-fdiv.cpp

@@ -0,0 +1,46 @@
+// RUN: %{build} -Wno-error=unused-command-line-argument -foffload-fp32-prec-div -foffload-fp32-prec-sqrt -o %t_with.out
+// RUN: %{build} -Wno-error=unused-command-line-argument -foffload-fp32-prec-div -o %t_with_div.out
+// RUN: %{build} -Wno-error=unused-command-line-argument -foffload-fp32-prec-sqrt -o %t_with_sqrt.out
+// RUN: %{build} -o %t_without.out
+
+// RUN: env SYCL_UR_TRACE=2 %{run} %t_with.out 2>&1 | FileCheck %if hip || cuda %{ --check-prefix=CHECK-WITHOUT %} %else %{ --check-prefix=CHECK-WITH %} %s
+// RUN: env SYCL_UR_TRACE=2 %{run} %t_with_div.out 2>&1 | FileCheck %if hip || cuda %{ --check-prefix=CHECK-WITHOUT %} %else %{ --check-prefix=CHECK-WITH %} %s
+// RUN: env SYCL_UR_TRACE=2 %{run} %t_with_sqrt.out 2>&1 | FileCheck %if hip || cuda %{ --check-prefix=CHECK-WITHOUT %} %else %{ --check-prefix=CHECK-WITH %} %s
+// RUN: env SYCL_UR_TRACE=2 %{run} %t_without.out 2>&1 | FileCheck --implicit-check-not=fp32-correctly-rounded-divide-sqrt %s
+
+// CHECK-WITH: <--- urProgramBuild
+// CHECK-WITH-SAME: fp32-correctly-rounded-divide-sqrt
+
+// CHECK-WITHOUT-NOT: <--- urProgramBuild{{.*}}fp32-correctly-rounded-divide-sqrt{{.*}} -> UR_RESULT_SUCCESS
+// CHECK-WITHOUT: <--- urProgramBuild{{.*}} -> UR_RESULT_SUCCESS
+
+#include <sycl/detail/core.hpp>
+
+int main() {
+  sycl::buffer<size_t, 1> Buffer(4);
+
+  sycl::queue Queue;
+
+  sycl::range<1> NumOfWorkItems{Buffer.size()};
+
+  Queue.submit([&](sycl::handler &cgh) {
+    sycl::accessor Accessor{Buffer, cgh, sycl::write_only};
+    cgh.parallel_for<class FillBuffer>(NumOfWorkItems, [=](sycl::id<1> WIid) {
+      Accessor[WIid] = WIid.get(0);
+    });
+  });
+
+  sycl::host_accessor HostAccessor{Buffer, sycl::read_only};
+
+  bool MismatchFound = false;
+  for (size_t I = 0; I < Buffer.size(); ++I) {
+    if (HostAccessor[I] != I) {
+      std::cout << "The result is incorrect for element: " << I
+                << " , expected: " << I << " , got: " << HostAccessor[I]
+                << std::endl;
+      MismatchFound = true;
+    }
+  }
+
+  return MismatchFound;
+}