Author: tra Date: Wed Dec 16 12:51:59 2015 New Revision: 255802 URL: http://llvm.org/viewvc/llvm-project?rev=255802&view=rev Log: [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 Added: cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h - copied, changed from r255655, cfe/trunk/lib/Headers/cuda_runtime.h Removed: cfe/trunk/lib/Headers/cuda_runtime.h Modified: cfe/trunk/lib/Driver/ToolChains.cpp cfe/trunk/lib/Headers/CMakeLists.txt cfe/trunk/test/Driver/cuda-detect.cu Modified: cfe/trunk/lib/Driver/ToolChains.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Driver/ToolChains.cpp?rev=255802&r1=255801&r2=255802&view=diff ============================================================================== --- cfe/trunk/lib/Driver/ToolChains.cpp (original) +++ cfe/trunk/lib/Driver/ToolChains.cpp Wed Dec 16 12:51:59 2015 @@ -4116,7 +4116,7 @@ void Linux::AddCudaIncludeArgs(const Arg 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"); } } Modified: cfe/trunk/lib/Headers/CMakeLists.txt URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/CMakeLists.txt?rev=255802&r1=255801&r2=255802&view=diff ============================================================================== --- cfe/trunk/lib/Headers/CMakeLists.txt (original) +++ cfe/trunk/lib/Headers/CMakeLists.txt Wed Dec 16 12:51:59 2015 @@ -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 Copied: cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h (from r255655, cfe/trunk/lib/Headers/cuda_runtime.h) URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h?p2=cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h&p1=cfe/trunk/lib/Headers/cuda_runtime.h&r1=255655&r2=255802&rev=255802&view=diff ============================================================================== --- cfe/trunk/lib/Headers/cuda_runtime.h (original) +++ cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h Wed Dec 16 12:51:59 2015 @@ -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)) #endif #endif // __CUDA__ -#endif // __CLANG_CUDA_RUNTIME_H__ +#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__ Removed: cfe/trunk/lib/Headers/cuda_runtime.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/cuda_runtime.h?rev=255801&view=auto ============================================================================== --- cfe/trunk/lib/Headers/cuda_runtime.h (original) +++ cfe/trunk/lib/Headers/cuda_runtime.h (removed) @@ -1,179 +0,0 @@ -/*===---- 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__ Modified: cfe/trunk/test/Driver/cuda-detect.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Driver/cuda-detect.cu?rev=255802&r1=255801&r2=255802&view=diff ============================================================================== --- cfe/trunk/test/Driver/cuda-detect.cu (original) +++ cfe/trunk/test/Driver/cuda-detect.cu Wed Dec 16 12:51:59 2015 @@ -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" _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
