CUDA host device code with two code paths
Summary:
Allow CUDA host device functions with two code paths using __CUDA_ARCH__
to differentiate between code path being compiled.
For example:
__host__ __device__ void host_device_function(void) {
#ifdef __CUDA_ARCH__
device_only_function();
#else
host_only_function();
#endif
}
Patch by Jacques Pienaar.
Reviewed By: rnk
Differential Revision: http://reviews.llvm.org/D6457
llvm-svn: 223271
diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index f671a2f..3550ac2 100644
--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
@@ -870,6 +870,13 @@
Builder.defineMacro("_OPENMP", "201307");
}
+ // CUDA device path compilaton
+ if (LangOpts.CUDAIsDevice) {
+ // The CUDA_ARCH value is set for the GPU target specified in the NVPTX
+ // backend's target defines.
+ Builder.defineMacro("__CUDA_ARCH__");
+ }
+
// Get other target #defines.
TI.getTargetDefines(LangOpts, Builder);
}