[CUDA] Added a wrapper header for inclusion of stock CUDA headers.

[tra]

Summary

Header files that come with CUDA are assuming split host/device compilation and are not usable by clang out of the box.
With a bit of preprocessor magic it's possible to twist them into a form usable by clang after D13170 and D13144 land.

[tra]

Fixed typos and whitespace nits.
use #pragma push_macro for CUDACC_RTC, too.

[eliben]

lib/Headers/clang_cuda_support.h

54

If this includes CUDA headers, maybe you can #error out if the CUDA version isn't good?

[tra]

Include cuda.h and #error if we see wrong CUDA_VERSION.

[eliben]

lgtm

[echristo]

Bikeshed: it's part of the clang headers, do we really need "clang" in the header name?

-eric

[tra]

The (vague) idea was to make clear that the header is *not* part of cuda
distribution.

That said, the file could use a better name.

Do any of these sound better?

• fix_cuda_headers.h
• adapt_cuda_headers.h
• cuda_shim.h

--Artem

[echristo]

Why not just call it cuda.h and use #include next for it and then fix it up?

-eric

[tra]

cuda_runtime.h may be a better choice. nvcc -includes it during both host
and device compilation and this wrapper file is intended to serve a similar
purpose and will probably end up being -included by cc1 in the end.

--Artem

[tra]

Renamed wrapper to cuda_runtime.h
Similarly to nvcc, automatically add "-include cuda_runtime.h" to CC1 invocations unless -nocudainc is specified.

[tra]

Changed header wrapping strategy. Previous version was attempting to
make CUDA headers work for host/device compilations separately. In the
end host and device compilations ended up with different view of
CUDA-provided functions. While it mostly worked, that is not what we
really want. What we want is to have identical view of device-specific
functions in both cases and let function overloading handle name clashes
between host and device functions.

This wrapper now always includes CUDA headers exactly the same way during
host and device compilation passes and produces identical preprocessed
content during host and device side compilation for sm_35 GPUs. Device
compilation passes for older GPUs will see a smaller subset of device
functions supported by particular GPU.

As a bonus this wrapper works with CUDA 7.5 now.

[echristo]

I'm ignoring the content of the header, but this seems to be a not terrible way to do things. I gather that cuda_runtime.h is something that's typically included by the driver by nvidia and not the client?

Also, tests?

-eric

[tra]

In D13171#272397, @echristo wrote:

I'm ignoring the content of the header, but this seems to be a not terrible way to do things. I gather that cuda_runtime.h is something that's typically included by the driver by nvidia and not the client?

Correct. cuda_runtime.h (and all it pulls in) is -include'd under the hood by nvcc.

Also, tests?

I'll add a test to verify that "-include cuda_runtime.h" shows up on cc1 command line where/when it's expected.

What would be a good way to test the wrapper itself within clang tree without real CUDA headers?

I've done fair amount of manual testing outside of clang source tree.

• manual comparison of preprocessed output from cuda_runtime.h between host and device passes.
• compiled 39 out of 46 thrust examples and verified that they produce output identical to nvcc-compiled binaries.

[echristo]

In D13171#272441, @tra wrote:

In D13171#272397, @echristo wrote:

I'm ignoring the content of the header, but this seems to be a not terrible way to do things. I gather that cuda_runtime.h is something that's typically included by the driver by nvidia and not the client?

Correct. cuda_runtime.h (and all it pulls in) is -include'd under the hood by nvcc.

Also, tests?

I'll add a test to verify that "-include cuda_runtime.h" shows up on cc1 command line where/when it's expected.

Ick.

What would be a good way to test the wrapper itself within clang tree without real CUDA headers?

Hrm. Maybe a set of inputs that stub out things? Hard really.

I've done fair amount of manual testing outside of clang source tree.

• manual comparison of preprocessed output from cuda_runtime.h between host and device passes.
• compiled 39 out of 46 thrust examples and verified that they produce output identical to nvcc-compiled binaries.

Cool.

LGTM with those changes and give a thought at how to test this in tree better.

Thanks!

-eric

[tra]

Added test cases for force-including of cuda_runtime.h
Tweaked inclusion of one header due to use of default arguments.

