[SYCL][Driver] Force precise division rounding for precise model (#12107)

This commit makes the driver pass
`-cl-fp32-correctly-rounded-divide-sqrt` as backend compiler options for
SPIR-V targets when the `-fpp-model=precise` option is used.

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -1246,6 +1246,11 @@ void SYCLToolChain::AddImpliedTargetArgs(const llvm::Triple &Triple,
     if (Arg *A = Args.getLastArg(options::OPT_O_Group))
       if (A->getOption().matches(options::OPT_O0))
         BeArgs.push_back("-cl-opt-disable");
+  // In precise floating-point mode we pass the OpenCL flag forcing division to
+  // be correctly rounded.
+  if (Arg *A = Args.getLastArg(options::OPT_ffp_model_EQ))
+    if (StringRef{A->getValue()}.equals("precise"))
+      BeArgs.push_back("-cl-fp32-correctly-rounded-divide-sqrt");
   StringRef RegAllocModeOptName = "-ftarget-register-alloc-mode=";
   if (Arg *A = Args.getLastArg(options::OPT_ftarget_register_alloc_mode_EQ)) {
     StringRef RegAllocModeVal = A->getValue(0);

clang/test/Driver/sycl-offload-aot.cpp

@@ -220,6 +220,18 @@
 // RUN:   | FileCheck -check-prefix=CHK-TOOLS-IMPLIED-OPTS-GEN %s
 // CHK-TOOLS-IMPLIED-OPTS-GEN: ocloc{{.*}} "-options" "-g -cl-opt-disable" "-DFOO1" "-DFOO2"
 
+// RUN:   %clang -### -target x86_64-unknown-linux-gnu -fsycl -fsycl-targets=spir64_fpga-unknown-unknown -ffp-model=precise -Xsycl-target-backend "-DFOO1 -DFOO2" %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-TOOLS-IMPLIED-ROUNDING-FPGA %s
+// CHK-TOOLS-IMPLIED-ROUNDING-FPGA: opencl-aot{{.*}} "--bo=-cl-fp32-correctly-rounded-divide-sqrt" "-DFOO1" "-DFOO2"
+
+// RUN:   %clang -### -target x86_64-unknown-linux-gnu -fsycl -fsycl-targets=spir64_x86_64-unknown-unknown -ffp-model=precise -Xsycl-target-backend "-DFOO1 -DFOO2" %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-TOOLS-IMPLIED-ROUNDING-CPU %s
+// CHK-TOOLS-IMPLIED-ROUNDING-CPU: opencl-aot{{.*}} "--bo=-cl-fp32-correctly-rounded-divide-sqrt" "-DFOO1" "-DFOO2"
+
+// RUN:   %clang -### -target x86_64-unknown-linux-gnu -fsycl -fsycl-targets=spir64_gen-unknown-unknown -ffp-model=precise -Xsycl-target-backend "-DFOO1 -DFOO2" %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-TOOLS-IMPLIED-ROUNDING-GEN %s
+// CHK-TOOLS-IMPLIED-ROUNDING-GEN: ocloc{{.*}} "-options" "-cl-fp32-correctly-rounded-divide-sqrt" "-DFOO1" "-DFOO2"
+
 /// Check -Xsycl-target-linker option passing
 // RUN:   %clang -### -target x86_64-unknown-linux-gnu -fsycl -fsycl-targets=spir64_fpga-unknown-unknown -Xshardware -Xsycl-target-linker "-DFOO1 -DFOO2" %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=CHK-TOOLS-FPGA-OPTS2 %s

clang/test/Driver/sycl-offload.c

@@ -424,6 +424,11 @@
 // CHK-TOOLS-IMPLIED-OPTS-O0-NOT: clang-offload-wrapper{{.*}} "-compile-opts={{.*}}-cl-opt-disable"
 // CHK-TOOLS-IMPLIED-OPTS-O2-NOT: clang-offload-wrapper{{.*}} "-compile-opts={{.*}}-cl-opt-disable"
 
+/// Check for implied options (-ffp-model=precise)
+// RUN:   %clang -### -target x86_64-unknown-linux-gnu -fsycl -fsycl-targets=spir64 -ffp-model=precise %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-TOOLS-IMPLIED-ROUNDING %s
+// CHK-TOOLS-IMPLIED-ROUNDING: clang-offload-wrapper{{.*}} "-compile-opts={{.*}}-cl-fp32-correctly-rounded-divide-sqrt
+
 // RUN:   %clang -### -target x86_64-unknown-linux-gnu -fsycl -fsycl-targets=spir64-unknown-unknown -Xsycl-target-linker "-DFOO1 -DFOO2" %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=CHK-TOOLS-OPTS2 %s
 // CHK-TOOLS-OPTS2: clang-offload-wrapper{{.*}} "-link-opts=-DFOO1 -DFOO2"

sycl/test-e2e/Basic/float_division_precise.cpp

@@ -0,0 +1,39 @@
+// RUN: %{build} -ffp-model=precise -o %t.out
+// RUN: %{run} %t.out
+
+// Tests that -ffp-model=precise causes floating point division to be the same
+// on device and host.
+
+#include <sycl.hpp>
+
+constexpr size_t NumElems = 1024;
+
+int main() {
+  sycl::queue Q;
+  float *InData = sycl::malloc_shared<float>(NumElems, Q);
+  float *OutData = sycl::malloc_shared<float>(NumElems, Q);
+
+  for (size_t I = 0; I < NumElems; ++I) {
+    InData[I] = float(I) + 1.0f;
+    OutData[I] = 0.0f;
+  }
+
+  Q.parallel_for(sycl::range<1>(NumElems), [=](sycl::id<1> Idx) {
+     OutData[Idx] = InData[Idx] / InData[NumElems - Idx - 1];
+   }).wait_and_throw();
+
+  size_t NumFails = 0;
+  for (size_t I = 0; I < NumElems; ++I) {
+    float Expected = InData[I] / InData[NumElems - I - 1];
+    if (OutData[I] != Expected) {
+      std::cout << "Unexpected result for element " << I << ": " << OutData[I]
+                << " != " << Expected << std::endl;
+      ++NumFails;
+    }
+  }
+
+  sycl::free(InData, Q);
+  sycl::free(OutData, Q);
+
+  return NumFails;
+}