2015-12-17 02:51:59 +08:00
|
|
|
/*===---- __clang_cuda_runtime_wrapper.h - CUDA runtime support -------------===
|
2015-11-18 06:28:52 +08:00
|
|
|
*
|
|
|
|
* 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.
|
|
|
|
*
|
|
|
|
*===-----------------------------------------------------------------------===
|
|
|
|
*/
|
|
|
|
|
2015-12-17 02:51:59 +08:00
|
|
|
/*
|
|
|
|
* 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__
|
2015-11-18 06:28:52 +08:00
|
|
|
|
|
|
|
#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.
|
2015-12-18 06:25:22 +08:00
|
|
|
#include <cmath>
|
2016-01-27 07:37:29 +08:00
|
|
|
#include <cstdlib>
|
2016-02-12 10:22:53 +08:00
|
|
|
#include <stdlib.h>
|
2015-11-18 06:28:52 +08:00
|
|
|
|
|
|
|
// Preserve common macros that will be changed below by us or by CUDA
|
|
|
|
// headers.
|
|
|
|
#pragma push_macro("__THROW")
|
|
|
|
#pragma push_macro("__CUDA_ARCH__")
|
|
|
|
|
2015-12-17 02:51:59 +08:00
|
|
|
// 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.
|
2015-11-18 06:28:52 +08:00
|
|
|
#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__
|
2016-02-04 04:53:58 +08:00
|
|
|
#define __COMMON_FUNCTIONS_H__
|
2015-11-18 06:28:52 +08:00
|
|
|
|
|
|
|
#undef __CUDACC__
|
|
|
|
#define __CUDABE__
|
|
|
|
// Disables definitions of device-side runtime support stubs in
|
|
|
|
// cuda_device_runtime_api.h
|
2016-02-12 10:22:53 +08:00
|
|
|
#include "driver_types.h"
|
2015-11-18 06:28:52 +08:00
|
|
|
#include "host_config.h"
|
|
|
|
#include "host_defines.h"
|
|
|
|
|
|
|
|
#undef __CUDABE__
|
|
|
|
#define __CUDACC__
|
2015-12-17 02:51:59 +08:00
|
|
|
#include "cuda_runtime.h"
|
2015-11-18 06:28:52 +08:00
|
|
|
|
|
|
|
#undef __CUDACC__
|
|
|
|
#define __CUDABE__
|
|
|
|
|
2015-12-17 02:51:59 +08:00
|
|
|
// CUDA headers use __nvvm_memcpy and __nvvm_memset which Clang does
|
2015-11-18 06:28:52 +08:00
|
|
|
// not have at the moment. Emulate them with a builtin memcpy/memset.
|
2016-02-12 10:22:53 +08:00
|
|
|
#define __nvvm_memcpy(s, d, n, a) __builtin_memcpy(s, d, n)
|
|
|
|
#define __nvvm_memset(d, c, n, a) __builtin_memset(d, c, n)
|
2015-11-18 06:28:52 +08:00
|
|
|
|
|
|
|
#include "crt/device_runtime.h"
|
2016-02-12 10:22:53 +08:00
|
|
|
#include "crt/host_runtime.h"
|
2015-11-18 06:28:52 +08:00
|
|
|
// 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
|
|
|
|
|
2015-12-18 06:25:22 +08:00
|
|
|
// We need decls for functions in CUDA's libdevice with __device__
|
2015-11-18 06:28:52 +08:00
|
|
|
// 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__")
|
|
|
|
|
2015-12-18 06:25:22 +08:00
|
|
|
// Pull in host-only functions that are only available when neither
|
|
|
|
// __CUDACC__ nor __CUDABE__ are defined.
|
|
|
|
#undef __MATH_FUNCTIONS_HPP__
|
|
|
|
#undef __CUDABE__
|
|
|
|
#include "math_functions.hpp"
|
|
|
|
// 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.
|
2016-02-12 10:22:53 +08:00
|
|
|
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);
|
2015-12-18 06:25:22 +08:00
|
|
|
}
|
2016-02-12 10:22:53 +08:00
|
|
|
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); }
|
2015-12-18 06:25:22 +08:00
|
|
|
|
2015-11-18 06:28:52 +08:00
|
|
|
// For some reason single-argument variant is not always declared by
|
|
|
|
// CUDA headers. Alas, device_functions.hpp included below needs it.
|
2016-02-12 10:22:53 +08:00
|
|
|
static inline __device__ void __brkpt(int __c) { __brkpt(); }
|
2015-11-18 06:28:52 +08:00
|
|
|
|
|
|
|
// 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_atomic_functions.hpp"
|
2016-02-12 10:22:53 +08:00
|
|
|
#include "device_functions.hpp"
|
2015-11-18 06:28:52 +08:00
|
|
|
#include "sm_20_atomic_functions.hpp"
|
|
|
|
#include "sm_20_intrinsics.hpp"
|
2016-02-12 10:22:53 +08:00
|
|
|
#include "sm_32_atomic_functions.hpp"
|
2015-11-18 06:28:52 +08:00
|
|
|
// 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__
|
|
|
|
|
|
|
|
#if defined(__CUDA_ARCH__)
|
2015-12-18 06:25:22 +08:00
|
|
|
// We need to emit IR declaration for non-existing __nvvm_reflect() to
|
2015-11-18 06:28:52 +08:00
|
|
|
// let backend know that it should be treated as const nothrow
|
2015-12-18 06:25:22 +08:00
|
|
|
// function which is what NVVMReflect pass expects to see.
|
2015-11-18 06:28:52 +08:00
|
|
|
extern "C" __device__ __attribute__((const)) int __nvvm_reflect(const void *);
|
|
|
|
static __device__ __attribute__((used)) int __nvvm_reflect_anchor() {
|
|
|
|
return __nvvm_reflect("NONE");
|
|
|
|
}
|
|
|
|
#endif
|
|
|
|
|
2016-02-04 04:53:58 +08:00
|
|
|
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.
|
2016-02-12 10:22:53 +08:00
|
|
|
__device__ int vprintf(const char *, const char *);
|
2016-02-04 04:53:58 +08:00
|
|
|
__device__ void free(void *) __attribute((nothrow));
|
|
|
|
__device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc));
|
2016-02-12 10:22:53 +08:00
|
|
|
__device__ void __assertfail(const char *__message, const char *__file,
|
|
|
|
unsigned __line, const char *__function,
|
|
|
|
size_t __charSize) __attribute__((noreturn));
|
2016-02-04 04:53:58 +08:00
|
|
|
|
|
|
|
// In order for standard assert() macro on linux to work we need to
|
|
|
|
// provide device-side __assert_fail()
|
2016-02-12 10:22:53 +08:00
|
|
|
__device__ static inline void __assert_fail(const char *__message,
|
|
|
|
const char *__file, unsigned __line,
|
|
|
|
const char *__function) {
|
|
|
|
__assertfail(__message, __file, __line, __function, sizeof(char));
|
2016-02-04 04:53:58 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
// 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
|
|
|
|
|
2016-02-25 05:49:33 +08:00
|
|
|
// Out-of-line implementations from 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 uint3() const {
|
|
|
|
uint3 ret;
|
|
|
|
ret.x = x;
|
|
|
|
ret.y = y;
|
|
|
|
ret.z = z;
|
|
|
|
return ret;
|
|
|
|
}
|
|
|
|
|
|
|
|
__device__ inline __cuda_builtin_blockIdx_t::operator uint3() const {
|
|
|
|
uint3 ret;
|
|
|
|
ret.x = x;
|
|
|
|
ret.y = y;
|
|
|
|
ret.z = z;
|
|
|
|
return ret;
|
|
|
|
}
|
|
|
|
|
|
|
|
__device__ inline __cuda_builtin_blockDim_t::operator dim3() const {
|
|
|
|
return dim3(x, y, z);
|
|
|
|
}
|
|
|
|
|
|
|
|
__device__ inline __cuda_builtin_gridDim_t::operator dim3() const {
|
|
|
|
return dim3(x, y, z);
|
|
|
|
}
|
|
|
|
|
2016-01-27 07:37:29 +08:00
|
|
|
#include <__clang_cuda_cmath.h>
|
|
|
|
|
2016-02-25 05:49:31 +08:00
|
|
|
// 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 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")
|
|
|
|
|
2015-11-18 06:28:52 +08:00
|
|
|
#endif // __CUDA__
|
2015-12-17 02:51:59 +08:00
|
|
|
#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__
|