Subversion Repositories QNX 8.QNX8 LLVM/Clang compiler suite

/*===---- __clang_cuda_runtime_wrapper.h - CUDA runtime support -------------===

 *

 * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.

 * See https://llvm.org/LICENSE.txt for license information.

 * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception

 *

 *===-----------------------------------------------------------------------===

 */

/*

 * 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__)

// Include some forward declares that must come before cmath.

#include <__clang_cuda_math_forward_declares.h>

// Define __CUDACC__ early as libstdc++ standard headers with GNU extensions

// enabled depend on it to avoid using __float128, which is unsupported in

// CUDA.

#define __CUDACC__

// Include some standard headers to avoid CUDA headers including them

// while some required macros (like __THROW) are in a weird state.

#include <cmath>

#include <cstdlib>

#include <stdlib.h>

#include <string.h>

#undef __CUDACC__

// 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 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"

#elif CUDA_VERSION < 7000

#error "Unsupported CUDA version!"

#endif

#pragma push_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__")

#if CUDA_VERSION >= 10000

#define __CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__

#endif

// Make largest subset of device functions available during host

// compilation.

#ifndef __CUDA_ARCH__

#define __CUDA_ARCH__ 9999

#endif

#include "__clang_cuda_builtin_vars.h"

// No need for device_launch_parameters.h as __clang_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__

#define __COMMON_FUNCTIONS_H__

// device_functions_decls is replaced by __clang_cuda_device_functions.h

// included below.

#define __DEVICE_FUNCTIONS_DECLS_H__

#undef __CUDACC__

#if CUDA_VERSION < 9000

#define __CUDABE__

#else

#define __CUDACC__

#define __CUDA_LIBDEVICE__

#endif

// Disables definitions of device-side runtime support stubs in

// cuda_device_runtime_api.h

#include "host_defines.h"

#undef __CUDACC__

#include "driver_types.h"

#include "host_config.h"

// Temporarily replace "nv_weak" with weak, so __attribute__((nv_weak)) in

// cuda_device_runtime_api.h ends up being __attribute__((weak)) which is the

// functional equivalent of what we need.

#pragma push_macro("nv_weak")

#define nv_weak weak

#undef __CUDABE__

#undef __CUDA_LIBDEVICE__

#define __CUDACC__

#include "cuda_runtime.h"

#pragma pop_macro("nv_weak")

#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)

#if CUDA_VERSION < 9000

#include "crt/device_runtime.h"

#endif

#include "crt/host_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_new

#undef __cxa_vec_new2

#undef __cxa_vec_new3

#undef __cxa_vec_delete2

#undef __cxa_vec_delete

#undef __cxa_vec_delete3

#undef __cxa_pure_virtual

// math_functions.hpp expects this host function be defined on MacOS, but it

// ends up not being there because of the games we play here.  Just define it

// ourselves; it's simple enough.

#ifdef __APPLE__

inline __host__ double __signbitd(double x) {

  return std::signbit(x);

}

#endif

// CUDA 9.1 no longer provides declarations for libdevice functions, so we need

// to provide our own.

#include <__clang_cuda_libdevice_declares.h>

// Wrappers for many device-side standard library functions, incl. math

// functions, became compiler builtins in CUDA-9 and have been removed from the

// CUDA headers. Clang now provides its own implementation of the wrappers.

#if CUDA_VERSION >= 9000

#include <__clang_cuda_device_functions.h>

#include <__clang_cuda_math.h>

#endif

// __THROW is redefined to be empty by device_functions_decls.h in CUDA. Clang's

// counterpart does not do it, so we need to make it empty here to keep

// following CUDA includes happy.

#undef __THROW

#define __THROW

// CUDA 8.0.41 relies on __USE_FAST_MATH__ and __CUDA_PREC_DIV's values.

// Previous versions used to check whether they are defined or not.

// CU_DEVICE_INVALID macro is only defined in 8.0.41, so we use it

// here to detect the switch.

#if defined(CU_DEVICE_INVALID)

#if !defined(__USE_FAST_MATH__)

#define __USE_FAST_MATH__ 0

#endif

#if !defined(__CUDA_PREC_DIV)

#define __CUDA_PREC_DIV 0

#endif

#endif

// Temporarily poison __host__ macro to ensure it's not used by any of

// the headers we're about to include.

#pragma push_macro("__host__")

#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))

#if CUDA_VERSION < 9000

#include "device_functions.hpp"

#endif

// math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we

// get the slow-but-accurate or fast-but-inaccurate versions of functions like

// sin and exp.  This is controlled in clang by -fcuda-approx-transcendentals.

//

// device_functions.hpp uses __USE_FAST_MATH__ for a different purpose (fast vs.

// slow divides), so we need to scope our define carefully here.

#pragma push_macro("__USE_FAST_MATH__")

#if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__)

#define __USE_FAST_MATH__ 1

#endif

#if CUDA_VERSION >= 9000

#include "crt/math_functions.hpp"

#else

#include "math_functions.hpp"

#endif

#pragma pop_macro("__USE_FAST_MATH__")

#if CUDA_VERSION < 9000

#include "math_functions_dbl_ptx3.hpp"

#endif

#pragma pop_macro("__forceinline__")

// Pull in host-only functions that are only available when neither

// __CUDACC__ nor __CUDABE__ are defined.

#undef __MATH_FUNCTIONS_HPP__

#undef __CUDABE__

#if CUDA_VERSION < 9000

#include "math_functions.hpp"

#endif

// Alas, additional overloads for these functions are hard to get to.

// Considering that we only need these overloads for a few functions,

// we can provide them here.

static inline float rsqrt(float __a) { return rsqrtf(__a); }

static inline float rcbrt(float __a) { return rcbrtf(__a); }

static inline float sinpi(float __a) { return sinpif(__a); }

static inline float cospi(float __a) { return cospif(__a); }

static inline void sincospi(float __a, float *__b, float *__c) {

  return sincospif(__a, __b, __c);

}

static inline float erfcinv(float __a) { return erfcinvf(__a); }

static inline float normcdfinv(float __a) { return normcdfinvf(__a); }

static inline float normcdf(float __a) { return normcdff(__a); }

static inline float erfcx(float __a) { return erfcxf(__a); }

#if CUDA_VERSION < 9000

// 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(); }

#endif

// 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__

#if CUDA_VERSION >= 9000

// Some atomic functions became compiler builtins in CUDA-9 , so we need their

// declarations.

#include "device_atomic_functions.h"

#endif

#undef __DEVICE_FUNCTIONS_HPP__

#include "device_atomic_functions.hpp"

#if CUDA_VERSION >= 9000

#include "crt/device_functions.hpp"

#include "crt/device_double_functions.hpp"

#else

#include "device_functions.hpp"

#define __CUDABE__

#include "device_double_functions.h"

#undef __CUDABE__

#endif

#include "sm_20_atomic_functions.hpp"

// Predicate functions used in `__builtin_assume` need to have no side effect.

// However, sm_20_intrinsics.hpp doesn't define them with neither pure nor

// const attribute. Rename definitions from sm_20_intrinsics.hpp and re-define

// them as pure ones.

#pragma push_macro("__isGlobal")

#pragma push_macro("__isShared")

#pragma push_macro("__isConstant")

#pragma push_macro("__isLocal")

#define __isGlobal __ignored_cuda___isGlobal

#define __isShared __ignored_cuda___isShared

#define __isConstant __ignored_cuda___isConstant

#define __isLocal __ignored_cuda___isLocal

#include "sm_20_intrinsics.hpp"

#pragma pop_macro("__isGlobal")

#pragma pop_macro("__isShared")

#pragma pop_macro("__isConstant")

#pragma pop_macro("__isLocal")

#pragma push_macro("__DEVICE__")

#define __DEVICE__ static __device__ __forceinline__ __attribute__((const))

__DEVICE__ unsigned int __isGlobal(const void *p) {

  return __nvvm_isspacep_global(p);

}

__DEVICE__ unsigned int __isShared(const void *p) {

  return __nvvm_isspacep_shared(p);

}

__DEVICE__ unsigned int __isConstant(const void *p) {

  return __nvvm_isspacep_const(p);

}

__DEVICE__ unsigned int __isLocal(const void *p) {

  return __nvvm_isspacep_local(p);

}

#pragma pop_macro("__DEVICE__")

#include "sm_32_atomic_functions.hpp"

// Don't include sm_30_intrinsics.h and sm_32_intrinsics.h.  These define the

// __shfl and __ldg intrinsics using inline (volatile) asm, but we want to

// define them using builtins so that the optimizer can reason about and across

// these instructions.  In particular, using intrinsics for ldg gets us the

// [addr+imm] addressing mode, which, although it doesn't actually exist in the

// hardware, seems to generate faster machine code because ptxas can more easily

// reason about our code.

#if CUDA_VERSION >= 8000

#pragma push_macro("__CUDA_ARCH__")

#undef __CUDA_ARCH__

#include "sm_60_atomic_functions.hpp"

#include "sm_61_intrinsics.hpp"

#pragma pop_macro("__CUDA_ARCH__")

#endif

#undef __MATH_FUNCTIONS_HPP__

// math_functions.hpp defines ::signbit as a __host__ __device__ function.  This

// conflicts with libstdc++'s constexpr ::signbit, so we have to rename

// math_function.hpp's ::signbit.  It's guarded by #undef signbit, but that's

// conditional on __GNUC__.  :)

#pragma push_macro("signbit")

#pragma push_macro("__GNUC__")

#undef __GNUC__

#define signbit __ignored_cuda_signbit

// CUDA-9 omits device-side definitions of some math functions if it sees

// include guard from math.h wrapper from libstdc++. We have to undo the header

// guard temporarily to get the definitions we need.

#pragma push_macro("_GLIBCXX_MATH_H")

#pragma push_macro("_LIBCPP_VERSION")

#if CUDA_VERSION >= 9000

#undef _GLIBCXX_MATH_H

// We also need to undo another guard that checks for libc++ 3.8+

#ifdef _LIBCPP_VERSION

#define _LIBCPP_VERSION 3700

#endif

#endif

#if CUDA_VERSION >= 9000

#include "crt/math_functions.hpp"

#else

#include "math_functions.hpp"

#endif

#pragma pop_macro("_GLIBCXX_MATH_H")

#pragma pop_macro("_LIBCPP_VERSION")

#pragma pop_macro("__GNUC__")

#pragma pop_macro("signbit")

#pragma pop_macro("__host__")

// __clang_cuda_texture_intrinsics.h must be included first in order to provide

// implementation for __nv_tex_surf_handler that CUDA's headers depend on.

// The implementation requires c++11 and only works with CUDA-9 or newer.

#if __cplusplus >= 201103L && CUDA_VERSION >= 9000

// clang-format off

#include <__clang_cuda_texture_intrinsics.h>

// clang-format on

#else

#if CUDA_VERSION >= 9000

// Provide a hint that texture support needs C++11.

template <typename T> struct __nv_tex_needs_cxx11 {

  const static bool value = false;

};

template <class T>

__host__ __device__ void __nv_tex_surf_handler(const char *name, T *ptr,

                                               cudaTextureObject_t obj,

                                               float x) {

  _Static_assert(__nv_tex_needs_cxx11<T>::value,

                 "Texture support requires C++11");

}

#else

// Textures in CUDA-8 and older are not supported by clang.There's no

// convenient way to intercept texture use in these versions, so we can't

// produce a meaningful error. The source code that attempts to use textures

// will continue to fail as it does now.

#endif // CUDA_VERSION

#endif // __cplusplus >= 201103L && CUDA_VERSION >= 9000

#include "texture_fetch_functions.h"

#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__

extern "C" {

// Device-side CUDA system calls.

// http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls

// We need these declarations and wrappers for device-side

// malloc/free/printf calls to work without relying on

// -fcuda-disable-target-call-checks option.

__device__ int vprintf(const char *, const char *);

__device__ void free(void *) __attribute((nothrow));

__device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc));

// __assertfail() used to have a `noreturn` attribute. Unfortunately that

// contributed to triggering the longstanding bug in ptxas when assert was used

// in sufficiently convoluted code. See

// https://bugs.llvm.org/show_bug.cgi?id=27738 for the details.

__device__ void __assertfail(const char *__message, const char *__file,

                             unsigned __line, const char *__function,

                             size_t __charSize);

// In order for standard assert() macro on linux to work we need to

// provide device-side __assert_fail()

__device__ static inline void __assert_fail(const char *__message,

                                            const char *__file, unsigned __line,

                                            const char *__function) {

  __assertfail(__message, __file, __line, __function, sizeof(char));

}

// Clang will convert printf into vprintf, but we still need

// device-side declaration for it.

__device__ int printf(const char *, ...);

} // extern "C"

// We also need device-side std::malloc and std::free.

namespace std {

__device__ static inline void free(void *__ptr) { ::free(__ptr); }

__device__ static inline void *malloc(size_t __size) {

  return ::malloc(__size);

}

} // namespace std

// Out-of-line implementations from __clang_cuda_builtin_vars.h.  These need to

// come after we've pulled in the definition of uint3 and dim3.

__device__ inline __cuda_builtin_threadIdx_t::operator dim3() const {

  return dim3(x, y, z);

}

__device__ inline __cuda_builtin_threadIdx_t::operator uint3() const {

  return {x, y, z};

}

__device__ inline __cuda_builtin_blockIdx_t::operator dim3() const {

  return dim3(x, y, z);

}

__device__ inline __cuda_builtin_blockIdx_t::operator uint3() const {

  return {x, y, z};

}

__device__ inline __cuda_builtin_blockDim_t::operator dim3() const {

  return dim3(x, y, z);

}

__device__ inline __cuda_builtin_blockDim_t::operator uint3() const {

  return {x, y, z};

}

__device__ inline __cuda_builtin_gridDim_t::operator dim3() const {

  return dim3(x, y, z);

}

__device__ inline __cuda_builtin_gridDim_t::operator uint3() const {

  return {x, y, z};

}

#include <__clang_cuda_cmath.h>

#include <__clang_cuda_intrinsics.h>

#include <__clang_cuda_complex_builtins.h>

// curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host

// mode, giving them their "proper" types of dim3 and uint3.  This is

// incompatible with the types we give in __clang_cuda_builtin_vars.h.  As as

// hack, force-include the header (nvcc doesn't include it by default) but

// redefine dim3 and uint3 to our builtin types.  (Thankfully dim3 and uint3 are

// only used here for the redeclarations of blockDim and threadIdx.)

#pragma push_macro("dim3")

#pragma push_macro("uint3")

#define dim3 __cuda_builtin_blockDim_t

#define uint3 __cuda_builtin_threadIdx_t

#include "curand_mtgp32_kernel.h"

#pragma pop_macro("dim3")

#pragma pop_macro("uint3")

#pragma pop_macro("__USE_FAST_MATH__")

#pragma pop_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__")

// CUDA runtime uses this undocumented function to access kernel launch

// configuration. The declaration is in crt/device_functions.h but that file

// includes a lot of other stuff we don't want. Instead, we'll provide our own

// declaration for it here.

#if CUDA_VERSION >= 9020

extern "C" unsigned __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim,

                                                size_t sharedMem = 0,

                                                void *stream = 0);

#endif

#endif // __CUDA__

#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__