r270484 - [CUDA] Add -fcuda-approx-transcendentals flag.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Mon May 23 13:19:57 PDT 2016
Author: jlebar
Date: Mon May 23 15:19:56 2016
New Revision: 270484
URL: http://llvm.org/viewvc/llvm-project?rev=270484&view=rev
Log:
[CUDA] Add -fcuda-approx-transcendentals flag.
Summary:
This lets us emit e.g. sin.approx.f32. See
http://docs.nvidia.com/cuda/parallel-thread-execution/#floating-point-instructions-sin
Reviewers: rnk
Subscribers: tra, cfe-commits
Differential Revision: http://reviews.llvm.org/D20493
Added:
cfe/trunk/test/Preprocessor/cuda-approx-transcendentals.cu
Modified:
cfe/trunk/include/clang/Basic/LangOptions.def
cfe/trunk/include/clang/Driver/Options.td
cfe/trunk/lib/Driver/ToolChains.cpp
cfe/trunk/lib/Frontend/CompilerInvocation.cpp
cfe/trunk/lib/Frontend/InitPreprocessor.cpp
cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h
Modified: cfe/trunk/include/clang/Basic/LangOptions.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/LangOptions.def?rev=270484&r1=270483&r2=270484&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/LangOptions.def (original)
+++ cfe/trunk/include/clang/Basic/LangOptions.def Mon May 23 15:19:56 2016
@@ -190,6 +190,7 @@ LANGOPT(CUDAIsDevice , 1, 0, "compi
LANGOPT(CUDAAllowVariadicFunctions, 1, 0, "allowing variadic functions in CUDA device code")
LANGOPT(CUDAHostDeviceConstexpr, 1, 1, "treating unattributed constexpr functions as __host__ __device__")
LANGOPT(CUDADeviceFlushDenormalsToZero, 1, 0, "flushing denormals to zero")
+LANGOPT(CUDADeviceApproxTranscendentals, 1, 0, "using approximate transcendental functions")
LANGOPT(SizedDeallocation , 1, 0, "enable sized deallocation functions")
LANGOPT(ConceptsTS , 1, 0, "enable C++ Extensions for Concepts")
Modified: cfe/trunk/include/clang/Driver/Options.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Driver/Options.td?rev=270484&r1=270483&r2=270484&view=diff
==============================================================================
--- cfe/trunk/include/clang/Driver/Options.td (original)
+++ cfe/trunk/include/clang/Driver/Options.td Mon May 23 15:19:56 2016
@@ -395,6 +395,9 @@ def cuda_path_EQ : Joined<["--"], "cuda-
def fcuda_flush_denormals_to_zero : Flag<["-"], "fcuda-flush-denormals-to-zero">,
Flags<[CC1Option]>, HelpText<"Flush denormal floating point values to zero in CUDA device mode.">;
def fno_cuda_flush_denormals_to_zero : Flag<["-"], "fno-cuda-flush-denormals-to-zero">;
+def fcuda_approx_transcendentals : Flag<["-"], "fcuda-approx-transcendentals">,
+ Flags<[CC1Option]>, HelpText<"Use approximate transcendental functions">;
+def fno_cuda_approx_transcendentals : Flag<["-"], "fno-cuda-approx-transcendentals">;
def dA : Flag<["-"], "dA">, Group<d_Group>;
def dD : Flag<["-"], "dD">, Group<d_Group>, Flags<[CC1Option]>,
HelpText<"Print macro definitions in -E mode in addition to normal output">;
Modified: cfe/trunk/lib/Driver/ToolChains.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Driver/ToolChains.cpp?rev=270484&r1=270483&r2=270484&view=diff
==============================================================================
--- cfe/trunk/lib/Driver/ToolChains.cpp (original)
+++ cfe/trunk/lib/Driver/ToolChains.cpp Mon May 23 15:19:56 2016
@@ -4502,6 +4502,10 @@ CudaToolChain::addClangTargetOptions(con
options::OPT_fno_cuda_flush_denormals_to_zero, false))
CC1Args.push_back("-fcuda-flush-denormals-to-zero");
+ if (DriverArgs.hasFlag(options::OPT_fcuda_approx_transcendentals,
+ options::OPT_fno_cuda_approx_transcendentals, false))
+ CC1Args.push_back("-fcuda-approx-transcendentals");
+
if (DriverArgs.hasArg(options::OPT_nocudalib))
return;
Modified: cfe/trunk/lib/Frontend/CompilerInvocation.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Frontend/CompilerInvocation.cpp?rev=270484&r1=270483&r2=270484&view=diff
==============================================================================
--- cfe/trunk/lib/Frontend/CompilerInvocation.cpp (original)
+++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp Mon May 23 15:19:56 2016
@@ -1616,6 +1616,9 @@ static void ParseLangArgs(LangOptions &O
if (Opts.CUDAIsDevice && Args.hasArg(OPT_fcuda_flush_denormals_to_zero))
Opts.CUDADeviceFlushDenormalsToZero = 1;
+ if (Opts.CUDAIsDevice && Args.hasArg(OPT_fcuda_approx_transcendentals))
+ Opts.CUDADeviceApproxTranscendentals = 1;
+
if (Opts.ObjC1) {
if (Arg *arg = Args.getLastArg(OPT_fobjc_runtime_EQ)) {
StringRef value = arg->getValue();
Modified: cfe/trunk/lib/Frontend/InitPreprocessor.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Frontend/InitPreprocessor.cpp?rev=270484&r1=270483&r2=270484&view=diff
==============================================================================
--- cfe/trunk/lib/Frontend/InitPreprocessor.cpp (original)
+++ cfe/trunk/lib/Frontend/InitPreprocessor.cpp Mon May 23 15:19:56 2016
@@ -938,6 +938,12 @@ static void InitializePredefinedMacros(c
Builder.defineMacro("__CUDA_ARCH__");
}
+ // We need to communicate this to our CUDA header wrapper, which in turn
+ // informs the proper CUDA headers of this choice.
+ if (LangOpts.CUDADeviceApproxTranscendentals || LangOpts.FastMath) {
+ Builder.defineMacro("__CLANG_CUDA_APPROX_TRANSCENDENTALS__");
+ }
+
// OpenCL definitions.
if (LangOpts.OpenCL) {
#define OPENCLEXT(Ext) \
Modified: cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h?rev=270484&r1=270483&r2=270484&view=diff
==============================================================================
--- cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h (original)
+++ cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h Mon May 23 15:19:56 2016
@@ -142,7 +142,20 @@
#pragma push_macro("__forceinline__")
#define __forceinline__ __device__ __inline__ __attribute__((always_inline))
#include "device_functions.hpp"
+
+// math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we
+// get the slow-but-accurate or fast-but-inaccurate versions of functions like
+// sin and exp. This is controlled in clang by -fcuda-approx-transcendentals.
+//
+// device_functions.hpp uses __USE_FAST_MATH__ for a different purpose (fast vs.
+// slow divides), so we need to scope our define carefully here.
+#pragma push_macro("__USE_FAST_MATH__")
+#if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__)
+#define __USE_FAST_MATH__
+#endif
#include "math_functions.hpp"
+#pragma pop_macro("__USE_FAST_MATH__")
+
#include "math_functions_dbl_ptx3.hpp"
#pragma pop_macro("__forceinline__")
@@ -296,6 +309,7 @@ __device__ inline __cuda_builtin_gridDim
#include "curand_mtgp32_kernel.h"
#pragma pop_macro("dim3")
#pragma pop_macro("uint3")
+#pragma pop_macro("__USE_FAST_MATH__")
#endif // __CUDA__
#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__
Added: cfe/trunk/test/Preprocessor/cuda-approx-transcendentals.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Preprocessor/cuda-approx-transcendentals.cu?rev=270484&view=auto
==============================================================================
--- cfe/trunk/test/Preprocessor/cuda-approx-transcendentals.cu (added)
+++ cfe/trunk/test/Preprocessor/cuda-approx-transcendentals.cu Mon May 23 15:19:56 2016
@@ -0,0 +1,8 @@
+// RUN: %clang --cuda-host-only -nocudainc -target i386-unknown-linux-gnu -x cuda -E -dM -o - /dev/null | FileCheck --check-prefix HOST %s
+// RUN: %clang --cuda-device-only -nocudainc -target i386-unknown-linux-gnu -x cuda -E -dM -o - /dev/null | FileCheck --check-prefix DEVICE-NOFAST %s
+// RUN: %clang -fcuda-approx-transcendentals --cuda-device-only -nocudainc -target i386-unknown-linux-gnu -x cuda -E -dM -o - /dev/null | FileCheck --check-prefix DEVICE-FAST %s
+// RUN: %clang -ffast-math --cuda-device-only -nocudainc -target i386-unknown-linux-gnu -x cuda -E -dM -o - /dev/null | FileCheck --check-prefix DEVICE-FAST %s
+
+// HOST-NOT: __CLANG_CUDA_APPROX_TRANSCENDENTALS__
+// DEVICE-NOFAST-NOT: __CLANG_CUDA_APPROX_TRANSCENDENTALS__
+// DEVICE-FAST: __CLANG_CUDA_APPROX_TRANSCENDENTALS__
More information about the cfe-commits
mailing list