rename flag, apply module flag only for sycl

Author: pgorlani

clang/include/clang/Driver/Options.td

@@ -985,10 +985,6 @@ defm cuda_prec_sqrt : BoolFOption<"cuda-prec-sqrt",
   PosFlag<SetTrue, [CC1Option], "Specify">,
   NegFlag<SetFalse, [], "Don't specify">,
   BothFlags<[], " that sqrt is correctly rounded (for CUDA devices)">>;
-defm nvvm_cuda_approx_tanhf : BoolFOption<"cuda-approx-tanhf",
-  TargetOpts<"NVVMCudaApproxTanhf">, DefaultFalse,
-  PosFlag<SetTrue, [CC1Option], "Only SYCL. Use the built-in fast approximation of tanhf function. Device needs to have a compute capability >= 8.0">,
-  NegFlag<SetFalse>>;
 def rocm_path_EQ : Joined<["--"], "rocm-path=">, Group<i_Group>,
   HelpText<"ROCm installation path, used for finding and automatically linking required bitcode libraries.">;
 def hip_path_EQ : Joined<["--"], "hip-path=">, Group<i_Group>,
@@ -4730,7 +4726,10 @@ def fno_sycl_device_lib_EQ : CommaJoined<["-"], "fno-sycl-device-lib=">, Group<s
   Values<"libc, libm-fp32, libm-fp64, all">, HelpText<"Control exclusion of "
   "device libraries from device binary linkage. Valid arguments "
   "are libc, libm-fp32, libm-fp64, all">;
-
+defm nvvm_cuda_approx_tanh : BoolFOption<"sycl-cuda-approx-tanh",
+  TargetOpts<"NVVMCudaApproxTanhf">, DefaultFalse,
+  PosFlag<SetTrue, [CC1Option], "Use the built-in fast approximation of tanh function. Device needs to have a compute capability >= 8.0">,
+  NegFlag<SetFalse>>;
 //===----------------------------------------------------------------------===//
 // FLangOption + CoreOption + NoXarchOption
 //===----------------------------------------------------------------------===//

clang/lib/CodeGen/CodeGenModule.cpp

@@ -780,6 +780,9 @@ void CodeGenModule::Release() {
                                   llvm::DenormalMode::IEEE);
     getModule().addModuleFlag(llvm::Module::Override, "nvvm-reflect-prec-sqrt",
                               getTarget().getTargetOpts().NVVMCudaPrecSqrt);
+  }
+
+  if ( LangOpts.isSYCL() && getTriple().isNVPTX()) {
     getModule().addModuleFlag(llvm::Module::Override,
                               "nvvm-reflect-approx-tanhf",
                               getTarget().getTargetOpts().NVVMCudaApproxTanhf);

clang/test/CodeGenCUDA/flush-denormals.cu

@@ -44,8 +44,8 @@ extern "C" __device__ void foo() {}
 // FTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
 // NOFTZ-NOT: "denormal-fp-math-f32"
 
-// PTXFTZ:!llvm.module.flags = !{{{.*}}, [[MODFLAG:![0-9]+]], {{.*}}, {{.*}}}
+// PTXFTZ:!llvm.module.flags = !{{{.*}}, [[MODFLAG:![0-9]+]], {{.*}}}
 // PTXFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 1}
 
-// PTXNOFTZ:!llvm.module.flags = !{{{.*}}, [[MODFLAG:![0-9]+]], {{.*}}, {{.*}}}
+// PTXNOFTZ:!llvm.module.flags = !{{{.*}}, [[MODFLAG:![0-9]+]], {{.*}}}
 // PTXNOFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 0}

clang/test/CodeGenCUDA/nvvm-reflect-approx-tanh.cu

@@ -1,9 +1,9 @@
-// RUN: %clang_cc1 -fcuda-is-device -triple nvptx64-nvidia-cuda -emit-llvm  -fcuda-approx-tanhf %s -o -| FileCheck --check-prefix=CHECK-ON %s
-// RUN: %clang_cc1 -fcuda-is-device -triple nvptx64-nvidia-cuda -emit-llvm %s -o -| FileCheck --check-prefix=CHECK-OFF %s
+// RUN: %clang_cc1 -fsycl-is-device -triple nvptx64-nvidia-cuda -emit-llvm  -fsycl-cuda-approx-tanh %s -o -| FileCheck --check-prefix=CHECK-ON %s
+// RUN: %clang_cc1 -fsycl-is-device -triple nvptx64-nvidia-cuda -emit-llvm %s -o -| FileCheck --check-prefix=CHECK-OFF %s
 
 #include "Inputs/cuda.h"
 
-// Check that the -fcuda-approx-tanhf flag correctly sets the nvvm-reflect module flags.
+// Check that the -fsycl-cuda-approx-tanh flag correctly sets the nvvm-reflect module flags.
 
 extern "C" __device__ void foo() {}