[CUDA] renamed cuda_runtime.h wrapper to __cuda_runtime.h

Currently it's easy to break CUDA compilation by passing
"-isystem /path/to/cuda/include" to compiler which leads to
compiler including real cuda_runtime.h from there instead
of the wrapper we need.

Renaming the wrapper ensures that we can include the wrapper
regardless of user-specified include paths and files.

Differential Revision: http://reviews.llvm.org/D15534

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@255802 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/Driver/ToolChains.cpp b/lib/Driver/ToolChains.cpp
index 2d882eb5368506d8527c2a6b7c346ca612a1635e..0921bc19947db3abefd0df03c5cbd66031d446b1 100644 (file)
--- a/lib/Driver/ToolChains.cpp
+++ b/lib/Driver/ToolChains.cpp
@@ -4116,7 +4116,7 @@ void Linux::AddCudaIncludeArgs(const ArgList &DriverArgs,
   if (CudaInstallation.isValid()) {
     addSystemInclude(DriverArgs, CC1Args, CudaInstallation.getIncludePath());
     CC1Args.push_back("-include");
-    CC1Args.push_back("cuda_runtime.h");
+    CC1Args.push_back("__clang_cuda_runtime_wrapper.h");
   }
 }
 

diff --git a/lib/Headers/CMakeLists.txt b/lib/Headers/CMakeLists.txt
index 139afa2751a6a388470ddaaa5d1b69018bb423dc..9393f69d41fab8b84d3a48ac65e6a1bc931e8178 100644 (file)
--- a/lib/Headers/CMakeLists.txt
+++ b/lib/Headers/CMakeLists.txt
@@ -15,9 +15,9 @@ set(files
   avxintrin.h
   bmi2intrin.h
   bmiintrin.h
+  __clang_cuda_runtime_wrapper.h
   cpuid.h
   cuda_builtin_vars.h
-  cuda_runtime.h
   emmintrin.h
   f16cintrin.h
   float.h

diff --git a/lib/Headers/cuda_runtime.h b/lib/Headers/__clang_cuda_runtime_wrapper.h
similarity index 84%
rename from lib/Headers/cuda_runtime.h
rename to lib/Headers/__clang_cuda_runtime_wrapper.h
index 1153690e9fa9c6032d033d10c94c4d408a7dac7d..a88606a5eb6bd950403529e968068c9ca2391029 100644 (file)
--- a/lib/Headers/cuda_runtime.h
+++ b/lib/Headers/__clang_cuda_runtime_wrapper.h
@@ -1,4 +1,4 @@
-/*===---- cuda_runtime.h - CUDA runtime support ----------------------------===
+/*===---- __clang_cuda_runtime_wrapper.h - CUDA runtime support -------------===
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -21,8 +21,24 @@
  *===-----------------------------------------------------------------------===
  */
 
-#ifndef __CLANG_CUDA_RUNTIME_H__
-#define __CLANG_CUDA_RUNTIME_H__
+/*
+ * WARNING: This header is intended to be directly -include'd by
+ * the compiler and is not supposed to be included by users.
+ *
+ * CUDA headers are implemented in a way that currently makes it
+ * impossible for user code to #include directly when compiling with
+ * Clang. They present different view of CUDA-supplied functions
+ * depending on where in NVCC's compilation pipeline the headers are
+ * included. Neither of these modes provides function definitions with
+ * correct attributes, so we use preprocessor to force the headers
+ * into a form that Clang can use.
+ *
+ * Similarly to NVCC which -include's cuda_runtime.h, Clang -include's
+ * this file during every CUDA compilation.
+ */
+
+#ifndef __CLANG_CUDA_RUNTIME_WRAPPER_H__
+#define __CLANG_CUDA_RUNTIME_WRAPPER_H__
 
 #if defined(__CUDA__) && defined(__clang__)
 
@@ -35,9 +51,9 @@
 #pragma push_macro("__THROW")
 #pragma push_macro("__CUDA_ARCH__")
 
-// WARNING: Preprocessor hacks below are based on specific of
-// implementation of CUDA-7.x headers and are expected to break with
-// any other version of CUDA headers.
+// WARNING: Preprocessor hacks below are based on specific details of
+// CUDA-7.x headers and are not expected to work with any other
+// version of CUDA headers.
 #include "cuda.h"
 #if !defined(CUDA_VERSION)
 #error "cuda.h did not define CUDA_VERSION"
@@ -76,12 +92,12 @@
 
 #undef __CUDABE__
 #define __CUDACC__
-#include_next "cuda_runtime.h"
+#include "cuda_runtime.h"
 
 #undef __CUDACC__
 #define __CUDABE__
 
-// CUDA headers use __nvvm_memcpy and __nvvm_memset which clang does
+// CUDA headers use __nvvm_memcpy and __nvvm_memset which Clang does
 // not have at the moment. Emulate them with a builtin memcpy/memset.
 #define __nvvm_memcpy(s,d,n,a) __builtin_memcpy(s,d,n)
 #define __nvvm_memset(d,c,n,a) __builtin_memset(d,c,n)
@@ -176,4 +192,4 @@ static __device__ __attribute__((used)) int __nvvm_reflect_anchor() {
 #endif
 
 #endif // __CUDA__
-#endif // __CLANG_CUDA_RUNTIME_H__
+#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__

diff --git a/test/Driver/cuda-detect.cu b/test/Driver/cuda-detect.cu
index 160ca11964fc20613d0c55e1347af631a0182221..d8fba06605e93ff1e97f2df2f5319fd1f55ea4ad 100644 (file)
--- a/test/Driver/cuda-detect.cu
+++ b/test/Driver/cuda-detect.cu
@@ -39,7 +39,7 @@
 // RUN:   -nocudalib --cuda-path=%S/Inputs/CUDA/usr/local/cuda %s 2>&1 \
 // RUN:   | FileCheck %s -check-prefix COMMON -check-prefix NOLIBDEVICE
 // Verify that we don't add include paths, link with libdevice or
-// -include cuda_runtime without valid CUDA installation.
+// -include __clang_cuda_runtime_wrapper.h without valid CUDA installation.
 // RUN: %clang -### -v --target=i386-unknown-linux --cuda-gpu-arch=sm_35 \
 // RUN:   --cuda-path=%S/no-cuda-there %s 2>&1 \
 // RUN:   | FileCheck %s -check-prefix COMMON \
@@ -59,6 +59,6 @@
 // NOLIBDEVICE-NOT: "-target-feature" "+ptx42"
 // CUDAINC-SAME: "-internal-isystem" "{{.*}}/Inputs/CUDA/usr/local/cuda/include"
 // NOCUDAINC-NOT: "-internal-isystem" "{{.*}}/cuda/include"
-// CUDAINC-SAME: "-include" "cuda_runtime.h"
-// NOCUDAINC-NOT: "-include" "cuda_runtime.h"
+// CUDAINC-SAME: "-include" "__clang_cuda_runtime_wrapper.h"
+// NOCUDAINC-NOT: "-include" "__clang_cuda_runtime_wrapper.h"
 // COMMON-SAME: "-x" "cuda"