sdk/build-tools/25.0.2/renderscript/clang-include/__clang_cuda_runtime_wrapper.h

   1 /*===---- __clang_cuda_runtime_wrapper.h - CUDA runtime support -------------===
   2  *
   3  * Permission is hereby granted, free of charge, to any person obtaining a copy
   4  * of this software and associated documentation files (the "Software"), to deal
   5  * in the Software without restriction, including without limitation the rights
   6  * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
   7  * copies of the Software, and to permit persons to whom the Software is
   8  * furnished to do so, subject to the following conditions:
   9  *
  10  * The above copyright notice and this permission notice shall be included in
  11  * all copies or substantial portions of the Software.
  12  *
  13  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
  14  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
  15  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
  16  * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
  17  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
  18  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
  19  * THE SOFTWARE.
  20  *
  21  *===-----------------------------------------------------------------------===
  22  */
  23
  24 /*
  25  * WARNING: This header is intended to be directly -include'd by
  26  * the compiler and is not supposed to be included by users.
  27  *
  28  * CUDA headers are implemented in a way that currently makes it
  29  * impossible for user code to #include directly when compiling with
  30  * Clang. They present different view of CUDA-supplied functions
  31  * depending on where in NVCC's compilation pipeline the headers are
  32  * included. Neither of these modes provides function definitions with
  33  * correct attributes, so we use preprocessor to force the headers
  34  * into a form that Clang can use.
  35  *
  36  * Similarly to NVCC which -include's cuda_runtime.h, Clang -include's
  37  * this file during every CUDA compilation.
  38  */
  39
  40 #ifndef __CLANG_CUDA_RUNTIME_WRAPPER_H__
  41 #define __CLANG_CUDA_RUNTIME_WRAPPER_H__
  42
  43 #if defined(__CUDA__) && defined(__clang__)
  44
  45 // Include some standard headers to avoid CUDA headers including them
  46 // while some required macros (like __THROW) are in a weird state.
  47 #include <stdlib.h>
  48 #include <cmath>
  49
  50 // Preserve common macros that will be changed below by us or by CUDA
  51 // headers.
  52 #pragma push_macro("__THROW")
  53 #pragma push_macro("__CUDA_ARCH__")
  54
  55 // WARNING: Preprocessor hacks below are based on specific details of
  56 // CUDA-7.x headers and are not expected to work with any other
  57 // version of CUDA headers.
  58 #include "cuda.h"
  59 #if !defined(CUDA_VERSION)
  60 #error "cuda.h did not define CUDA_VERSION"
  61 #elif CUDA_VERSION < 7000 || CUDA_VERSION > 7050
  62 #error "Unsupported CUDA version!"
  63 #endif
  64
  65 // Make largest subset of device functions available during host
  66 // compilation -- SM_35 for the time being.
  67 #ifndef __CUDA_ARCH__
  68 #define __CUDA_ARCH__ 350
  69 #endif
  70
  71 #include "cuda_builtin_vars.h"
  72
  73 // No need for device_launch_parameters.h as cuda_builtin_vars.h above
  74 // has taken care of builtin variables declared in the file.
  75 #define __DEVICE_LAUNCH_PARAMETERS_H__
  76
  77 // {math,device}_functions.h only have declarations of the
  78 // functions. We don't need them as we're going to pull in their
  79 // definitions from .hpp files.
  80 #define __DEVICE_FUNCTIONS_H__
  81 #define __MATH_FUNCTIONS_H__
  82
  83 #undef __CUDACC__
  84 #define __CUDABE__
  85 // Disables definitions of device-side runtime support stubs in
  86 // cuda_device_runtime_api.h
  87 #define __CUDADEVRT_INTERNAL__
  88 #include "host_config.h"
  89 #include "host_defines.h"
  90 #include "driver_types.h"
  91 #include "common_functions.h"
  92 #undef __CUDADEVRT_INTERNAL__
  93
  94 #undef __CUDABE__
  95 #define __CUDACC__
  96 #include "cuda_runtime.h"
  97
  98 #undef __CUDACC__
  99 #define __CUDABE__
 100
 101 // CUDA headers use __nvvm_memcpy and __nvvm_memset which Clang does
 102 // not have at the moment. Emulate them with a builtin memcpy/memset.
 103 #define __nvvm_memcpy(s,d,n,a) __builtin_memcpy(s,d,n)
 104 #define __nvvm_memset(d,c,n,a) __builtin_memset(d,c,n)
 105
 106 #include "crt/host_runtime.h"
 107 #include "crt/device_runtime.h"
 108 // device_runtime.h defines __cxa_* macros that will conflict with
 109 // cxxabi.h.
 110 // FIXME: redefine these as __device__ functions.
 111 #undef __cxa_vec_ctor
 112 #undef __cxa_vec_cctor
 113 #undef __cxa_vec_dtor
 114 #undef __cxa_vec_new2
 115 #undef __cxa_vec_new3
 116 #undef __cxa_vec_delete2
 117 #undef __cxa_vec_delete
 118 #undef __cxa_vec_delete3
 119 #undef __cxa_pure_virtual
 120
 121 // We need decls for functions in CUDA's libdevice with __device__
 122 // attribute only. Alas they come either as __host__ __device__ or
 123 // with no attributes at all. To work around that, define __CUDA_RTC__
 124 // which produces HD variant and undef __host__ which gives us desided
 125 // decls with __device__ attribute.
 126 #pragma push_macro("__host__")
 127 #define __host__
 128 #define __CUDACC_RTC__
 129 #include "device_functions_decls.h"
 130 #undef __CUDACC_RTC__
 131
 132 // Temporarily poison __host__ macro to ensure it's not used by any of
 133 // the headers we're about to include.
 134 #define __host__ UNEXPECTED_HOST_ATTRIBUTE
 135
 136 // device_functions.hpp and math_functions*.hpp use 'static
 137 // __forceinline__' (with no __device__) for definitions of device
 138 // functions. Temporarily redefine __forceinline__ to include
 139 // __device__.
 140 #pragma push_macro("__forceinline__")
 141 #define __forceinline__ __device__ __inline__ __attribute__((always_inline))
 142 #include "device_functions.hpp"
 143 #include "math_functions.hpp"
 144 #include "math_functions_dbl_ptx3.hpp"
 145 #pragma pop_macro("__forceinline__")
 146
 147 // Pull in host-only functions that are only available when neither
 148 // __CUDACC__ nor __CUDABE__ are defined.
 149 #undef __MATH_FUNCTIONS_HPP__
 150 #undef __CUDABE__
 151 #include "math_functions.hpp"
 152 // Alas, additional overloads for these functions are hard to get to.
 153 // Considering that we only need these overloads for a few functions,
 154 // we can provide them here.
 155 static inline float rsqrt(float a) { return rsqrtf(a); }
 156 static inline float rcbrt(float a) { return rcbrtf(a); }
 157 static inline float sinpi(float a) { return sinpif(a); }
 158 static inline float cospi(float a) { return cospif(a); }
 159 static inline void sincospi(float a, float *b, float *c) {
 160   return sincospi(a, b, c);
 161 }
 162 static inline float erfcinv(float a) { return erfcinvf(a); }
 163 static inline float normcdfinv(float a) { return normcdfinvf(a); }
 164 static inline float normcdf(float a) { return normcdff(a); }
 165 static inline float erfcx(float a) { return erfcxf(a); }
 166
 167 // For some reason single-argument variant is not always declared by
 168 // CUDA headers. Alas, device_functions.hpp included below needs it.
 169 static inline __device__ void __brkpt(int c) { __brkpt(); }
 170
 171 // Now include *.hpp with definitions of various GPU functions.  Alas,
 172 // a lot of thins get declared/defined with __host__ attribute which
 173 // we don't want and we have to define it out. We also have to include
 174 // {device,math}_functions.hpp again in order to extract the other
 175 // branch of #if/else inside.
 176
 177 #define __host__
 178 #undef __CUDABE__
 179 #define __CUDACC__
 180 #undef __DEVICE_FUNCTIONS_HPP__
 181 #include "device_functions.hpp"
 182 #include "device_atomic_functions.hpp"
 183 #include "sm_20_atomic_functions.hpp"
 184 #include "sm_32_atomic_functions.hpp"
 185 #include "sm_20_intrinsics.hpp"
 186 // sm_30_intrinsics.h has declarations that use default argument, so
 187 // we have to include it and it will in turn include .hpp
 188 #include "sm_30_intrinsics.h"
 189 #include "sm_32_intrinsics.hpp"
 190 #undef __MATH_FUNCTIONS_HPP__
 191 #include "math_functions.hpp"
 192 #pragma pop_macro("__host__")
 193
 194 #include "texture_indirect_functions.h"
 195
 196 // Restore state of __CUDA_ARCH__ and __THROW we had on entry.
 197 #pragma pop_macro("__CUDA_ARCH__")
 198 #pragma pop_macro("__THROW")
 199
 200 // Set up compiler macros expected to be seen during compilation.
 201 #undef __CUDABE__
 202 #define __CUDACC__
 203 #define __NVCC__
 204
 205 #if defined(__CUDA_ARCH__)
 206 // We need to emit IR declaration for non-existing __nvvm_reflect() to
 207 // let backend know that it should be treated as const nothrow
 208 // function which is what NVVMReflect pass expects to see.
 209 extern "C" __device__ __attribute__((const)) int __nvvm_reflect(const void *);
 210 static __device__ __attribute__((used)) int __nvvm_reflect_anchor() {
 211   return __nvvm_reflect("NONE");
 212 }
 213 #endif
 214
 215 #endif // __CUDA__
 216 #endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__