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

[tra]

Summary

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.

The file is only intended to be -include'd by clang and should never be included by users, hence '__'.

[tra]

Changed name to __clang_cuda_runtime_wrapper.h
Added comments in the header explaining intended use.

[chandlerc]

The substance of the patch LGTM. My nit picking is just on the wording of the comment. =] Submit whenever.

lib/Headers/cuda_runtime.h

1–2

"by compiler" -> "by the compiler"

4–5

You say above that they'll be included by clang directly?

I think instead of "impossible to use them by clang directly" you want to say something more along the lines of "impossible for user code to #include directly when compiling with clang".

6–10

I would consistently capitalize "CUDA" when not talking about a particular header like cuda_runtime.h, and "Clang" and "NVCC" unless talking about running a command.

Some other nits:

"included from" -> "included"

"so we have to abuse preprocessor in order to" -> "so we use the preprocessor to"

"shape CUDA headers into something" -> "force the headers into a form that"

[jhen]

lib/Headers/cuda_runtime.h

0

Now that the name of this header has been changed, would it be appropriate to change this #include_next to a simple #include?

[tra]

Incorporated Chandler's suggestions.
Fixed #include_next -> #include.

lib/Headers/cuda_runtime.h

0

Fixed.

lib/Driver/ToolChains.cpp

 void Linux::AddCudaIncludeArgs(const ArgList &DriverArgs,
                                ArgStringList &CC1Args) const {
   if (DriverArgs.hasArg(options::OPT_nocudainc))
     return;
 
   if (CudaInstallation.isValid()) {
     addSystemInclude(DriverArgs, CC1Args, CudaInstallation.getIncludePath());
     CC1Args.push_back("-include");
-    CC1Args.push_back("cuda_runtime.h");
+    CC1Args.push_back("__cuda_runtime.h");
   }
 }
 
 bool Linux::isPIEDefault() const { return getSanitizerArgs().requiresPIE(); }
 
 SanitizerMask Linux::getSupportedSanitizers() const {
   const bool IsX86 = getTriple().getArch() == llvm::Triple::x86;
   const bool IsX86_64 = getTriple().getArch() == llvm::Triple::x86_64;

lib/Headers/CMakeLists.txt

set(files
   avx512vlintrin.h
   avx512dqintrin.h
   avx512vldqintrin.h
   avxintrin.h
   bmi2intrin.h
   bmiintrin.h
   cpuid.h
   cuda_builtin_vars.h
-  cuda_runtime.h
+  __cuda_runtime.h
   emmintrin.h
   f16cintrin.h
   float.h
   fma4intrin.h
   fmaintrin.h
   fxsrintrin.h
   htmintrin.h
   htmxlintrin.h

lib/Headers/__cuda_runtime.h

 /*===---- cuda_runtime.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
  * in the Software without restriction, including without limitation the rights
  * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
  * copies of the Software, and to permit persons to whom the Software is
  * furnished to do so, subject to the following conditions:
  *
  * The above copyright notice and this permission notice shall be included in
  * all copies or substantial portions of the Software.
  *
  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
  * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
  * THE SOFTWARE.
  *
  *===-----------------------------------------------------------------------===
  */
 
 #ifndef __CLANG_CUDA_RUNTIME_H__
 #define __CLANG_CUDA_RUNTIME_H__
 
 #if defined(__CUDA__) && defined(__clang__)
 
 // Include some standard headers to avoid CUDA headers including them
 // while some required macros (like __THROW) are in a weird state.
 #include <stdlib.h>
 
 // Preserve common macros that will be changed below by us or by CUDA
 // headers.
 #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.
 #include "cuda.h"
 #if !defined(CUDA_VERSION)
 #error "cuda.h did not define CUDA_VERSION"
 #elif CUDA_VERSION < 7000 || CUDA_VERSION > 7050
 #error "Unsupported CUDA version!"
 #endif
 
 // Make largest subset of device functions available during host
 // compilation -- SM_35 for the time being.
 #ifndef __CUDA_ARCH__
 #define __CUDA_ARCH__ 350
 #endif
 
 #include "cuda_builtin_vars.h"
 
 // No need for device_launch_parameters.h as cuda_builtin_vars.h above
 // has taken care of builtin variables declared in the file.
 #define __DEVICE_LAUNCH_PARAMETERS_H__
 
 // {math,device}_functions.h only have declarations of the
 // functions. We don't need them as we're going to pull in their
 // definitions from .hpp files.
 #define __DEVICE_FUNCTIONS_H__
 #define __MATH_FUNCTIONS_H__
 
 #undef __CUDACC__
 #define __CUDABE__
 // Disables definitions of device-side runtime support stubs in
 // cuda_device_runtime_api.h
 #define __CUDADEVRT_INTERNAL__
 #include "host_config.h"
 #include "host_defines.h"
 #include "driver_types.h"
 #include "common_functions.h"
 #undef __CUDADEVRT_INTERNAL__
 
 #undef __CUDABE__
 #define __CUDACC__
 #include_next "cuda_runtime.h"
 
 #undef __CUDACC__
 #define __CUDABE__
 
 // 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)
 
 #include "crt/host_runtime.h"
 #include "crt/device_runtime.h"
 // device_runtime.h defines __cxa_* macros that will conflict with
 // cxxabi.h.
 // FIXME: redefine these as __device__ functions.
 #undef __cxa_vec_ctor
 #undef __cxa_vec_cctor
 #undef __cxa_vec_dtor
 #undef __cxa_vec_new2
 #undef __cxa_vec_new3
 #undef __cxa_vec_delete2
 #undef __cxa_vec_delete
 #undef __cxa_vec_delete3
 #undef __cxa_pure_virtual
 
 // We need decls for functions in CUDA's libdevice woth __device__
 // attribute only. Alas they come either as __host__ __device__ or
 // with no attributes at all. To work around that, define __CUDA_RTC__
 // which produces HD variant and undef __host__ which gives us desided
 // decls with __device__ attribute.
 #pragma push_macro("__host__")
 #define __host__
 #define __CUDACC_RTC__
 #include "device_functions_decls.h"
 #undef __CUDACC_RTC__
 
 // Temporarily poison __host__ macro to ensure it's not used by any of
 // the headers we're about to include.
 #define __host__ UNEXPECTED_HOST_ATTRIBUTE
 
 // device_functions.hpp and math_functions*.hpp use 'static
 // __forceinline__' (with no __device__) for definitions of device
 // functions. Temporarily redefine __forceinline__ to include
 // __device__.
 #pragma push_macro("__forceinline__")
 #define __forceinline__ __device__ __inline__ __attribute__((always_inline))
 #include "device_functions.hpp"
 #include "math_functions.hpp"
 #include "math_functions_dbl_ptx3.hpp"
 #pragma pop_macro("__forceinline__")
 
 // For some reason single-argument variant is not always declared by
 // CUDA headers. Alas, device_functions.hpp included below needs it.
 static inline __device__ void __brkpt(int c) { __brkpt(); }
 
 // Now include *.hpp with definitions of various GPU functions.  Alas,
 // a lot of thins get declared/defined with __host__ attribute which
 // we don't want and we have to define it out. We also have to include
 // {device,math}_functions.hpp again in order to extract the other
 // branch of #if/else inside.
 
 #define __host__
 #undef __CUDABE__
 #define __CUDACC__
 #undef __DEVICE_FUNCTIONS_HPP__
 #include "device_functions.hpp"
 #include "device_atomic_functions.hpp"
 #include "sm_20_atomic_functions.hpp"
 #include "sm_32_atomic_functions.hpp"
 #include "sm_20_intrinsics.hpp"
 // sm_30_intrinsics.h has declarations that use default argument, so
 // we have to include it and it will in turn include .hpp
 #include "sm_30_intrinsics.h"
 #include "sm_32_intrinsics.hpp"
 #undef __MATH_FUNCTIONS_HPP__
 #include "math_functions.hpp"
 #pragma pop_macro("__host__")
 
 #include "texture_indirect_functions.h"
 
 // Restore state of __CUDA_ARCH__ and __THROW we had on entry.
 #pragma pop_macro("__CUDA_ARCH__")
 #pragma pop_macro("__THROW")
 
 // Set up compiler macros expected to be seen during compilation.
 #undef __CUDABE__
 #define __CUDACC__
 #define __NVCC__
 
 #if defined(__CUDA_ARCH__)
 // We need to emit IR declaration for non-existing __nvvm_reflect to
 // let backend know that it should be treated as const nothrow
 // function which is implicitly assumed by NVVMReflect pass.
 extern "C" __device__ __attribute__((const)) int __nvvm_reflect(const void *);
 static __device__ __attribute__((used)) int __nvvm_reflect_anchor() {
   return __nvvm_reflect("NONE");
 }
 #endif
 
 #endif // __CUDA__
 #endif // __CLANG_CUDA_RUNTIME_H__

lib/Headers/cuda_runtime.h

-/*===---- cuda_runtime.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
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in
- * all copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
- * THE SOFTWARE.
- *
- *===-----------------------------------------------------------------------===
- */
-
-#ifndef __CLANG_CUDA_RUNTIME_H__
-#define __CLANG_CUDA_RUNTIME_H__
-
-#if defined(__CUDA__) && defined(__clang__)
-
-// Include some standard headers to avoid CUDA headers including them
-// while some required macros (like __THROW) are in a weird state.
-#include <stdlib.h>
-
-// Preserve common macros that will be changed below by us or by CUDA
-// headers.
-#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.
-#include "cuda.h"
-#if !defined(CUDA_VERSION)
-#error "cuda.h did not define CUDA_VERSION"
-#elif CUDA_VERSION < 7000 || CUDA_VERSION > 7050
-#error "Unsupported CUDA version!"
-#endif
-
-// Make largest subset of device functions available during host
-// compilation -- SM_35 for the time being.
-#ifndef __CUDA_ARCH__
-#define __CUDA_ARCH__ 350
-#endif
-
-#include "cuda_builtin_vars.h"
-
-// No need for device_launch_parameters.h as cuda_builtin_vars.h above
-// has taken care of builtin variables declared in the file.
-#define __DEVICE_LAUNCH_PARAMETERS_H__
-
-// {math,device}_functions.h only have declarations of the
-// functions. We don't need them as we're going to pull in their
-// definitions from .hpp files.
-#define __DEVICE_FUNCTIONS_H__
-#define __MATH_FUNCTIONS_H__
-
-#undef __CUDACC__
-#define __CUDABE__
-// Disables definitions of device-side runtime support stubs in
-// cuda_device_runtime_api.h
-#define __CUDADEVRT_INTERNAL__
-#include "host_config.h"
-#include "host_defines.h"
-#include "driver_types.h"
-#include "common_functions.h"
-#undef __CUDADEVRT_INTERNAL__
-
-#undef __CUDABE__
-#define __CUDACC__
-#include_next "cuda_runtime.h"
-
-#undef __CUDACC__
-#define __CUDABE__
-
-// 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)
-
-#include "crt/host_runtime.h"
-#include "crt/device_runtime.h"
-// device_runtime.h defines __cxa_* macros that will conflict with
-// cxxabi.h.
-// FIXME: redefine these as __device__ functions.
-#undef __cxa_vec_ctor
-#undef __cxa_vec_cctor
-#undef __cxa_vec_dtor
-#undef __cxa_vec_new2
-#undef __cxa_vec_new3
-#undef __cxa_vec_delete2
-#undef __cxa_vec_delete
-#undef __cxa_vec_delete3
-#undef __cxa_pure_virtual
-
-// We need decls for functions in CUDA's libdevice woth __device__
-// attribute only. Alas they come either as __host__ __device__ or
-// with no attributes at all. To work around that, define __CUDA_RTC__
-// which produces HD variant and undef __host__ which gives us desided
-// decls with __device__ attribute.
-#pragma push_macro("__host__")
-#define __host__
-#define __CUDACC_RTC__
-#include "device_functions_decls.h"
-#undef __CUDACC_RTC__
-
-// Temporarily poison __host__ macro to ensure it's not used by any of
-// the headers we're about to include.
-#define __host__ UNEXPECTED_HOST_ATTRIBUTE
-
-// device_functions.hpp and math_functions*.hpp use 'static
-// __forceinline__' (with no __device__) for definitions of device
-// functions. Temporarily redefine __forceinline__ to include
-// __device__.
-#pragma push_macro("__forceinline__")
-#define __forceinline__ __device__ __inline__ __attribute__((always_inline))
-#include "device_functions.hpp"
-#include "math_functions.hpp"
-#include "math_functions_dbl_ptx3.hpp"
-#pragma pop_macro("__forceinline__")
-
-// For some reason single-argument variant is not always declared by
-// CUDA headers. Alas, device_functions.hpp included below needs it.
-static inline __device__ void __brkpt(int c) { __brkpt(); }
-
-// Now include *.hpp with definitions of various GPU functions.  Alas,
-// a lot of thins get declared/defined with __host__ attribute which
-// we don't want and we have to define it out. We also have to include
-// {device,math}_functions.hpp again in order to extract the other
-// branch of #if/else inside.
-
-#define __host__
-#undef __CUDABE__
-#define __CUDACC__
-#undef __DEVICE_FUNCTIONS_HPP__
-#include "device_functions.hpp"
-#include "device_atomic_functions.hpp"
-#include "sm_20_atomic_functions.hpp"
-#include "sm_32_atomic_functions.hpp"
-#include "sm_20_intrinsics.hpp"
-// sm_30_intrinsics.h has declarations that use default argument, so
-// we have to include it and it will in turn include .hpp
-#include "sm_30_intrinsics.h"
-#include "sm_32_intrinsics.hpp"
-#undef __MATH_FUNCTIONS_HPP__
-#include "math_functions.hpp"
-#pragma pop_macro("__host__")
-
-#include "texture_indirect_functions.h"
-
-// Restore state of __CUDA_ARCH__ and __THROW we had on entry.
-#pragma pop_macro("__CUDA_ARCH__")
-#pragma pop_macro("__THROW")
-
-// Set up compiler macros expected to be seen during compilation.
-#undef __CUDABE__
-#define __CUDACC__
-#define __NVCC__
-
-#if defined(__CUDA_ARCH__)
-// We need to emit IR declaration for non-existing __nvvm_reflect to
-// let backend know that it should be treated as const nothrow
-// function which is implicitly assumed by NVVMReflect pass.
-extern "C" __device__ __attribute__((const)) int __nvvm_reflect(const void *);
-static __device__ __attribute__((used)) int __nvvm_reflect_anchor() {
-  return __nvvm_reflect("NONE");
-}
-#endif
-
-#endif // __CUDA__
-#endif // __CLANG_CUDA_RUNTIME_H__

test/Driver/cuda-detect.cu

 // RUN: %clang -### -v --target=i386-unknown-linux --cuda-gpu-arch=sm_30 \
 // RUN:   --cuda-path=%S/Inputs/CUDA/usr/local/cuda %s 2>&1 \
 // RUN:   | FileCheck %s -check-prefix COMMON -check-prefix NOLIBDEVICE
 // .. or if we explicitly passed -nocudalib
 // RUN: %clang -### -v --target=i386-unknown-linux --cuda-gpu-arch=sm_35 \
 // 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 __cuda_runtime 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 \
 // RUN:     -check-prefix NOCUDAINC -check-prefix NOLIBDEVICE
 
 // CHECK: Found CUDA installation: {{.*}}/Inputs/CUDA/usr/local/cuda
 // NOCUDA-NOT: Found CUDA installation:
 
 // COMMON: "-triple" "nvptx-nvidia-cuda"
 // COMMON-SAME: "-fcuda-is-device"
 // LIBDEVICE-SAME: "-mlink-cuda-bitcode"
 // NOLIBDEVICE-NOT: "-mlink-cuda-bitcode"
 // LIBDEVICE21-SAME: libdevice.compute_20.10.bc
 // LIBDEVICE35-SAME: libdevice.compute_35.10.bc
 // NOLIBDEVICE-NOT: libdevice.compute_{{.*}}.bc
 // LIBDEVICE-SAME: "-target-feature" "+ptx42"
 // 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" "__cuda_runtime.h"
+// NOCUDAINC-NOT: "-include" "__cuda_runtime.h"
 // COMMON-SAME: "-x" "cuda"