lib/Headers/CMakeLists.txt

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

lib/Headers/clang_cuda_support.h

+#ifndef __CLANG_CUDA_SUPPORT_H__
+#define __CLANG_CUDA_SUPPORT_H__
+
+#if defined(__PTX__)
+
+// WARNING: Preprocessor hacks below are based on specific of
+// implementation of CUDA-7.0 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
+#error "Unsupported CUDA version!"
+#endif
+
+#define __NVCC__ 1
+#if defined(__CUDA_ARCH__)
+#define __CUDABE__ 1
+#else
+#define __CUDACC__ 1
+#endif
+
+// Fake include guards to prevent inclusion of some CUDA headers.
+#define __HOST_DEFINES_H__
+#define __DEVICE_LAUNCH_PARAMETERS_H__
+#define __TEXTURE_INDIRECT_FUNCTIONS_HPP__
+#define __SURFACE_INDIRECT_FUNCTIONS_HPP__
+
+// Standard CUDA attributes
+#define __constant__ __attribute__((constant))
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+#define __host__ __attribute__((host))
+#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
+#define __shared__ __attribute__((shared))
+
+// Additional macros used throughout CUDA headers.
+#define __align__(x) __attribute__((aligned(x)))
+#define __builtin_align__(x) __align__(x)
+#define __cudart_builtin__
+#define __device_builtin__
+#define __forceinline__ __inline__ __attribute__((always_inline))
+
+#define CUDARTAPI
+#define _CRTIMP
+
+// Texture and surface types are not supported yet.
+#define __device_builtin_surface_type__
+#define __device_builtin_texture_type__
+
+// Include support for built-in variables.
+#include "cuda_builtin_vars.h"
+
+// CUDA headers were implemented with the assumption of split-mode
+// compilation and present CUDA functions differently for host and
+// device mode. Typically in host mode they provide declarations with
+// __device__ attribute attached. In device mode we get definitions
+// but *without* __device__ attribute. This does not work well in
+// combined compilation mode used by clang, so we have to trick CUDA
+// headers into something we can use.
+
+// libdevice functions in device_functions_decls.h either come with
+// __host__ __device__ attributes or with none at all. Temporarily
+// undefine __host__ so only __device__ is applied.
+#pragma push_macro("__CUDACC_RTC__")
+#pragma push_macro("__host__")
+#define __CUDACC_RTC__
+#define __host__
+#include "device_functions_decls.h"
+#pragma pop_macro("__host__")
+#pragma pop_macro("__CUDACC_RTC__")
+
+#include "cuda_runtime.h"
+#include "crt/device_runtime.h"
+
+#if defined(__CUDA_ARCH__)
+// 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.h"
+#include "math_functions.h"
+#pragma pop_macro("__forceinline__")
+#else
+#include "device_functions.h"
+#include "math_functions.h"
+#endif
+
+#if defined(__CUDA_ARCH__)
+// Definitions for device specific functions are provided only if
+// __CUDACC__ is defined. Alas, they've already been transitively
+// included by device_functions.h and are now behind include guards.
+// We need to temporarily define __CUDACC__, undo include guards and
+// include the files with implementation of these functions.
+
+#pragma push_macro("__CUDACC__")
+#define __CUDACC__ 1
+
+#undef __DEVICE_ATOMIC_FUNCTIONS_HPP__
+#include "device_atomic_functions.hpp"
+
+#undef __SM_20_ATOMIC_FUNCTIONS_HPP__
+#include "sm_20_atomic_functions.hpp"
+#undef __SM_32_ATOMIC_FUNCTIONS_HPP__
+#include "sm_32_atomic_functions.hpp"
+
+#undef __SM_20_INTRINSICS_HPP__
+#include "sm_20_intrinsics.hpp"
+#undef __SM_30_INTRINSICS_HPP__
+#include "sm_30_intrinsics.hpp"
+#undef __SM_32_INTRINSICS_HPP__
+#include "sm_32_intrinsics.hpp"
+
+#pragma pop_macro("__CUDACC__")
+#endif // __CUDA_ARCH__
+#endif // __PTX__
+#endif // __CLANG_CUDA_SUPPORT_H__