[Driver][SYCL] Deprecate -fsycl-fp32-prec-sqrt (#17257)

mdtoguchi · web-flow · commit 44632006cfd8 · 2025-03-19T07:14:00.000-07:00
We have introduced a new set of options -foffload-fp32-prec-div and
-foffload-fp32-prec-sqrt. The -foffload-fp32-prec-sqrt option provides
the same functionality as -fsycl-fp32-prec-sqrt. Deprecate
-fsycl-fp32-prec-sqrt, making it an alias for -foffload-fp32-prec-sqrt.

Use of -fsycl-fp32-prec-sqrt will also emit a deprecation message:

option '-fsycl-fp32-prec-sqrt' is deprecated and will be removed in a
future release, use '-foffload-fp32-prec-sqrt' instead

Upon updating this behavior, an existing deprecated option
-fsycl-use-bitcode is also impacted and will emit the new diagnostic.

Update the UserManual to reflect the replacement option for both
-fsycl-fp32-prec-sqrt and -fsycl-use-bitcode.
diff --git a/clang/include/clang/Basic/CodeGenOptions.def b/clang/include/clang/Basic/CodeGenOptions.def
@@ -197,7 +197,6 @@ CODEGENOPT(OpenCLCorrectlyRoundedDivSqrt, 1, 0) ///< -cl-fp32-correctly-rounded-
 CODEGENOPT(HIPCorrectlyRoundedDivSqrt, 1, 1) ///< -fno-hip-fp32-correctly-rounded-divide-sqrt
 CODEGENOPT(DisableBlockSignatureString, 1, 0) ///< Set when -fdisable-block-signature-string is enabled.
 CODEGENOPT(HIPSaveKernelArgName, 1, 0) ///< Set when -fhip-kernel-arg-name is enabled.
-CODEGENOPT(SYCLFp32PrecSqrt, 1, 0) ///< -fsycl-fp32-prec-sqrt
 CODEGENOPT(UniqueInternalLinkageNames, 1, 0) ///< Internal Linkage symbols get unique names.
 CODEGENOPT(SplitMachineFunctions, 1, 0) ///< Split machine functions using profile information.
 CODEGENOPT(PPCUseFullRegisterNames, 1, 0) ///< Print full register names in assembly
diff --git a/clang/include/clang/Basic/DiagnosticDriverKinds.td b/clang/include/clang/Basic/DiagnosticDriverKinds.td
@@ -566,7 +566,7 @@ def warn_drv_object_size_disabled_O0 : Warning<
 def warn_drv_deprecated_option : Warning<
   "option '%0' is deprecated, use '%1' directly instead">, InGroup<Deprecated>;
 def warn_drv_deprecated_option_release : Warning<
-  "option '%0' is deprecated and will be removed in a future release">,
+  "option '%0' is deprecated and will be removed in a future release%select{|, use '%2' instead}1">,
   InGroup<Deprecated>;
 def warn_drv_deprecated_argument_option_release : Warning<
   "argument '%0' for option '%1' is deprecated and will be removed in a "
diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td
@@ -6943,9 +6943,9 @@ def fsycl_device_lib_jit_link : Flag<["-"], "fsycl-device-lib-jit-link">,
 def fno_sycl_device_lib_jit_link : Flag<["-"], "fno-sycl-device-lib-jit-link">,
   HelpText<"Disables sycl device library jit link (experimental)">;
 def fsycl_fp32_prec_sqrt : Flag<["-"], "fsycl-fp32-prec-sqrt">,
+  Alias<foffload_fp32_prec_sqrt>, Flags<[Deprecated]>,
   Visibility<[ClangOption, CC1Option]>, HelpText<"SYCL only. Specify that "
-  "single precision floating-point sqrt is correctly rounded.">,
-  MarshallingInfoFlag<CodeGenOpts<"SYCLFp32PrecSqrt">>;
+  "single precision floating-point sqrt is correctly rounded. (deprecated)">;
 def fsycl_default_sub_group_size
     : Separate<["-"], "fsycl-default-sub-group-size">,
       HelpText<"Set the default sub group size for SYCL kernels">,
diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp
@@ -360,7 +360,16 @@ InputArgList Driver::ParseArgStrings(ArrayRef<const char *> ArgStrings,
     while (Used->getAlias())
       Used = Used->getAlias();
     if (Used->getOption().hasFlag(options::Deprecated)) {
-      Diag(diag::warn_drv_deprecated_option_release) << Used->getAsString(Args);
+      // Some deprecated options have a replacement option.  In these cases,
+      // add the replacement option string to the diagnostic.
+      SmallString<128> AliasOpt;
+      if (Used != A) {
+        AliasOpt = A->getSpelling();
+        if (A->getNumValues())
+          AliasOpt += A->getValue();
+      }
+      Diag(diag::warn_drv_deprecated_option_release)
+          << Used->getAsString(Args) << !AliasOpt.empty() << AliasOpt;
       ContainsError |= Diags.getDiagnosticLevel(
                            diag::warn_drv_deprecated_option_release,
                            SourceLocation()) > DiagnosticsEngine::Warning;
diff --git a/clang/lib/Driver/ToolChains/AMDGPU.cpp b/clang/lib/Driver/ToolChains/AMDGPU.cpp
@@ -1079,11 +1079,10 @@ ROCMToolChain::getCommonDeviceLibNames(
   bool FastRelaxedMath = DriverArgs.hasFlag(options::OPT_ffast_math,
                                             options::OPT_fno_fast_math, false);
   bool CorrectSqrt = false;
-  if (DeviceOffloadingKind == Action::OFK_SYCL) {
+  if (DeviceOffloadingKind == Action::OFK_SYCL)
     // When using SYCL, sqrt is only correctly rounded if the flag is specified
-    CorrectSqrt = DriverArgs.hasArg(options::OPT_fsycl_fp32_prec_sqrt) ||
-                  DriverArgs.hasArg(options::OPT_foffload_fp32_prec_sqrt);
-  } else
+    CorrectSqrt = DriverArgs.hasArg(options::OPT_foffload_fp32_prec_sqrt);
+  else
     CorrectSqrt = DriverArgs.hasFlag(
         options::OPT_fhip_fp32_correctly_rounded_divide_sqrt,
         options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt, true);
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -3144,17 +3144,18 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
     CmdArgs.push_back(A->getValue());
   }
 
-  auto addSPIRVArgs = [&](StringRef SPIRVArg) {
+  auto addSPIRVArgs = [&](const Arg *PrecArg) {
     if (IsFp32PrecDivSqrtAllowed) {
+      OptSpecifier Opt(PrecArg->getOption().getID());
       if (!FPAccuracy.empty())
-        EmitAccuracyDiag(D, JA, FPAccuracy, SPIRVArg);
-      if (SPIRVArg == "-fno-offload-fp32-prec-div")
+        EmitAccuracyDiag(D, JA, FPAccuracy, PrecArg->getSpelling());
+      if (Opt == options::OPT_fno_offload_fp32_prec_div)
         NoOffloadFP32PrecDiv = true;
-      else if (SPIRVArg == "-fno-offload-fp32-prec-sqrt")
+      else if (Opt == options::OPT_fno_offload_fp32_prec_sqrt)
         NoOffloadFP32PrecSqrt = true;
-      else if (SPIRVArg == "-foffload-fp32-prec-sqrt")
+      else if (Opt == options::OPT_foffload_fp32_prec_sqrt)
         NoOffloadFP32PrecSqrt = false;
-      else if (SPIRVArg == "-foffload-fp32-prec-div")
+      else if (Opt == options::OPT_foffload_fp32_prec_div)
         NoOffloadFP32PrecDiv = false;
     }
   };
@@ -3181,17 +3182,15 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
     default: continue;
 
     case options::OPT_foffload_fp32_prec_div:
-      addSPIRVArgs("-foffload-fp32-prec-div");
-      break;
     case options::OPT_foffload_fp32_prec_sqrt:
-      addSPIRVArgs("-foffload-fp32-prec-sqrt");
-      break;
     case options::OPT_fno_offload_fp32_prec_div:
-      addSPIRVArgs("-fno-offload-fp32-prec-div");
-      break;
     case options::OPT_fno_offload_fp32_prec_sqrt:
-      addSPIRVArgs("-fno-offload-fp32-prec-sqrt");
-      break;
+      if (IsFp32PrecDivSqrtAllowed) {
+        addSPIRVArgs(A);
+        break;
+      }
+      // Skip claim, as we didn't use the option.
+      continue;
     case options::OPT_fcx_limited_range:
       if (GccRangeComplexOption.empty()) {
         if (Range != LangOptions::ComplexRangeKind::CX_Basic)
diff --git a/clang/lib/Driver/ToolChains/Cuda.cpp b/clang/lib/Driver/ToolChains/Cuda.cpp
@@ -965,8 +965,7 @@ void CudaToolChain::addClangTargetOptions(
   if (DeviceOffloadingKind == Action::OFK_SYCL) {
     SYCLInstallation.addSYCLIncludeArgs(DriverArgs, CC1Args);
 
-    if (DriverArgs.hasArg(options::OPT_fsycl_fp32_prec_sqrt) ||
-        DriverArgs.hasArg(options::OPT_foffload_fp32_prec_sqrt))
+    if (DriverArgs.hasArg(options::OPT_foffload_fp32_prec_sqrt))
       CC1Args.push_back("-fcuda-prec-sqrt");
 
     bool FastRelaxedMath = DriverArgs.hasFlag(
diff --git a/clang/test/Driver/sycl-deprecated.cpp b/clang/test/Driver/sycl-deprecated.cpp
@@ -1,8 +1,13 @@
 /// Test for any deprecated options
 // RUN: %clangxx -fsycl-explicit-simd %s -### 2>&1 | FileCheck %s -DOPTION=-fsycl-explicit-simd
 // RUN: %clangxx -fno-sycl-explicit-simd %s -### 2>&1 | FileCheck %s -DOPTION=-fno-sycl-explicit-simd
-// RUN: %clangxx -fsycl -fsycl-use-bitcode %s -### 2>&1 | FileCheck %s -DOPTION=-fsycl-use-bitcode
-// RUN: %clangxx -fsycl -fno-sycl-use-bitcode %s -### 2>&1 | FileCheck %s -DOPTION=-fno-sycl-use-bitcode
 // RUN: %clangxx -fsycl -fsycl-allow-device-dependencies %s -### 2>&1 | FileCheck %s -DOPTION=-fsycl-allow-device-dependencies
 // RUN: %clangxx -fsycl -fno-sycl-allow-device-dependencies %s -### 2>&1 | FileCheck %s -DOPTION=-fno-sycl-allow-device-dependencies
 // CHECK: option '[[OPTION]]' is deprecated and will be removed in a future release
+
+// RUN: %clangxx -fsycl -fsycl-use-bitcode %s -### 2>&1 \
+// RUN: | FileCheck %s --check-prefix=CHECK_REPLACE -DOPTION=-fsycl-use-bitcode -DOPTION_REPLACE=-fsycl-device-obj=llvmir
+// RUN: %clangxx -fsycl -fno-sycl-use-bitcode %s -### 2>&1 \
+// RUN: | FileCheck %s --check-prefix=CHECK_REPLACE -DOPTION=-fno-sycl-use-bitcode -DOPTION_REPLACE=-fsycl-device-obj=spirv
+// RUN: %clangxx -fsycl -fsycl-fp32-prec-sqrt %s -### 2>&1 | FileCheck %s --check-prefix=CHECK_REPLACE -DOPTION=-fsycl-fp32-prec-sqrt -DOPTION_REPLACE=-foffload-fp32-prec-sqrt
+// CHECK_REPLACE: option '[[OPTION]]' is deprecated and will be removed in a future release, use '[[OPTION_REPLACE]]' instead
diff --git a/clang/test/Driver/sycl-no-prec-sqrt.cpp b/clang/test/Driver/sycl-no-prec-sqrt.cpp
diff --git a/clang/test/Driver/sycl-specific-args-diagnostics.cpp b/clang/test/Driver/sycl-specific-args-diagnostics.cpp
@@ -81,11 +81,17 @@
 // RUN:   %clang_cl -### -fsycl-force-inline-kernel-lambda  %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=WARNING-UNUSED-ARG -DOPT=-fsycl-force-inline-kernel-lambda %s
 
-// Warning should be emitted when using -fsycl-fp32-prec-sqrt without -fsycl
-// RUN:   %clang -### -fsycl-fp32-prec-sqrt  %s 2>&1 \
-// RUN:   | FileCheck -check-prefix=WARNING-UNUSED-ARG -DOPT=-fsycl-fp32-prec-sqrt %s
-// RUN:   %clang_cl -### -fsycl-fp32-prec-sqrt  %s 2>&1 \
-// RUN:   | FileCheck -check-prefix=WARNING-UNUSED-ARG -DOPT=-fsycl-fp32-prec-sqrt %s
+// Warning should be emitted when using -foffload-fp32-prec-sqrt without -fsycl
+// RUN:   %clang -### -foffload-fp32-prec-sqrt %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=WARNING-UNUSED-ARG -DOPT=-foffload-fp32-prec-sqrt %s
+// RUN:   %clang_cl -### -foffload-fp32-prec-sqrt %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=WARNING-UNUSED-ARG -DOPT=-foffload-fp32-prec-sqrt %s
+
+// Warning should be emitted when using -foffload-fp32-prec-div without -fsycl
+// RUN:   %clang -### -foffload-fp32-prec-div %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=WARNING-UNUSED-ARG -DOPT=-foffload-fp32-prec-div %s
+// RUN:   %clang_cl -### -foffload-fp32-prec-div %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=WARNING-UNUSED-ARG -DOPT=-foffload-fp32-prec-div %s
 
 // Warning should be emitted when using -fsycl-id-queries-fit-in-int without -fsycl
 // RUN:   %clang -### -fsycl-id-queries-fit-in-int  %s 2>&1 \
diff --git a/sycl/doc/UsersManual.md b/sycl/doc/UsersManual.md
@@ -343,6 +343,7 @@ and not recommended to use in production environment.
     Emit SYCL device code in LLVM-IR bitcode format. When disabled, SPIR-V is
     emitted.
     Enabled by default.
+    This option is replaced with -fsycl-device-obj=<arg>.
 
 **`-fsycl-device-obj=<arg>`** [EXPERIMENTAL]
 
@@ -382,6 +383,7 @@ and not recommended to use in production environment.
     Enable use of correctly rounded `sycl::sqrt` function as defined by IEE754.
     Without this flag, the default precision requirement for `sycl::sqrt` is 3
     ULP.
+    This option is replaced with -foffload-fp32-prec-sqrt.
 
     NOTE: This flag is currently only supported with the CUDA and HIP targets.
 
