From b208e876583162225f640f8925722256b7b82ea7 Mon Sep 17 00:00:00 2001 From: Artem Belevich Date: Wed, 16 Dec 2015 18:51:59 +0000 Subject: [PATCH] [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 --- lib/Driver/ToolChains.cpp | 2 +- lib/Headers/CMakeLists.txt | 2 +- ...ntime.h => __clang_cuda_runtime_wrapper.h} | 34 ++++++++++++++----- test/Driver/cuda-detect.cu | 6 ++-- 4 files changed, 30 insertions(+), 14 deletions(-) rename lib/Headers/{cuda_runtime.h => __clang_cuda_runtime_wrapper.h} (84%) diff --git a/lib/Driver/ToolChains.cpp b/lib/Driver/ToolChains.cpp index 2d882eb536..0921bc1994 100644 --- 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 139afa2751..9393f69d41 100644 --- 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 1153690e9f..a88606a5eb 100644 --- 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 160ca11964..d8fba06605 100644 --- 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" -- 2.40.0