From bbc017851815da300ce9d6315204a73d68754a1c Mon Sep 17 00:00:00 2001 From: Reid Kleckner Date: Wed, 3 Dec 2014 21:53:36 +0000 Subject: 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 --- clang/lib/Frontend/InitPreprocessor.cpp | 7 +++++++ 1 file changed, 7 insertions(+) (limited to 'clang/lib/Frontend/InitPreprocessor.cpp') diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp index f671a2f6..3550ac2 100644 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -870,6 +870,13 @@ static void InitializePredefinedMacros(const TargetInfo &TI, 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); } -- cgit v1.1