ÿØÿà JFIF ÿÛ „ ( %!1!%*+...983,7(-.- PKSR]��sx""complexnu�[��/*===---- complex - CUDA wrapper for ------------------------------=== * * 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. * *===-----------------------------------------------------------------------=== */ #ifndef __CLANG_CUDA_WRAPPERS_COMPLEX #define __CLANG_CUDA_WRAPPERS_COMPLEX // Wrapper around that forces its functions to be __host__ // __device__. // First, include host-only headers we think are likely to be included by // , so that the pragma below only applies to itself. #if __cplusplus >= 201103L #include #endif #include #include #include // Next, include our wrapper, to ensure that device overloads of // std::min/max are available. #include #pragma clang force_cuda_host_device begin // When compiling for device, ask libstdc++ to use its own implements of // complex functions, rather than calling builtins (which resolve to library // functions that don't exist when compiling CUDA device code). // // This is a little dicey, because it causes libstdc++ to define a different // set of overloads on host and device. // // // Present only when compiling for host. // __host__ __device__ void complex sin(const complex& x) { // return __builtin_csinf(x); // } // // // Present when compiling for host and for device. // template // void __host__ __device__ complex sin(const complex& x) { // return complex(sin(x.real()) * cosh(x.imag()), // cos(x.real()), sinh(x.imag())); // } // // This is safe because when compiling for device, all function calls in // __host__ code to sin() will still resolve to *something*, even if they don't // resolve to the same function as they resolve to when compiling for host. We // don't care that they don't resolve to the right function because we won't // codegen this host code when compiling for device. #pragma push_macro("_GLIBCXX_USE_C99_COMPLEX") #pragma push_macro("_GLIBCXX_USE_C99_COMPLEX_TR1") #define _GLIBCXX_USE_C99_COMPLEX 0 #define _GLIBCXX_USE_C99_COMPLEX_TR1 0 // Work around a compatibility issue with libstdc++ 11.1.0 // https://bugs.llvm.org/show_bug.cgi?id=50383 #pragma push_macro("__failed_assertion") #if _GLIBCXX_RELEASE == 11 #define __failed_assertion __cuda_failed_assertion #endif #include_next #pragma pop_macro("__failed_assertion") #pragma pop_macro("_GLIBCXX_USE_C99_COMPLEX_TR1") #pragma pop_macro("_GLIBCXX_USE_C99_COMPLEX") #pragma clang force_cuda_host_device end #endif // include guard PKSR]Tzm�bbbits/basic_string.tccnu�[��// CUDA headers define __noinline__ which interferes with libstdc++'s use of // `__attribute((__noinline__))`. In order to avoid compilation error, // temporarily unset __noinline__ when we include affected libstdc++ header. #pragma push_macro("__noinline__") #undef __noinline__ #include_next "bits/basic_string.tcc" #pragma pop_macro("__noinline__") PKSR]-``bits/basic_string.hnu�[��// CUDA headers define __noinline__ which interferes with libstdc++'s use of // `__attribute((__noinline__))`. In order to avoid compilation error, // temporarily unset __noinline__ when we include affected libstdc++ header. #pragma push_macro("__noinline__") #undef __noinline__ #include_next "bits/basic_string.h" #pragma pop_macro("__noinline__") PKSR]{͸ccbits/shared_ptr_base.hnu�[��// CUDA headers define __noinline__ which interferes with libstdc++'s use of // `__attribute((__noinline__))`. In order to avoid compilation error, // temporarily unset __noinline__ when we include affected libstdc++ header. #pragma push_macro("__noinline__") #undef __noinline__ #include_next "bits/shared_ptr_base.h" #pragma pop_macro("__noinline__") PKSR]��bits/c++config.hnu�[��// libstdc++ uses the non-constexpr function std::__glibcxx_assert_fail() // to trigger compilation errors when the __glibcxx_assert(cond) macro // is used in a constexpr context. // Compilation fails when using code from the libstdc++ (such as std::array) on // device code, since these assertions invoke a non-constexpr host function from // device code. // // To work around this issue, we declare our own device version of the function #ifndef __CLANG_CUDA_WRAPPERS_BITS_CPP_CONFIG #define __CLANG_CUDA_WRAPPERS_BITS_CPP_CONFIG #include_next #ifdef _LIBCPP_BEGIN_NAMESPACE_STD _LIBCPP_BEGIN_NAMESPACE_STD #else namespace std { #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION _GLIBCXX_BEGIN_NAMESPACE_VERSION #endif #pragma push_macro("CUDA_NOEXCEPT") #if __cplusplus >= 201103L #define CUDA_NOEXCEPT noexcept #else #define CUDA_NOEXCEPT #endif __attribute__((device, noreturn)) inline void __glibcxx_assert_fail(const char *file, int line, const char *function, const char *condition) CUDA_NOEXCEPT { #ifdef _GLIBCXX_VERBOSE_ASSERT if (file && function && condition) __builtin_printf("%s:%d: %s: Assertion '%s' failed.\n", file, line, function, condition); else if (function) __builtin_printf("%s: Undefined behavior detected.\n", function); #endif __builtin_abort(); } #endif __attribute__((device, noreturn, __always_inline__, __visibility__("default"))) inline void __glibcxx_assert_fail() CUDA_NOEXCEPT { __builtin_abort(); } #pragma pop_macro("CUDA_NOEXCEPT") #ifdef _LIBCPP_END_NAMESPACE_STD _LIBCPP_END_NAMESPACE_STD #else #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION _GLIBCXX_END_NAMESPACE_VERSION #endif } // namespace std #endif #endif PKSR]�zR��newnu�[��/*===---- new - CUDA wrapper for -------------------------------------=== * * 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. * *===-----------------------------------------------------------------------=== */ #ifndef __CLANG_CUDA_WRAPPERS_NEW #define __CLANG_CUDA_WRAPPERS_NEW #include_next #if !defined(__device__) // The header has been included too early from the standard C++ library // and CUDA-specific macros are not available yet. // Undo the include guard and try again later. #undef __CLANG_CUDA_WRAPPERS_NEW #else #pragma push_macro("CUDA_NOEXCEPT") #if __cplusplus >= 201103L #define CUDA_NOEXCEPT noexcept #else #define CUDA_NOEXCEPT #endif // Device overrides for non-placement new and delete. __device__ inline void *operator new(__SIZE_TYPE__ size) { if (size == 0) { size = 1; } return ::malloc(size); } __device__ inline void *operator new(__SIZE_TYPE__ size, const std::nothrow_t &) CUDA_NOEXCEPT { return ::operator new(size); } __device__ inline void *operator new[](__SIZE_TYPE__ size) { return ::operator new(size); } __device__ inline void *operator new[](__SIZE_TYPE__ size, const std::nothrow_t &) { return ::operator new(size); } __device__ inline void operator delete(void* ptr) CUDA_NOEXCEPT { if (ptr) { ::free(ptr); } } __device__ inline void operator delete(void *ptr, const std::nothrow_t &) CUDA_NOEXCEPT { ::operator delete(ptr); } __device__ inline void operator delete[](void* ptr) CUDA_NOEXCEPT { ::operator delete(ptr); } __device__ inline void operator delete[](void *ptr, const std::nothrow_t &) CUDA_NOEXCEPT { ::operator delete(ptr); } // Sized delete, C++14 only. #if __cplusplus >= 201402L __device__ inline void operator delete(void *ptr, __SIZE_TYPE__ size) CUDA_NOEXCEPT { ::operator delete(ptr); } __device__ inline void operator delete[](void *ptr, __SIZE_TYPE__ size) CUDA_NOEXCEPT { ::operator delete(ptr); } #endif // Device overrides for placement new and delete. #if !(_LIBCPP_STD_VER >= 26 || __cpp_lib_constexpr_new >= 202406L) __device__ inline void *operator new(__SIZE_TYPE__, void *__ptr) CUDA_NOEXCEPT { return __ptr; } __device__ inline void *operator new[](__SIZE_TYPE__, void *__ptr) CUDA_NOEXCEPT { return __ptr; } #endif __device__ inline void operator delete(void *, void *) CUDA_NOEXCEPT {} __device__ inline void operator delete[](void *, void *) CUDA_NOEXCEPT {} #pragma pop_macro("CUDA_NOEXCEPT") #endif // __device__ #endif // include guard PKSR]L �&��cmathnu�[��/*===---- cmath - CUDA wrapper for ---------------------------------=== * * 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. * *===-----------------------------------------------------------------------=== */ #ifndef __CLANG_CUDA_WRAPPERS_CMATH #define __CLANG_CUDA_WRAPPERS_CMATH #if __has_include_next() #include_next #else #error "Could not find standard C++ header 'cmath'. Add -v to your compilation command to check the include paths being searched. You may need to install the appropriate standard C++ library package corresponding to the search path." #endif #if defined(_LIBCPP_STD_VER) // libc++ will need long double variants of these functions, but CUDA does not // provide them. We'll provide their declarations, which should allow the // headers to parse, but would not allow accidental use of them on a GPU. __attribute__((device)) long double logb(long double); __attribute__((device)) long double scalbn(long double, int); #endif // _LIBCPP_STD_VER #endif // include guard PKSR]*s ; algorithmnu�[��/*===---- algorithm - CUDA wrapper for -------------------------=== * * 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. * *===-----------------------------------------------------------------------=== */ #ifndef __CLANG_CUDA_WRAPPERS_ALGORITHM #define __CLANG_CUDA_WRAPPERS_ALGORITHM // This header defines __device__ overloads of std::min/max. // // Ideally we'd declare these functions only if we're <= C++11. In C++14, // these functions are constexpr, and so are implicitly __host__ __device__. // // However, the compiler being in C++14 mode does not imply that the standard // library supports C++14. There is no macro we can test to check that the // stdlib has constexpr std::min/max. Thus we have to unconditionally define // our device overloads. // // A host+device function cannot be overloaded, and a constexpr function // implicitly become host device if there's no explicitly host or device // overload preceding it. So the simple thing to do would be to declare our // device min/max overloads, and then #include_next . This way our // device overloads would come first, and so if we have a C++14 stdlib, its // min/max won't become host+device and conflict with our device overloads. // // But that also doesn't work. libstdc++ is evil and declares std::min/max in // an internal header that is included *before* . Thus by the time // we're inside of this file, std::min/max may already have been declared, and // thus we can't prevent them from becoming host+device if they're constexpr. // // Therefore we perpetrate the following hack: We mark our __device__ overloads // with __attribute__((enable_if(true, ""))). This causes the signature of the // function to change without changing anything else about it. (Except that // overload resolution will prefer it over the __host__ __device__ version // rather than considering them equally good). #include_next // We need to define these overloads in exactly the namespace our standard // library uses (including the right inline namespace), otherwise they won't be // picked up by other functions in the standard library (e.g. functions in // ). Thus the ugliness below. #ifdef _LIBCPP_BEGIN_NAMESPACE_STD _LIBCPP_BEGIN_NAMESPACE_STD #else namespace std { #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION _GLIBCXX_BEGIN_NAMESPACE_VERSION #endif #endif #pragma push_macro("_CPP14_CONSTEXPR") #if __cplusplus >= 201402L #define _CPP14_CONSTEXPR constexpr #else #define _CPP14_CONSTEXPR #endif template __attribute__((enable_if(true, ""))) inline _CPP14_CONSTEXPR __host__ __device__ const __T & max(const __T &__a, const __T &__b, __Cmp __cmp) { return __cmp(__a, __b) ? __b : __a; } template __attribute__((enable_if(true, ""))) inline _CPP14_CONSTEXPR __host__ __device__ const __T & max(const __T &__a, const __T &__b) { return __a < __b ? __b : __a; } template __attribute__((enable_if(true, ""))) inline _CPP14_CONSTEXPR __host__ __device__ const __T & min(const __T &__a, const __T &__b, __Cmp __cmp) { return __cmp(__b, __a) ? __b : __a; } template __attribute__((enable_if(true, ""))) inline _CPP14_CONSTEXPR __host__ __device__ const __T & min(const __T &__a, const __T &__b) { return __b < __a ? __b : __a; } #pragma pop_macro("_CPP14_CONSTEXPR") #ifdef _LIBCPP_END_NAMESPACE_STD _LIBCPP_END_NAMESPACE_STD #else #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION _GLIBCXX_END_NAMESPACE_VERSION #endif } // namespace std #endif #endif // __CLANG_CUDA_WRAPPERS_ALGORITHM PKSR]��sx""complexnu�[��PKSR]Tzm�bbYbits/basic_string.tccnu�[��PKSR]-``bits/basic_string.hnu�[��PKSR]{͸cc�bits/shared_ptr_base.hnu�[��PKSR]��Lbits/c++config.hnu�[��PKSR]�zR��Wnewnu�[��PKSR]L �&��%)cmathnu�[��PKSR]*s ; W1algorithmnu�[��PKf�C