diff --git a/sycl/CMakeLists.txt b/sycl/CMakeLists.txt index db7bccae6b1d6..d16d718442397 100644 --- a/sycl/CMakeLists.txt +++ b/sycl/CMakeLists.txt @@ -5,7 +5,6 @@ project(sycl-solution) set(CMAKE_CXX_STANDARD 11) set(CMAKE_CXX_STANDARD_REQUIRED ON) set(CMAKE_CXX_EXTENSIONS OFF) -set(CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS ON) option(SYCL_ENABLE_WERROR "Treat all warnings as errors in SYCL project" OFF) # enable all warnings by default diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index 1d758901d05f7..b93b72d9cf47b 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -9,6 +9,7 @@ #pragma once #include #include +#include #include #include #include @@ -256,16 +257,17 @@ OpGroupAsyncCopyLocalToGlobal(__spv::Scope::Flag Execution, dataT *Dest, return nullptr; } -extern void __spirv_ocl_prefetch(const char *Ptr, size_t NumBytes) noexcept; +extern __SYCL_EXPORT void __spirv_ocl_prefetch(const char *Ptr, + size_t NumBytes) noexcept; -__SYCL_CONVERGENT__ extern SYCL_EXTERNAL void +__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void __spirv_ControlBarrier(__spv::Scope Execution, __spv::Scope Memory, uint32_t Semantics) noexcept; -__SYCL_CONVERGENT__ extern SYCL_EXTERNAL void +__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept; -__SYCL_CONVERGENT__ extern SYCL_EXTERNAL void +__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void __spirv_GroupWaitEvents(__spv::Scope Execution, uint32_t NumEvents, __ocl_event_t *WaitEvents) noexcept; diff --git a/sycl/include/CL/sycl/context.hpp b/sycl/include/CL/sycl/context.hpp index 599ea43e4576f..1dfe69f7336ad 100644 --- a/sycl/include/CL/sycl/context.hpp +++ b/sycl/include/CL/sycl/context.hpp @@ -8,9 +8,11 @@ #pragma once #include +#include #include #include #include + #include // 4.6.2 Context class @@ -23,7 +25,7 @@ namespace detail { class context_impl; } -class context { +class __SYCL_EXPORT context { public: /// Constructs a SYCL context instance using an instance of default_selector. /// diff --git a/sycl/include/CL/sycl/detail/accessor_impl.hpp b/sycl/include/CL/sycl/detail/accessor_impl.hpp index 83bb2fadbed88..f9cffa5344b9f 100644 --- a/sycl/include/CL/sycl/detail/accessor_impl.hpp +++ b/sycl/include/CL/sycl/detail/accessor_impl.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include #include @@ -37,8 +38,7 @@ template class AccessorImplDevice { range MemRange; bool operator==(const AccessorImplDevice &Rhs) const { - return (Offset == Rhs.Offset && - AccessRange == Rhs.AccessRange && + return (Offset == Rhs.Offset && AccessRange == Rhs.AccessRange && MemRange == Rhs.MemRange); } }; @@ -59,7 +59,7 @@ template class LocalAccessorBaseDevice { } }; -class AccessorImplHost { +class __SYCL_EXPORT AccessorImplHost { public: AccessorImplHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange, access::mode AccessMode, detail::SYCLMemObjI *SYCLMemObject, @@ -129,7 +129,7 @@ class AccessorBaseHost { AccessorImplPtr impl; }; -class LocalAccessorImplHost { +class __SYCL_EXPORT LocalAccessorImplHost { public: LocalAccessorImplHost(sycl::range<3> Size, int Dims, int ElemSize) : MSize(Size), MDims(Dims), MElemSize(ElemSize), @@ -185,7 +185,7 @@ class LocalAccessorBaseHost { using Requirement = AccessorImplHost; -void addHostAccessorAndWait(Requirement *Req); +void __SYCL_EXPORT addHostAccessorAndWait(Requirement *Req); } // namespace detail } // namespace sycl diff --git a/sycl/include/CL/sycl/detail/buffer_impl.hpp b/sycl/include/CL/sycl/detail/buffer_impl.hpp index 54e0c2e5ccb76..b1561928218e5 100644 --- a/sycl/include/CL/sycl/detail/buffer_impl.hpp +++ b/sycl/include/CL/sycl/detail/buffer_impl.hpp @@ -12,6 +12,7 @@ #include #include #include +#include #include #include #include @@ -38,7 +39,7 @@ using buffer_allocator = detail::sycl_memory_object_allocator; namespace detail { -class buffer_impl final : public SYCLMemObjT { +class __SYCL_EXPORT buffer_impl final : public SYCLMemObjT { using BaseT = SYCLMemObjT; using typename BaseT::MemObjType; @@ -104,7 +105,7 @@ class buffer_impl final : public SYCLMemObjT { std::move(Allocator)) {} void *allocateMem(ContextImplPtr Context, bool InitFromUserData, - void *HostPtr, RT::PiEvent &OutEventToWait) override; + void *HostPtr, RT::PiEvent &OutEventToWait) override; MemObjType getType() const override { return MemObjType::BUFFER; } diff --git a/sycl/include/CL/sycl/detail/cg.hpp b/sycl/include/CL/sycl/detail/cg.hpp index a87daa3e8e154..6795ff3eeca7b 100644 --- a/sycl/include/CL/sycl/detail/cg.hpp +++ b/sycl/include/CL/sycl/detail/cg.hpp @@ -10,6 +10,7 @@ #include #include +#include #include #include #include @@ -55,7 +56,7 @@ class interop_handler { private: cl_command_queue MQueue; std::vector MMemObjs; - cl_mem getMemImpl(detail::Requirement* Req) const; + __SYCL_EXPORT cl_mem getMemImpl(detail::Requirement *Req) const; }; namespace detail { diff --git a/sycl/include/CL/sycl/detail/common.hpp b/sycl/include/CL/sycl/detail/common.hpp index 141d1fffdd7b8..b1310a103c0c9 100644 --- a/sycl/include/CL/sycl/detail/common.hpp +++ b/sycl/include/CL/sycl/detail/common.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include // Suppress a compiler warning about undefined CL_TARGET_OPENCL_VERSION // Khronos ICD supports only latest OpenCL version @@ -84,7 +85,7 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { namespace detail { -const char *stringifyErrorCode(cl_int error); +__SYCL_EXPORT const char *stringifyErrorCode(cl_int error); static inline std::string codeToString(cl_int code) { return std::string(std::to_string(code) + " (" + stringifyErrorCode(code) + diff --git a/sycl/include/CL/sycl/detail/common_info.hpp b/sycl/include/CL/sycl/detail/common_info.hpp index bd05ec22b9db2..b7fe443cdd2b8 100644 --- a/sycl/include/CL/sycl/detail/common_info.hpp +++ b/sycl/include/CL/sycl/detail/common_info.hpp @@ -7,14 +7,15 @@ //===----------------------------------------------------------------------===// #pragma once +#include #include __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { namespace detail { -vector_class split_string(const string_class &str, - char delimeter); +vector_class __SYCL_EXPORT split_string(const string_class &str, + char delimeter); } // namespace detail } // namespace sycl diff --git a/sycl/include/CL/sycl/detail/defines.hpp b/sycl/include/CL/sycl/detail/defines.hpp index c778ce9603fc9..352b027890154 100644 --- a/sycl/include/CL/sycl/detail/defines.hpp +++ b/sycl/include/CL/sycl/detail/defines.hpp @@ -31,11 +31,3 @@ #ifndef SYCL_EXTERNAL #define SYCL_EXTERNAL #endif - -#if __cplusplus >= 201402 -#define __SYCL_DEPRECATED__(message) [[deprecated(message)]] -#elif !defined _MSC_VER -#define __SYCL_DEPRECATED__(message) __attribute__((deprecated(message))) -#else -#define __SYCL_DEPRECATED__(message) -#endif diff --git a/sycl/include/CL/sycl/detail/export.hpp b/sycl/include/CL/sycl/detail/export.hpp new file mode 100644 index 0000000000000..68b15af73456a --- /dev/null +++ b/sycl/include/CL/sycl/detail/export.hpp @@ -0,0 +1,38 @@ +//==---------------- export.hpp - SYCL standard header file ----------------==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#pragma once + +#ifndef SYCL_DEVICE_ONLY +#ifndef __SYCL_EXPORT +#ifdef _WIN32 + +// MSVC discourages export of classes, that use STL class in API. This +// results in a warning, treated as compile error. Silence C4251 to workaround. +#pragma warning(disable : 4251) +#pragma warning(disable : 4275) + +#define DLL_LOCAL + +#if __SYCL_BUILD_SYCL_DLL +#define __SYCL_EXPORT __declspec(dllexport) +#define __SYCL_EXPORT_DEPRECATED(x) __declspec(dllexport, deprecated(x)) +#else +#define __SYCL_EXPORT __declspec(dllimport) +#define __SYCL_EXPORT_DEPRECATED(x) __declspec(dllimport, deprecated(x)) +#endif +#else + +#define DLL_LOCAL __attribute__((visibility("hidden"))) + +#define __SYCL_EXPORT __attribute__((visibility("default"))) +#define __SYCL_EXPORT_DEPRECATED(x) \ + __attribute__((visibility("default"), deprecated(x))) +#endif +#endif +#endif diff --git a/sycl/include/CL/sycl/detail/generic_type_traits.hpp b/sycl/include/CL/sycl/detail/generic_type_traits.hpp index 78331225d20d5..eb8bfd53aa574 100644 --- a/sycl/include/CL/sycl/detail/generic_type_traits.hpp +++ b/sycl/include/CL/sycl/detail/generic_type_traits.hpp @@ -604,8 +604,7 @@ class is_same_vector_size_impl { public: static constexpr bool value = - IsSizeEqual ? is_same_vector_size_impl::value - : false; + IsSizeEqual ? is_same_vector_size_impl::value : false; }; template diff --git a/sycl/include/CL/sycl/detail/helpers.hpp b/sycl/include/CL/sycl/detail/helpers.hpp index d3a520b1645ab..024801274b006 100644 --- a/sycl/include/CL/sycl/detail/helpers.hpp +++ b/sycl/include/CL/sycl/detail/helpers.hpp @@ -12,6 +12,7 @@ #include #include #include +#include #include #include @@ -44,11 +45,11 @@ inline void memcpy(void *Dst, const void *Src, size_t Size) { class context_impl; // The function returns list of events that can be passed to OpenCL API as // dependency list and waits for others. -std::vector +__SYCL_EXPORT std::vector getOrWaitEvents(std::vector DepEvents, std::shared_ptr Context); -void waitEvents(std::vector DepEvents); +__SYCL_EXPORT void waitEvents(std::vector DepEvents); class Builder { public: diff --git a/sycl/include/CL/sycl/detail/host_profiling_info.hpp b/sycl/include/CL/sycl/detail/host_profiling_info.hpp index 8a98de238e10a..4aa33f7728b53 100644 --- a/sycl/include/CL/sycl/detail/host_profiling_info.hpp +++ b/sycl/include/CL/sycl/detail/host_profiling_info.hpp @@ -9,13 +9,14 @@ #pragma once #include +#include __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { namespace detail { /// Profiling info for the host execution. -class HostProfilingInfo { +class __SYCL_EXPORT HostProfilingInfo { cl_ulong StartTime = 0; cl_ulong EndTime = 0; diff --git a/sycl/include/CL/sycl/detail/image_accessor_util.hpp b/sycl/include/CL/sycl/detail/image_accessor_util.hpp index f0a4cefbe43b0..0d712cb73125f 100644 --- a/sycl/include/CL/sycl/detail/image_accessor_util.hpp +++ b/sycl/include/CL/sycl/detail/image_accessor_util.hpp @@ -13,6 +13,7 @@ #ifndef __SYCL_DEVICE_ONLY__ #include +#include #include #include #include @@ -99,21 +100,25 @@ getImageOffset(const vec &Coords, const id<3> ImgPitch, // Process cl_float4 Coordinates and return the appropriate Pixel Coordinates to // read from based on Addressing Mode for Nearest filter mode. -cl_int4 getPixelCoordNearestFiltMode(cl_float4, const addressing_mode, - const range<3>); +__SYCL_EXPORT cl_int4 getPixelCoordNearestFiltMode(cl_float4, + const addressing_mode, + const range<3>); // Process cl_float4 Coordinates and return the appropriate Pixel Coordinates to // read from based on Addressing Mode for Linear filter mode. -cl_int8 getPixelCoordLinearFiltMode(cl_float4, const addressing_mode, - const range<3>, cl_float4 &); +__SYCL_EXPORT cl_int8 getPixelCoordLinearFiltMode(cl_float4, + const addressing_mode, + const range<3>, cl_float4 &); // Check if PixelCoord are out of range for Sampler with clamp adressing mode. -bool isOutOfRange(const cl_int4 PixelCoord, const addressing_mode SmplAddrMode, - const range<3> ImgRange); +__SYCL_EXPORT bool isOutOfRange(const cl_int4 PixelCoord, + const addressing_mode SmplAddrMode, + const range<3> ImgRange); // Get Border Color for the image_channel_order, the border color values are // only used when the sampler has clamp addressing mode. -cl_float4 getBorderColor(const image_channel_order ImgChannelOrder); +__SYCL_EXPORT cl_float4 +getBorderColor(const image_channel_order ImgChannelOrder); // Reads data from a pixel at Ptr location, based on the number of Channels in // Order and returns the data. diff --git a/sycl/include/CL/sycl/detail/image_impl.hpp b/sycl/include/CL/sycl/detail/image_impl.hpp index 285227c2d43e7..2ff2f67b12902 100644 --- a/sycl/include/CL/sycl/detail/image_impl.hpp +++ b/sycl/include/CL/sycl/detail/image_impl.hpp @@ -10,6 +10,7 @@ #include #include +#include #include #include #include @@ -37,18 +38,23 @@ namespace detail { using image_allocator = aligned_allocator; // utility function: Returns the Number of Channels for a given Order. -uint8_t getImageNumberChannels(image_channel_order Order); +__SYCL_EXPORT uint8_t getImageNumberChannels(image_channel_order Order); // utility function: Returns the number of bytes per image element -uint8_t getImageElementSize(uint8_t NumChannels, image_channel_type Type); +__SYCL_EXPORT uint8_t getImageElementSize(uint8_t NumChannels, + image_channel_type Type); -RT::PiMemImageChannelOrder convertChannelOrder(image_channel_order Order); +__SYCL_EXPORT RT::PiMemImageChannelOrder +convertChannelOrder(image_channel_order Order); -image_channel_order convertChannelOrder(RT::PiMemImageChannelOrder Order); +__SYCL_EXPORT image_channel_order +convertChannelOrder(RT::PiMemImageChannelOrder Order); -RT::PiMemImageChannelType convertChannelType(image_channel_type Type); +__SYCL_EXPORT RT::PiMemImageChannelType +convertChannelType(image_channel_type Type); -image_channel_type convertChannelType(RT::PiMemImageChannelType Type); +__SYCL_EXPORT image_channel_type +convertChannelType(RT::PiMemImageChannelType Type); // validImageDataT: cl_int4, cl_uint4, cl_float4, cl_half4 template @@ -59,7 +65,8 @@ template using EnableIfImgAccDataT = typename std::enable_if::value, DataT>::type; -template class image_impl final : public SYCLMemObjT { +template +class __SYCL_EXPORT image_impl final : public SYCLMemObjT { using BaseT = SYCLMemObjT; using typename BaseT::MemObjType; diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp index 4bd64ed4e699e..7f1bcf496a55e 100644 --- a/sycl/include/CL/sycl/detail/kernel_desc.hpp +++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp @@ -9,7 +9,10 @@ #pragma once #include -#include // for DLL_LOCAL used in int. header +#include +#include // for DLL_LOCAL used in int. header + +#include __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { diff --git a/sycl/include/CL/sycl/detail/memory_manager.hpp b/sycl/include/CL/sycl/detail/memory_manager.hpp index d161ea283c584..d07010d9ed8c2 100644 --- a/sycl/include/CL/sycl/detail/memory_manager.hpp +++ b/sycl/include/CL/sycl/detail/memory_manager.hpp @@ -10,6 +10,7 @@ #include #include +#include #include #include @@ -31,7 +32,7 @@ using ContextImplPtr = std::shared_ptr; // The class contains methods that work with memory. All operations with // device memory should go through MemoryManager. -class MemoryManager { +class __SYCL_EXPORT MemoryManager { public: // The following method releases memory allocation of memory object. // Depending on the context it releases memory on host or on device. @@ -133,7 +134,6 @@ class MemoryManager { static void prefetch_usm(void *Ptr, QueueImplPtr Queue, size_t Len, std::vector DepEvents, RT::PiEvent &OutEvent); - }; } // namespace detail } // namespace sycl diff --git a/sycl/include/CL/sycl/detail/os_util.hpp b/sycl/include/CL/sycl/detail/os_util.hpp index 354998d86dca6..19eca9f377ef2 100644 --- a/sycl/include/CL/sycl/detail/os_util.hpp +++ b/sycl/include/CL/sycl/detail/os_util.hpp @@ -11,6 +11,7 @@ #pragma once #include +#include #include #include @@ -36,26 +37,6 @@ #error "Unsupported compiler or OS" #endif // _WIN32 -#if defined(SYCL_RT_OS_WINDOWS) - -#define DLL_LOCAL -// If SYCL headers are included to build SYCL library then the macro is used -// to set dllexport attribute for global variables/functions/classes. -// Otherwise, the macro is used used to set dllimport for the same global -// variables/functions/classes. -#if defined(__SYCL_BUILD_SYCL_DLL) -#define __SYCL_EXPORTED __declspec(dllexport) -#else -#define __SYCL_EXPORTED __declspec(dllimport) -#endif - -#elif defined(SYCL_RT_OS_POSIX_SUPPORT) - -#define DLL_LOCAL __attribute__((visibility("hidden"))) -#define __SYCL_EXPORTED - -#endif - __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { namespace detail { @@ -65,7 +46,7 @@ namespace detail { using OSModuleHandle = intptr_t; /// Groups the OS-dependent services. -class OSUtil { +class __SYCL_EXPORT OSUtil { public: /// Returns a module enclosing given address or nullptr. static OSModuleHandle getOSModuleHandle(const void *VirtAddr); @@ -74,7 +55,7 @@ class OSUtil { static std::string getCurrentDSODir(); /// Returns a directory component of a path. - static std::string getDirName(const char* Path); + static std::string getDirName(const char *Path); /// Module handle for the executable module - it is assumed there is always /// single one at most. @@ -85,9 +66,9 @@ class OSUtil { static constexpr OSModuleHandle DummyModuleHandle = -2; #ifdef SYCL_RT_OS_WINDOWS - static constexpr const char* DirSep = "\\"; + static constexpr const char *DirSep = "\\"; #else - static constexpr const char* DirSep = "/"; + static constexpr const char *DirSep = "/"; #endif /// Returns the amount of RAM available for the operating system. diff --git a/sycl/include/CL/sycl/detail/pi.hpp b/sycl/include/CL/sycl/detail/pi.hpp index 73ba98a4e4530..723a74f88acf8 100644 --- a/sycl/include/CL/sycl/detail/pi.hpp +++ b/sycl/include/CL/sycl/detail/pi.hpp @@ -14,11 +14,12 @@ #pragma once #include +#include #include #include -#include #include +#include #include #ifdef XPTI_ENABLE_INSTRUMENTATION @@ -48,9 +49,9 @@ namespace pi { #endif // Report error and no return (keeps compiler happy about no return statements). -[[noreturn]] void die(const char *Message); +[[noreturn]] __SYCL_EXPORT void die(const char *Message); -void assertion(bool Condition, const char *Message = nullptr); +__SYCL_EXPORT void assertion(bool Condition, const char *Message = nullptr); template void handleUnknownParamName(const char *functionName, T parameter) { diff --git a/sycl/include/CL/sycl/detail/sampler_impl.hpp b/sycl/include/CL/sycl/detail/sampler_impl.hpp index 36326d60a385a..1f3b013d78768 100644 --- a/sycl/include/CL/sycl/detail/sampler_impl.hpp +++ b/sycl/include/CL/sycl/detail/sampler_impl.hpp @@ -10,6 +10,7 @@ #include #include +#include #include @@ -21,7 +22,7 @@ enum class filtering_mode : unsigned int; enum class coordinate_normalization_mode : unsigned int; namespace detail { -class sampler_impl { +class __SYCL_EXPORT sampler_impl { public: #ifdef __SYCL_DEVICE_ONLY__ __ocl_sampler_t m_Sampler; diff --git a/sycl/include/CL/sycl/detail/stream_impl.hpp b/sycl/include/CL/sycl/detail/stream_impl.hpp index d91897264ec50..1f34361d27fc0 100644 --- a/sycl/include/CL/sycl/detail/stream_impl.hpp +++ b/sycl/include/CL/sycl/detail/stream_impl.hpp @@ -11,6 +11,7 @@ #include #include #include +#include #include #include #include @@ -77,7 +78,7 @@ using EnableIfSwizzleVec = typename std::enable_if::value, typename IsSwizzleOp::Type>::type; -class stream_impl { +class __SYCL_EXPORT stream_impl { public: using GlobalBufAccessorT = accessor +#include #include #include #include @@ -36,7 +37,7 @@ class aligned_allocator; using sycl_memory_object_allocator = aligned_allocator; // The class serves as a base for all SYCL memory objects. -class SYCLMemObjT : public SYCLMemObjI { +class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI { // The check for output iterator is commented out as it blocks set_final_data // with void * argument to be used. diff --git a/sycl/include/CL/sycl/detail/usm_impl.hpp b/sycl/include/CL/sycl/detail/usm_impl.hpp index c5b234ace7424..47ae9447bb4ec 100644 --- a/sycl/include/CL/sycl/detail/usm_impl.hpp +++ b/sycl/include/CL/sycl/detail/usm_impl.hpp @@ -9,6 +9,7 @@ #include #include +#include #include __SYCL_INLINE_NAMESPACE(cl) { @@ -16,13 +17,15 @@ namespace sycl { namespace detail { namespace usm { -void *alignedAlloc(size_t Alignment, size_t Bytes, const context &Ctxt, - const device &Dev, cl::sycl::usm::alloc Kind); +__SYCL_EXPORT void *alignedAlloc(size_t Alignment, size_t Bytes, + const context &Ctxt, const device &Dev, + cl::sycl::usm::alloc Kind); -void *alignedAllocHost(size_t Alignment, size_t Bytes, const context &Ctxt, - cl::sycl::usm::alloc Kind); +__SYCL_EXPORT void *alignedAllocHost(size_t Alignment, size_t Bytes, + const context &Ctxt, + cl::sycl::usm::alloc Kind); -void free(void *Ptr, const context &Ctxt); +__SYCL_EXPORT void free(void *Ptr, const context &Ctxt); } // namespace usm } // namespace detail diff --git a/sycl/include/CL/sycl/device.hpp b/sycl/include/CL/sycl/device.hpp index 7c78adf7cd417..c1ea233eaa68c 100644 --- a/sycl/include/CL/sycl/device.hpp +++ b/sycl/include/CL/sycl/device.hpp @@ -9,9 +9,11 @@ #pragma once #include +#include #include #include #include + #include #include @@ -22,7 +24,7 @@ class device_selector; namespace detail { class device_impl; } -class device { +class __SYCL_EXPORT device { public: /// Constructs a SYCL device instance as a host device. device(); diff --git a/sycl/include/CL/sycl/device_selector.hpp b/sycl/include/CL/sycl/device_selector.hpp index 67bb373298e94..f2ad228e2fd47 100644 --- a/sycl/include/CL/sycl/device_selector.hpp +++ b/sycl/include/CL/sycl/device_selector.hpp @@ -8,6 +8,8 @@ #pragma once +#include + // 4.6.1 Device selection class __SYCL_INLINE_NAMESPACE(cl) { @@ -16,7 +18,7 @@ namespace sycl { // Forward declarations class device; -class device_selector { +class __SYCL_EXPORT device_selector { public: virtual ~device_selector() = default; @@ -25,27 +27,27 @@ class device_selector { virtual int operator()(const device &device) const = 0; }; -class default_selector : public device_selector { +class __SYCL_EXPORT default_selector : public device_selector { public: int operator()(const device &dev) const override; }; -class gpu_selector : public device_selector { +class __SYCL_EXPORT gpu_selector : public device_selector { public: int operator()(const device &dev) const override; }; -class cpu_selector : public device_selector { +class __SYCL_EXPORT cpu_selector : public device_selector { public: int operator()(const device &dev) const override; }; -class accelerator_selector : public device_selector { +class __SYCL_EXPORT accelerator_selector : public device_selector { public: int operator()(const device &dev) const override; }; -class host_selector : public device_selector { +class __SYCL_EXPORT host_selector : public device_selector { public: int operator()(const device &dev) const override; }; diff --git a/sycl/include/CL/sycl/event.hpp b/sycl/include/CL/sycl/event.hpp index 86e24dda794fc..001f5ee209c15 100644 --- a/sycl/include/CL/sycl/event.hpp +++ b/sycl/include/CL/sycl/event.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include @@ -22,7 +23,7 @@ namespace detail { class event_impl; } -class event { +class __SYCL_EXPORT event { public: /// Constructs a ready SYCL event. /// diff --git a/sycl/include/CL/sycl/exception.hpp b/sycl/include/CL/sycl/exception.hpp index 781d108ec0dfa..aedc9eb02bebc 100644 --- a/sycl/include/CL/sycl/exception.hpp +++ b/sycl/include/CL/sycl/exception.hpp @@ -11,6 +11,7 @@ // 4.9.2 Exception Class Interface #include +#include #include #include @@ -24,7 +25,7 @@ class context; // Derive from std::exception so uncaught exceptions are printed in c++ default // exception handler. -class exception: public std::exception { +class __SYCL_EXPORT exception : public std::exception { public: exception() = default; diff --git a/sycl/include/CL/sycl/exception_list.hpp b/sycl/include/CL/sycl/exception_list.hpp index 9470ba14848ff..123ad8a3811bf 100644 --- a/sycl/include/CL/sycl/exception_list.hpp +++ b/sycl/include/CL/sycl/exception_list.hpp @@ -11,6 +11,7 @@ // 4.9.2 Exception Class Interface #include +#include #include #include @@ -23,7 +24,7 @@ namespace detail { class queue_impl; } -class exception_list { +class __SYCL_EXPORT exception_list { public: using value_type = exception_ptr_class; using reference = value_type &; diff --git a/sycl/include/CL/sycl/half_type.hpp b/sycl/include/CL/sycl/half_type.hpp index 7a65197c3478b..d836af37c9698 100644 --- a/sycl/include/CL/sycl/half_type.hpp +++ b/sycl/include/CL/sycl/half_type.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include @@ -31,7 +32,7 @@ namespace sycl { namespace detail { namespace host_half_impl { -class half { +class __SYCL_EXPORT half { public: half() = default; half(const half &) = default; diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index c9b4f0ed3179f..621602b032167 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -11,6 +11,7 @@ #include #include #include +#include #include #include #include @@ -19,7 +20,6 @@ #include #include #include - #include #include @@ -104,7 +104,7 @@ template struct get_kernel_name_t { using name = Type; }; -device getDeviceFromHandler(handler &); +__SYCL_EXPORT device getDeviceFromHandler(handler &); } // namespace detail /// 4.8.3 Command group handler class @@ -134,7 +134,7 @@ device getDeviceFromHandler(handler &); /// end. So, handler class contains all fields simultaneously, then during /// "finalization" it constructs CG object, that represents specific operation, /// passing fields that are required only. -class handler { +class __SYCL_EXPORT handler { private: /// Constructs SYCL handler from queue. /// diff --git a/sycl/include/CL/sycl/intel/fpga_reg.hpp b/sycl/include/CL/sycl/intel/fpga_reg.hpp index da43baa146b01..0078dd66c383c 100644 --- a/sycl/include/CL/sycl/intel/fpga_reg.hpp +++ b/sycl/include/CL/sycl/intel/fpga_reg.hpp @@ -30,6 +30,6 @@ template T fpga_reg(const T &t) { // Currently clang does not support nested namespace for attributes namespace intelfpga { template T fpga_reg(const T &t) { - return cl::sycl::intel::fpga_reg(t); -} + return cl::sycl::intel::fpga_reg(t); +} } diff --git a/sycl/include/CL/sycl/intel/function_pointer.hpp b/sycl/include/CL/sycl/intel/function_pointer.hpp index 69d32ac0e4940..2aa64cdd2e580 100644 --- a/sycl/include/CL/sycl/intel/function_pointer.hpp +++ b/sycl/include/CL/sycl/intel/function_pointer.hpp @@ -8,6 +8,7 @@ #pragma once +#include #include #include #include @@ -17,8 +18,8 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { namespace detail { -cl_ulong getDeviceFunctionPointerImpl(device &D, program &P, - const char *FuncName); +__SYCL_EXPORT cl_ulong getDeviceFunctionPointerImpl(device &D, program &P, + const char *FuncName); } namespace intel { diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp index 7d610b7983f50..8e4e100b93ba5 100644 --- a/sycl/include/CL/sycl/intel/sub_group.hpp +++ b/sycl/include/CL/sycl/intel/sub_group.hpp @@ -159,12 +159,12 @@ struct sub_group { /* --- vote / ballot functions --- */ - __SYCL_DEPRECATED__("Use sycl::intel::any_of instead.") + __SYCL_EXPORT_DEPRECATED("Use sycl::intel::any_of instead.") bool any(bool predicate) const { return __spirv_GroupAny(__spv::Scope::Subgroup, predicate); } - __SYCL_DEPRECATED__("Use sycl::intel::all_of instead.") + __SYCL_EXPORT_DEPRECATED("Use sycl::intel::all_of instead.") bool all(bool predicate) const { return __spirv_GroupAll(__spv::Scope::Subgroup, predicate); } @@ -176,13 +176,13 @@ struct sub_group { /* --- collectives --- */ template - __SYCL_DEPRECATED__("Use sycl::intel::broadcast instead.") + __SYCL_EXPORT_DEPRECATED("Use sycl::intel::broadcast instead.") EnableIfIsScalarArithmetic broadcast(T x, id<1> local_id) const { return detail::spirv::GroupBroadcast<__spv::Scope::Subgroup>(x, local_id); } template - __SYCL_DEPRECATED__("Use sycl::intel::reduce instead.") + __SYCL_EXPORT_DEPRECATED("Use sycl::intel::reduce instead.") EnableIfIsScalarArithmetic reduce(T x, BinaryOperation op) const { return detail::calc( @@ -190,13 +190,13 @@ struct sub_group { } template - __SYCL_DEPRECATED__("Use sycl::intel::reduce instead.") + __SYCL_EXPORT_DEPRECATED("Use sycl::intel::reduce instead.") EnableIfIsScalarArithmetic reduce(T x, T init, BinaryOperation op) const { return op(init, reduce(x, op)); } template - __SYCL_DEPRECATED__("Use sycl::intel::exclusive_scan instead.") + __SYCL_EXPORT_DEPRECATED("Use sycl::intel::exclusive_scan instead.") EnableIfIsScalarArithmetic exclusive_scan(T x, BinaryOperation op) const { return detail::calc( @@ -204,7 +204,7 @@ struct sub_group { } template - __SYCL_DEPRECATED__("Use sycl::intel::exclusive_scan instead.") + __SYCL_EXPORT_DEPRECATED("Use sycl::intel::exclusive_scan instead.") EnableIfIsScalarArithmetic exclusive_scan(T x, T init, BinaryOperation op) const { if (get_local_id().get(0) == 0) { @@ -218,7 +218,7 @@ struct sub_group { } template - __SYCL_DEPRECATED__("Use sycl::intel::inclusive_scan instead.") + __SYCL_EXPORT_DEPRECATED("Use sycl::intel::inclusive_scan instead.") EnableIfIsScalarArithmetic inclusive_scan(T x, BinaryOperation op) const { return detail::calc( @@ -226,7 +226,7 @@ struct sub_group { } template - __SYCL_DEPRECATED__("Use sycl::intel::inclusive_scan instead.") + __SYCL_EXPORT_DEPRECATED("Use sycl::intel::inclusive_scan instead.") EnableIfIsScalarArithmetic inclusive_scan(T x, BinaryOperation op, T init) const { if (get_local_id().get(0) == 0) { diff --git a/sycl/include/CL/sycl/kernel.hpp b/sycl/include/CL/sycl/kernel.hpp index 1aa8afec7703e..903bb59cf6a67 100644 --- a/sycl/include/CL/sycl/kernel.hpp +++ b/sycl/include/CL/sycl/kernel.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include @@ -23,7 +24,7 @@ namespace detail { class kernel_impl; } -class kernel { +class __SYCL_EXPORT kernel { public: /// Constructs a SYCL kernel instance from an OpenCL cl_kernel /// diff --git a/sycl/include/CL/sycl/ordered_queue.hpp b/sycl/include/CL/sycl/ordered_queue.hpp index eda3b48e18f66..fdbc97be7c491 100644 --- a/sycl/include/CL/sycl/ordered_queue.hpp +++ b/sycl/include/CL/sycl/ordered_queue.hpp @@ -10,6 +10,7 @@ #include #include +#include #include #include #include @@ -28,8 +29,8 @@ namespace detail { class queue_impl; } -class __SYCL_DEPRECATED__("Replaced by in_order queue property") ordered_queue { - +class __SYCL_EXPORT_DEPRECATED("Replaced by in_order queue property") + ordered_queue { public: explicit ordered_queue(const property_list &propList = {}) : ordered_queue(default_selector(), async_handler{}, propList) {} diff --git a/sycl/include/CL/sycl/platform.hpp b/sycl/include/CL/sycl/platform.hpp index 5841a0bbfd1d6..4f34cf43600e4 100644 --- a/sycl/include/CL/sycl/platform.hpp +++ b/sycl/include/CL/sycl/platform.hpp @@ -8,6 +8,7 @@ #pragma once #include +#include #include // 4.6.2 Platform class @@ -23,7 +24,7 @@ namespace detail { class platform_impl; } -class platform { +class __SYCL_EXPORT platform { public: /// Constructs a SYCL platform as a host platform. platform(); diff --git a/sycl/include/CL/sycl/program.hpp b/sycl/include/CL/sycl/program.hpp index e3aee274575d2..9e3d2e5714086 100644 --- a/sycl/include/CL/sycl/program.hpp +++ b/sycl/include/CL/sycl/program.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include #include @@ -28,7 +29,7 @@ class program_impl; enum class program_state { none, compiled, linked }; -class program { +class __SYCL_EXPORT program { public: program() = delete; diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp index 870854c682686..bea3df4857094 100644 --- a/sycl/include/CL/sycl/queue.hpp +++ b/sycl/include/CL/sycl/queue.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include #include @@ -30,7 +31,7 @@ namespace detail { class queue_impl; } -class queue { +class __SYCL_EXPORT queue { public: /// Constructs a SYCL queue instance using the device returned by an instance /// of default_selector. diff --git a/sycl/include/CL/sycl/sampler.hpp b/sycl/include/CL/sycl/sampler.hpp index a7741cb9bb50a..2f099aae25ce3 100644 --- a/sycl/include/CL/sycl/sampler.hpp +++ b/sycl/include/CL/sycl/sampler.hpp @@ -11,6 +11,7 @@ #include #include #include +#include #include __SYCL_INLINE_NAMESPACE(cl) { @@ -39,7 +40,7 @@ template #include __SYCL_INLINE_NAMESPACE(cl) { @@ -89,7 +90,7 @@ inline __width_manipulator__ setw(int Width) { return __width_manipulator__(Width); } -class stream { +class __SYCL_EXPORT stream { public: stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH); diff --git a/sycl/include/CL/sycl/usm.hpp b/sycl/include/CL/sycl/usm.hpp index 0a80524704b26..ba9e417b98499 100644 --- a/sycl/include/CL/sycl/usm.hpp +++ b/sycl/include/CL/sycl/usm.hpp @@ -7,6 +7,7 @@ // ===--------------------------------------------------------------------=== // #pragma once +#include #include #include @@ -17,43 +18,52 @@ namespace sycl { /// // Explicit USM /// -void *malloc_device(size_t size, const device &dev, const context &ctxt); -void *malloc_device(size_t size, const queue &q); +__SYCL_EXPORT void *malloc_device(size_t size, const device &dev, + const context &ctxt); +__SYCL_EXPORT void *malloc_device(size_t size, const queue &q); -void *aligned_alloc_device(size_t alignment, size_t size, const device &dev, - const context &ctxt); -void *aligned_alloc_device(size_t alignment, size_t size, const queue &q); +__SYCL_EXPORT void *aligned_alloc_device(size_t alignment, size_t size, + const device &dev, + const context &ctxt); +__SYCL_EXPORT void *aligned_alloc_device(size_t alignment, size_t size, + const queue &q); -void free(void *ptr, const context &ctxt); -void free(void *ptr, const queue &q); +__SYCL_EXPORT void free(void *ptr, const context &ctxt); +__SYCL_EXPORT void free(void *ptr, const queue &q); /// // Restricted USM /// -void *malloc_host(size_t size, const context &ctxt); -void *malloc_host(size_t size, const queue &q); +__SYCL_EXPORT void *malloc_host(size_t size, const context &ctxt); +__SYCL_EXPORT void *malloc_host(size_t size, const queue &q); -void *malloc_shared(size_t size, const device &dev, const context &ctxt); -void *malloc_shared(size_t size, const queue &q); +__SYCL_EXPORT void *malloc_shared(size_t size, const device &dev, + const context &ctxt); +__SYCL_EXPORT void *malloc_shared(size_t size, const queue &q); -void *aligned_alloc_host(size_t alignment, size_t size, const context &ctxt); -void *aligned_alloc_host(size_t alignment, size_t size, const queue &q); +__SYCL_EXPORT void *aligned_alloc_host(size_t alignment, size_t size, + const context &ctxt); +__SYCL_EXPORT void *aligned_alloc_host(size_t alignment, size_t size, + const queue &q); -void *aligned_alloc_shared(size_t alignment, size_t size, const device &dev, - const context &ctxt); -void *aligned_alloc_shared(size_t alignment, size_t size, const queue &q); +__SYCL_EXPORT void *aligned_alloc_shared(size_t alignment, size_t size, + const device &dev, + const context &ctxt); +__SYCL_EXPORT void *aligned_alloc_shared(size_t alignment, size_t size, + const queue &q); /// // single form /// -void *malloc(size_t size, const device &dev, const context &ctxt, - usm::alloc kind); -void *malloc(size_t size, const queue &q, usm::alloc kind); +__SYCL_EXPORT void *malloc(size_t size, const device &dev, const context &ctxt, + usm::alloc kind); +__SYCL_EXPORT void *malloc(size_t size, const queue &q, usm::alloc kind); -void *aligned_alloc(size_t alignment, size_t size, const device &dev, - const context &ctxt, usm::alloc kind); -void *aligned_alloc(size_t alignment, size_t size, const queue &q, - usm::alloc kind); +__SYCL_EXPORT void *aligned_alloc(size_t alignment, size_t size, + const device &dev, const context &ctxt, + usm::alloc kind); +__SYCL_EXPORT void *aligned_alloc(size_t alignment, size_t size, const queue &q, + usm::alloc kind); /// // Template forms @@ -150,14 +160,14 @@ T *aligned_alloc(size_t Alignment, size_t Count, const queue &Q, /// /// \param ptr is the USM pointer to query /// \param ctxt is the sycl context the ptr was allocated in -usm::alloc get_pointer_type(const void *ptr, const context &ctxt); +__SYCL_EXPORT usm::alloc get_pointer_type(const void *ptr, const context &ctxt); /// Queries the device against which the pointer was allocated /// Throws an invalid_object_error if ptr is a host allocation. /// /// \param ptr is the USM pointer to query /// \param ctxt is the sycl context the ptr was allocated in -device get_pointer_device(const void *ptr, const context &ctxt); +__SYCL_EXPORT device get_pointer_device(const void *ptr, const context &ctxt); } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/CL/sycl/usm/usm_allocator.hpp b/sycl/include/CL/sycl/usm/usm_allocator.hpp index 0a6e49845cde3..2967f1fe0195d 100644 --- a/sycl/include/CL/sycl/usm/usm_allocator.hpp +++ b/sycl/include/CL/sycl/usm/usm_allocator.hpp @@ -8,6 +8,7 @@ #pragma once #include +#include #include #include #include @@ -20,9 +21,10 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { // Forward declarations. -void *aligned_alloc(size_t alignment, size_t size, const device &dev, - const context &ctxt, usm::alloc kind); -void free(void *ptr, const context &ctxt); +__SYCL_EXPORT void *aligned_alloc(size_t alignment, size_t size, + const device &dev, const context &ctxt, + usm::alloc kind); +__SYCL_EXPORT void free(void *ptr, const context &ctxt); template class usm_allocator { diff --git a/sycl/plugins/CMakeLists.txt b/sycl/plugins/CMakeLists.txt index 791b4240dc005..7dbd3d76a0842 100644 --- a/sycl/plugins/CMakeLists.txt +++ b/sycl/plugins/CMakeLists.txt @@ -1,3 +1,5 @@ +set(CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS ON) + if(SYCL_BUILD_PI_CUDA) add_subdirectory(cuda) endif() diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index a8308ffd89add..a4801f58acff5 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -14,12 +14,18 @@ endif() function(add_sycl_rt_library LIB_NAME) # Add an optional argument so we can get the library name to # link with for Windows Debug version - cmake_parse_arguments(ARG "" "XPTI_LIB" "" ${ARGN}) + cmake_parse_arguments(ARG "" "XPTI_LIB" "COMPILE_OPTIONS;SOURCES" ${ARGN}) - add_library(${LIB_NAME} SHARED ${ARG_UNPARSED_ARGUMENTS}) + set(LIB_OBJ_NAME ${LIB_NAME}_object) + + add_library(${LIB_OBJ_NAME} OBJECT ${ARG_SOURCES}) + add_library(${LIB_NAME} SHARED $) + if (ARG_COMPILE_OPTIONS) + target_compile_options(${LIB_OBJ_NAME} PRIVATE ${ARG_COMPILE_OPTIONS}) + endif() #To-Do: Remove dependency on icd loader and opencl headers. - add_dependencies(${LIB_NAME} + add_dependencies(${LIB_OBJ_NAME} ocl-icd ocl-headers sycl-headers @@ -28,14 +34,16 @@ function(add_sycl_rt_library LIB_NAME) set_target_properties(${LIB_NAME} PROPERTIES LINKER_LANGUAGE CXX) if (SYCL_ENABLE_XPTI_TRACING) - target_compile_definitions(${LIB_NAME} PRIVATE XPTI_ENABLE_INSTRUMENTATION XPTI_STATIC_LIBRARY) + target_compile_definitions(${LIB_OBJ_NAME} PRIVATE XPTI_ENABLE_INSTRUMENTATION XPTI_STATIC_LIBRARY) target_link_libraries(${LIB_NAME} PRIVATE ${ARG_XPTI_LIB}) endif() if (MSVC) - target_compile_definitions(${LIB_NAME} PRIVATE __SYCL_BUILD_SYCL_DLL ) + target_compile_definitions(${LIB_OBJ_NAME} PRIVATE __SYCL_BUILD_SYCL_DLL ) target_link_libraries(${LIB_NAME} PRIVATE shlwapi) else() + target_compile_options(${LIB_OBJ_NAME} PUBLIC + -fvisibility=hidden -fvisibility-inlines-hidden) set(linker_script "${CMAKE_CURRENT_SOURCE_DIR}/ld-version-script.txt") target_link_libraries( ${LIB_NAME} PRIVATE "-Wl,--version-script=${linker_script}") @@ -46,17 +54,21 @@ function(add_sycl_rt_library LIB_NAME) endif() target_include_directories( - ${LIB_NAME} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} "${sycl_inc_dir}") + ${LIB_OBJ_NAME} + PRIVATE + ${CMAKE_CURRENT_SOURCE_DIR} + "${sycl_inc_dir}" + ${OpenCL_INCLUDE_DIRS} + ) target_link_libraries(${LIB_NAME} PRIVATE - OpenCL::Headers ${OpenCL_LIBRARIES} ${CMAKE_DL_LIBS} PUBLIC $<$:pi_cuda> ) - target_compile_definitions(${LIB_NAME} + target_compile_definitions(${LIB_OBJ_NAME} PUBLIC $<$:USE_PI_CUDA>) @@ -129,55 +141,33 @@ set(SYCL_SOURCES "$<$,$>:detail/posix_pi.cpp>" ) -if (SYCL_ENABLE_XPTI_TRACING) - add_sycl_rt_library(sycl ${SYCL_SOURCES} XPTI_LIB xpti) -else() - add_sycl_rt_library(sycl ${SYCL_SOURCES}) -endif() - if (MSVC) # MSVC provides two incompatible build variants for its CRT: release and debug # To avoid potential issues in user code we also need to provide two kinds # of SYCL Runtime Library for release and debug configurations. - set(SYCL_CXX_FLAGS "") - if (CMAKE_BUILD_TYPE MATCHES "Debug") - set(SYCL_CXX_FLAGS "${CMAKE_CXX_FLAGS_DEBUG}") - string(REPLACE "/MDd" "" SYCL_CXX_FLAGS "${SYCL_CXX_FLAGS}") - string(REPLACE "/MTd" "" SYCL_CXX_FLAGS "${SYCL_CXX_FLAGS}") - else() - if (CMAKE_BUILD_TYPE MATCHES "Release") - set(SYCL_CXX_FLAGS "${CMAKE_CXX_FLAGS_RELEASE}") - elseif (CMAKE_BUILD_TYPE MATCHES "RelWithDebInfo") - set(SYCL_CXX_FLAGS "${CMAKE_CXX_FLAGS_MINSIZEREL}") - elseif (CMAKE_BUILD_TYPE MATCHES "MinSizeRel") - set(SYCL_CXX_FLAGS "${CMAKE_CXX_FLAGS_RELWITHDEBINFO}") - endif() - string(REPLACE "/MD" "" SYCL_CXX_FLAGS "${SYCL_CXX_FLAGS}") - string(REPLACE "/MT" "" SYCL_CXX_FLAGS "${SYCL_CXX_FLAGS}") - endif() - - # target_compile_options requires list of options, not a string - string(REPLACE " " ";" SYCL_CXX_FLAGS "${SYCL_CXX_FLAGS}") - - set(SYCL_CXX_FLAGS_RELEASE "${SYCL_CXX_FLAGS};/MD") - set(SYCL_CXX_FLAGS_DEBUG "${SYCL_CXX_FLAGS};/MDd") - - # CMake automatically applies these flags to all targets. To override this - # behavior, options lists are reset. - set(CMAKE_CXX_FLAGS_RELEASE "") - set(CMAKE_CXX_FLAGS_MINSIZEREL "") - set(CMAKE_CXX_FLAGS_RELWITHDEBINFO "") - set(CMAKE_CXX_FLAGS_DEBUG "") - - target_compile_options(sycl PRIVATE ${SYCL_CXX_FLAGS_RELEASE}) + foreach(flag_var + CMAKE_CXX_FLAGS CMAKE_CXX_FLAGS_DEBUG CMAKE_CXX_FLAGS_RELEASE + CMAKE_CXX_FLAGS_MINSIZEREL CMAKE_CXX_FLAGS_RELWITHDEBINFO) + string(REGEX REPLACE "/MD" "" ${flag_var} "${${flag_var}}") + string(REGEX REPLACE "/MT" "" ${flag_var} "${${flag_var}}") + string(REGEX REPLACE "/MDd" "" ${flag_var} "${${flag_var}}") + string(REGEX REPLACE "/MTd" "" ${flag_var} "${${flag_var}}") + endforeach() if (SYCL_ENABLE_XPTI_TRACING) - add_sycl_rt_library(sycld ${SYCL_SOURCES} XPTI_LIB xptid) + add_sycl_rt_library(sycld XPTI_LIB xptid COMPILE_OPTIONS "/MDd" SOURCES ${SYCL_SOURCES}) else() - add_sycl_rt_library(sycld ${SYCL_SOURCES}) + add_sycl_rt_library(sycld COMPILE_OPTIONS "/MDd" SOURCES ${SYCL_SOURCES}) endif() - target_compile_options(sycld PRIVATE ${SYCL_CXX_FLAGS_DEBUG}) + set(SYCL_EXTRA_OPTS "/MD") + +endif() + +if (SYCL_ENABLE_XPTI_TRACING) + add_sycl_rt_library(sycl XPTI_LIB xpti COMPILE_OPTIONS ${SYCL_EXTRA_OPTS} SOURCES ${SYCL_SOURCES}) +else() + add_sycl_rt_library(sycl COMPILE_OPTIONS ${SYCL_EXTRA_OPTS} SOURCES ${SYCL_SOURCES}) endif() # Enable new IN_LIST operator. diff --git a/sycl/source/context.cpp b/sycl/source/context.cpp index 8a4cf83a92520..49b49914eabf6 100644 --- a/sycl/source/context.cpp +++ b/sycl/source/context.cpp @@ -73,7 +73,8 @@ context::context(cl_context ClContext, async_handler AsyncHandler) { } #define PARAM_TRAITS_SPEC(param_type, param, ret_type) \ - template <> ret_type context::get_info() const { \ + template <> \ + __SYCL_EXPORT ret_type context::get_info() const { \ return impl->get_info(); \ } diff --git a/sycl/source/detail/builtins_common.cpp b/sycl/source/detail/builtins_common.cpp index d9e4ad91cd4f0..130eb1ad1d9ed 100644 --- a/sycl/source/detail/builtins_common.cpp +++ b/sycl/source/detail/builtins_common.cpp @@ -26,32 +26,35 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace __host_std { namespace { -template inline T __fclamp(T x, T minval, T maxval) { +template __SYCL_EXPORT inline T __fclamp(T x, T minval, T maxval) { return std::fmin(std::fmax(x, minval), maxval); } -template inline T __degrees(T radians) { +template __SYCL_EXPORT inline T __degrees(T radians) { return (180 / M_PI) * radians; } -template inline T __mix(T x, T y, T a) { return x + (y - x) * a; } +template __SYCL_EXPORT inline T __mix(T x, T y, T a) { + return x + (y - x) * a; +} -template inline T __radians(T degrees) { +template __SYCL_EXPORT inline T __radians(T degrees) { return (M_PI / 180) * degrees; } -template inline T __step(T edge, T x) { +template __SYCL_EXPORT inline T __step(T edge, T x) { return (x < edge) ? 0.0 : 1.0; } -template inline T __smoothstep(T edge0, T edge1, T x) { +template +__SYCL_EXPORT inline T __smoothstep(T edge0, T edge1, T x) { T t; T v = (x - edge0) / (edge1 - edge0); t = __fclamp(v, T(0), T(1)); return t * t * (3 - 2 * t); } -template inline T __sign(T x) { +template __SYCL_EXPORT inline T __sign(T x) { if (std::isnan(d::cast_if_host_half(x))) return T(0.0); if (x > 0) @@ -66,15 +69,16 @@ template inline T __sign(T x) { // --------------- 4.13.5 Common functions. Host implementations --------------- // fclamp -s::cl_float fclamp(s::cl_float x, s::cl_float minval, - s::cl_float maxval) __NOEXC { +__SYCL_EXPORT s::cl_float fclamp(s::cl_float x, s::cl_float minval, + s::cl_float maxval) __NOEXC { return __fclamp(x, minval, maxval); } -s::cl_double fclamp(s::cl_double x, s::cl_double minval, - s::cl_double maxval) __NOEXC { +__SYCL_EXPORT s::cl_double fclamp(s::cl_double x, s::cl_double minval, + s::cl_double maxval) __NOEXC { return __fclamp(x, minval, maxval); } -s::cl_half fclamp(s::cl_half x, s::cl_half minval, s::cl_half maxval) __NOEXC { +__SYCL_EXPORT s::cl_half fclamp(s::cl_half x, s::cl_half minval, + s::cl_half maxval) __NOEXC { return __fclamp(x, minval, maxval); } MAKE_1V_2V_3V(fclamp, s::cl_float, s::cl_float, s::cl_float, s::cl_float) @@ -82,23 +86,27 @@ MAKE_1V_2V_3V(fclamp, s::cl_double, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V_3V(fclamp, s::cl_half, s::cl_half, s::cl_half, s::cl_half) // degrees -s::cl_float degrees(s::cl_float radians) __NOEXC { return __degrees(radians); } -s::cl_double degrees(s::cl_double radians) __NOEXC { +__SYCL_EXPORT s::cl_float degrees(s::cl_float radians) __NOEXC { + return __degrees(radians); +} +__SYCL_EXPORT s::cl_double degrees(s::cl_double radians) __NOEXC { + return __degrees(radians); +} +__SYCL_EXPORT s::cl_half degrees(s::cl_half radians) __NOEXC { return __degrees(radians); } -s::cl_half degrees(s::cl_half radians) __NOEXC { return __degrees(radians); } MAKE_1V(degrees, s::cl_float, s::cl_float) MAKE_1V(degrees, s::cl_double, s::cl_double) MAKE_1V(degrees, s::cl_half, s::cl_half) // fmin_common -s::cl_float fmin_common(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float fmin_common(s::cl_float x, s::cl_float y) __NOEXC { return std::fmin(x, y); } -s::cl_double fmin_common(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_double fmin_common(s::cl_double x, s::cl_double y) __NOEXC { return std::fmin(x, y); } -s::cl_half fmin_common(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_half fmin_common(s::cl_half x, s::cl_half y) __NOEXC { return std::fmin(x, y); } MAKE_1V_2V(fmin_common, s::cl_float, s::cl_float, s::cl_float) @@ -106,13 +114,13 @@ MAKE_1V_2V(fmin_common, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(fmin_common, s::cl_half, s::cl_half, s::cl_half) // fmax_common -s::cl_float fmax_common(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float fmax_common(s::cl_float x, s::cl_float y) __NOEXC { return std::fmax(x, y); } -s::cl_double fmax_common(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_double fmax_common(s::cl_double x, s::cl_double y) __NOEXC { return std::fmax(x, y); } -s::cl_half fmax_common(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_half fmax_common(s::cl_half x, s::cl_half y) __NOEXC { return std::fmax(x, y); } MAKE_1V_2V(fmax_common, s::cl_float, s::cl_float, s::cl_float) @@ -120,13 +128,15 @@ MAKE_1V_2V(fmax_common, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(fmax_common, s::cl_half, s::cl_half, s::cl_half) // mix -s::cl_float mix(s::cl_float x, s::cl_float y, s::cl_float a) __NOEXC { +__SYCL_EXPORT s::cl_float mix(s::cl_float x, s::cl_float y, + s::cl_float a) __NOEXC { return __mix(x, y, a); } -s::cl_double mix(s::cl_double x, s::cl_double y, s::cl_double a) __NOEXC { +__SYCL_EXPORT s::cl_double mix(s::cl_double x, s::cl_double y, + s::cl_double a) __NOEXC { return __mix(x, y, a); } -s::cl_half mix(s::cl_half x, s::cl_half y, s::cl_half a) __NOEXC { +__SYCL_EXPORT s::cl_half mix(s::cl_half x, s::cl_half y, s::cl_half a) __NOEXC { return __mix(x, y, a); } MAKE_1V_2V_3V(mix, s::cl_float, s::cl_float, s::cl_float, s::cl_float) @@ -134,23 +144,27 @@ MAKE_1V_2V_3V(mix, s::cl_double, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V_3V(mix, s::cl_half, s::cl_half, s::cl_half, s::cl_half) // radians -s::cl_float radians(s::cl_float degrees) __NOEXC { return __radians(degrees); } -s::cl_double radians(s::cl_double degrees) __NOEXC { +__SYCL_EXPORT s::cl_float radians(s::cl_float degrees) __NOEXC { + return __radians(degrees); +} +__SYCL_EXPORT s::cl_double radians(s::cl_double degrees) __NOEXC { + return __radians(degrees); +} +__SYCL_EXPORT s::cl_half radians(s::cl_half degrees) __NOEXC { return __radians(degrees); } -s::cl_half radians(s::cl_half degrees) __NOEXC { return __radians(degrees); } MAKE_1V(radians, s::cl_float, s::cl_float) MAKE_1V(radians, s::cl_double, s::cl_double) MAKE_1V(radians, s::cl_half, s::cl_half) // step -s::cl_float step(s::cl_float edge, s::cl_float x) __NOEXC { +__SYCL_EXPORT s::cl_float step(s::cl_float edge, s::cl_float x) __NOEXC { return __step(edge, x); } -s::cl_double step(s::cl_double edge, s::cl_double x) __NOEXC { +__SYCL_EXPORT s::cl_double step(s::cl_double edge, s::cl_double x) __NOEXC { return __step(edge, x); } -s::cl_half step(s::cl_half edge, s::cl_half x) __NOEXC { +__SYCL_EXPORT s::cl_half step(s::cl_half edge, s::cl_half x) __NOEXC { return __step(edge, x); } MAKE_1V_2V(step, s::cl_float, s::cl_float, s::cl_float) @@ -158,16 +172,16 @@ MAKE_1V_2V(step, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(step, s::cl_half, s::cl_half, s::cl_half) // smoothstep -s::cl_float smoothstep(s::cl_float edge0, s::cl_float edge1, - s::cl_float x) __NOEXC { +__SYCL_EXPORT s::cl_float smoothstep(s::cl_float edge0, s::cl_float edge1, + s::cl_float x) __NOEXC { return __smoothstep(edge0, edge1, x); } -s::cl_double smoothstep(s::cl_double edge0, s::cl_double edge1, - s::cl_double x) __NOEXC { +__SYCL_EXPORT s::cl_double smoothstep(s::cl_double edge0, s::cl_double edge1, + s::cl_double x) __NOEXC { return __smoothstep(edge0, edge1, x); } -s::cl_half smoothstep(s::cl_half edge0, s::cl_half edge1, - s::cl_half x) __NOEXC { +__SYCL_EXPORT s::cl_half smoothstep(s::cl_half edge0, s::cl_half edge1, + s::cl_half x) __NOEXC { return __smoothstep(edge0, edge1, x); } MAKE_1V_2V_3V(smoothstep, s::cl_float, s::cl_float, s::cl_float, s::cl_float) @@ -176,9 +190,9 @@ MAKE_1V_2V_3V(smoothstep, s::cl_double, s::cl_double, s::cl_double, MAKE_1V_2V_3V(smoothstep, s::cl_half, s::cl_half, s::cl_half, s::cl_half) // sign -s::cl_float sign(s::cl_float x) __NOEXC { return __sign(x); } -s::cl_double sign(s::cl_double x) __NOEXC { return __sign(x); } -s::cl_half sign(s::cl_half x) __NOEXC { return __sign(x); } +__SYCL_EXPORT s::cl_float sign(s::cl_float x) __NOEXC { return __sign(x); } +__SYCL_EXPORT s::cl_double sign(s::cl_double x) __NOEXC { return __sign(x); } +__SYCL_EXPORT s::cl_half sign(s::cl_half x) __NOEXC { return __sign(x); } MAKE_1V(sign, s::cl_float, s::cl_float) MAKE_1V(sign, s::cl_double, s::cl_double) MAKE_1V(sign, s::cl_half, s::cl_half) diff --git a/sycl/source/detail/builtins_geometric.cpp b/sycl/source/detail/builtins_geometric.cpp index 2d48b1062f7ac..605e9fc78f15f 100644 --- a/sycl/source/detail/builtins_geometric.cpp +++ b/sycl/source/detail/builtins_geometric.cpp @@ -19,19 +19,19 @@ namespace d = s::detail; __SYCL_INLINE_NAMESPACE(cl) { namespace __host_std { -s::cl_float Dot(s::cl_float2, s::cl_float2); -s::cl_float Dot(s::cl_float3, s::cl_float3); -s::cl_float Dot(s::cl_float4, s::cl_float4); -s::cl_double Dot(s::cl_double2, s::cl_double2); -s::cl_double Dot(s::cl_double3, s::cl_double3); -s::cl_double Dot(s::cl_double4, s::cl_double4); -s::cl_half Dot(s::cl_half2, s::cl_half2); -s::cl_half Dot(s::cl_half3, s::cl_half3); -s::cl_half Dot(s::cl_half4, s::cl_half4); - -s::cl_int All(s::cl_int2); -s::cl_int All(s::cl_int3); -s::cl_int All(s::cl_int4); +__SYCL_EXPORT s::cl_float Dot(s::cl_float2, s::cl_float2); +__SYCL_EXPORT s::cl_float Dot(s::cl_float3, s::cl_float3); +__SYCL_EXPORT s::cl_float Dot(s::cl_float4, s::cl_float4); +__SYCL_EXPORT s::cl_double Dot(s::cl_double2, s::cl_double2); +__SYCL_EXPORT s::cl_double Dot(s::cl_double3, s::cl_double3); +__SYCL_EXPORT s::cl_double Dot(s::cl_double4, s::cl_double4); +__SYCL_EXPORT s::cl_half Dot(s::cl_half2, s::cl_half2); +__SYCL_EXPORT s::cl_half Dot(s::cl_half3, s::cl_half3); +__SYCL_EXPORT s::cl_half Dot(s::cl_half4, s::cl_half4); + +__SYCL_EXPORT s::cl_int All(s::cl_int2); +__SYCL_EXPORT s::cl_int All(s::cl_int3); +__SYCL_EXPORT s::cl_int All(s::cl_int4); namespace { @@ -105,29 +105,35 @@ __fast_normalize(T t) { // --------------- 4.13.6 Geometric functions. Host implementations ------------ // cross -s::cl_float3 cross(s::cl_float3 p0, s::cl_float3 p1) __NOEXC { +__SYCL_EXPORT s::cl_float3 cross(s::cl_float3 p0, s::cl_float3 p1) __NOEXC { return __cross(p0, p1); } -s::cl_float4 cross(s::cl_float4 p0, s::cl_float4 p1) __NOEXC { +__SYCL_EXPORT s::cl_float4 cross(s::cl_float4 p0, s::cl_float4 p1) __NOEXC { return __cross(p0, p1); } -s::cl_double3 cross(s::cl_double3 p0, s::cl_double3 p1) __NOEXC { +__SYCL_EXPORT s::cl_double3 cross(s::cl_double3 p0, s::cl_double3 p1) __NOEXC { return __cross(p0, p1); } -s::cl_double4 cross(s::cl_double4 p0, s::cl_double4 p1) __NOEXC { +__SYCL_EXPORT s::cl_double4 cross(s::cl_double4 p0, s::cl_double4 p1) __NOEXC { return __cross(p0, p1); } -s::cl_half3 cross(s::cl_half3 p0, s::cl_half3 p1) __NOEXC { +__SYCL_EXPORT s::cl_half3 cross(s::cl_half3 p0, s::cl_half3 p1) __NOEXC { return __cross(p0, p1); } -s::cl_half4 cross(s::cl_half4 p0, s::cl_half4 p1) __NOEXC { +__SYCL_EXPORT s::cl_half4 cross(s::cl_half4 p0, s::cl_half4 p1) __NOEXC { return __cross(p0, p1); } // FMul -s::cl_float FMul(s::cl_float p0, s::cl_float p1) { return __FMul(p0, p1); } -s::cl_double FMul(s::cl_double p0, s::cl_double p1) { return __FMul(p0, p1); } -s::cl_float FMul(s::cl_half p0, s::cl_half p1) { return __FMul(p0, p1); } +__SYCL_EXPORT s::cl_float FMul(s::cl_float p0, s::cl_float p1) { + return __FMul(p0, p1); +} +__SYCL_EXPORT s::cl_double FMul(s::cl_double p0, s::cl_double p1) { + return __FMul(p0, p1); +} +__SYCL_EXPORT s::cl_float FMul(s::cl_half p0, s::cl_half p1) { + return __FMul(p0, p1); +} // Dot MAKE_GEO_1V_2V_RS(Dot, __FMul_impl, s::cl_float, s::cl_float, s::cl_float) @@ -135,89 +141,119 @@ MAKE_GEO_1V_2V_RS(Dot, __FMul_impl, s::cl_double, s::cl_double, s::cl_double) MAKE_GEO_1V_2V_RS(Dot, __FMul_impl, s::cl_half, s::cl_half, s::cl_half) // length -s::cl_float length(s::cl_float p) { return __length(p); } -s::cl_double length(s::cl_double p) { return __length(p); } -s::cl_half length(s::cl_half p) { return __length(p); } -s::cl_float length(s::cl_float2 p) { return __length(p); } -s::cl_float length(s::cl_float3 p) { return __length(p); } -s::cl_float length(s::cl_float4 p) { return __length(p); } -s::cl_double length(s::cl_double2 p) { return __length(p); } -s::cl_double length(s::cl_double3 p) { return __length(p); } -s::cl_double length(s::cl_double4 p) { return __length(p); } -s::cl_half length(s::cl_half2 p) { return __length(p); } -s::cl_half length(s::cl_half3 p) { return __length(p); } -s::cl_half length(s::cl_half4 p) { return __length(p); } +__SYCL_EXPORT s::cl_float length(s::cl_float p) { return __length(p); } +__SYCL_EXPORT s::cl_double length(s::cl_double p) { return __length(p); } +__SYCL_EXPORT s::cl_half length(s::cl_half p) { return __length(p); } +__SYCL_EXPORT s::cl_float length(s::cl_float2 p) { return __length(p); } +__SYCL_EXPORT s::cl_float length(s::cl_float3 p) { return __length(p); } +__SYCL_EXPORT s::cl_float length(s::cl_float4 p) { return __length(p); } +__SYCL_EXPORT s::cl_double length(s::cl_double2 p) { return __length(p); } +__SYCL_EXPORT s::cl_double length(s::cl_double3 p) { return __length(p); } +__SYCL_EXPORT s::cl_double length(s::cl_double4 p) { return __length(p); } +__SYCL_EXPORT s::cl_half length(s::cl_half2 p) { return __length(p); } +__SYCL_EXPORT s::cl_half length(s::cl_half3 p) { return __length(p); } +__SYCL_EXPORT s::cl_half length(s::cl_half4 p) { return __length(p); } // distance -s::cl_float distance(s::cl_float p0, s::cl_float p1) { return length(p0 - p1); } -s::cl_float distance(s::cl_float2 p0, s::cl_float2 p1) { +__SYCL_EXPORT s::cl_float distance(s::cl_float p0, s::cl_float p1) { + return length(p0 - p1); +} +__SYCL_EXPORT s::cl_float distance(s::cl_float2 p0, s::cl_float2 p1) { + return length(p0 - p1); +} +__SYCL_EXPORT s::cl_float distance(s::cl_float3 p0, s::cl_float3 p1) { + return length(p0 - p1); +} +__SYCL_EXPORT s::cl_float distance(s::cl_float4 p0, s::cl_float4 p1) { + return length(p0 - p1); +} +__SYCL_EXPORT s::cl_double distance(s::cl_double p0, s::cl_double p1) { return length(p0 - p1); } -s::cl_float distance(s::cl_float3 p0, s::cl_float3 p1) { +__SYCL_EXPORT s::cl_double distance(s::cl_double2 p0, s::cl_double2 p1) { return length(p0 - p1); } -s::cl_float distance(s::cl_float4 p0, s::cl_float4 p1) { +__SYCL_EXPORT s::cl_double distance(s::cl_double3 p0, s::cl_double3 p1) { return length(p0 - p1); } -s::cl_double distance(s::cl_double p0, s::cl_double p1) { +__SYCL_EXPORT s::cl_double distance(s::cl_double4 p0, s::cl_double4 p1) { return length(p0 - p1); } -s::cl_double distance(s::cl_double2 p0, s::cl_double2 p1) { +__SYCL_EXPORT s::cl_half distance(s::cl_half p0, s::cl_half p1) { return length(p0 - p1); } -s::cl_double distance(s::cl_double3 p0, s::cl_double3 p1) { +__SYCL_EXPORT s::cl_half distance(s::cl_half2 p0, s::cl_half2 p1) { return length(p0 - p1); } -s::cl_double distance(s::cl_double4 p0, s::cl_double4 p1) { +__SYCL_EXPORT s::cl_half distance(s::cl_half3 p0, s::cl_half3 p1) { + return length(p0 - p1); +} +__SYCL_EXPORT s::cl_half distance(s::cl_half4 p0, s::cl_half4 p1) { return length(p0 - p1); } -s::cl_half distance(s::cl_half p0, s::cl_half p1) { return length(p0 - p1); } -s::cl_half distance(s::cl_half2 p0, s::cl_half2 p1) { return length(p0 - p1); } -s::cl_half distance(s::cl_half3 p0, s::cl_half3 p1) { return length(p0 - p1); } -s::cl_half distance(s::cl_half4 p0, s::cl_half4 p1) { return length(p0 - p1); } // normalize -s::cl_float normalize(s::cl_float p) { return __normalize(p); } -s::cl_float2 normalize(s::cl_float2 p) { return __normalize(p); } -s::cl_float3 normalize(s::cl_float3 p) { return __normalize(p); } -s::cl_float4 normalize(s::cl_float4 p) { return __normalize(p); } -s::cl_double normalize(s::cl_double p) { return __normalize(p); } -s::cl_double2 normalize(s::cl_double2 p) { return __normalize(p); } -s::cl_double3 normalize(s::cl_double3 p) { return __normalize(p); } -s::cl_double4 normalize(s::cl_double4 p) { return __normalize(p); } -s::cl_half normalize(s::cl_half p) { return __normalize(p); } -s::cl_half2 normalize(s::cl_half2 p) { return __normalize(p); } -s::cl_half3 normalize(s::cl_half3 p) { return __normalize(p); } -s::cl_half4 normalize(s::cl_half4 p) { return __normalize(p); } +__SYCL_EXPORT s::cl_float normalize(s::cl_float p) { return __normalize(p); } +__SYCL_EXPORT s::cl_float2 normalize(s::cl_float2 p) { return __normalize(p); } +__SYCL_EXPORT s::cl_float3 normalize(s::cl_float3 p) { return __normalize(p); } +__SYCL_EXPORT s::cl_float4 normalize(s::cl_float4 p) { return __normalize(p); } +__SYCL_EXPORT s::cl_double normalize(s::cl_double p) { return __normalize(p); } +__SYCL_EXPORT s::cl_double2 normalize(s::cl_double2 p) { + return __normalize(p); +} +__SYCL_EXPORT s::cl_double3 normalize(s::cl_double3 p) { + return __normalize(p); +} +__SYCL_EXPORT s::cl_double4 normalize(s::cl_double4 p) { + return __normalize(p); +} +__SYCL_EXPORT s::cl_half normalize(s::cl_half p) { return __normalize(p); } +__SYCL_EXPORT s::cl_half2 normalize(s::cl_half2 p) { return __normalize(p); } +__SYCL_EXPORT s::cl_half3 normalize(s::cl_half3 p) { return __normalize(p); } +__SYCL_EXPORT s::cl_half4 normalize(s::cl_half4 p) { return __normalize(p); } // fast_length -s::cl_float fast_length(s::cl_float p) { return __fast_length(p); } -s::cl_float fast_length(s::cl_float2 p) { return __fast_length(p); } -s::cl_float fast_length(s::cl_float3 p) { return __fast_length(p); } -s::cl_float fast_length(s::cl_float4 p) { return __fast_length(p); } +__SYCL_EXPORT s::cl_float fast_length(s::cl_float p) { + return __fast_length(p); +} +__SYCL_EXPORT s::cl_float fast_length(s::cl_float2 p) { + return __fast_length(p); +} +__SYCL_EXPORT s::cl_float fast_length(s::cl_float3 p) { + return __fast_length(p); +} +__SYCL_EXPORT s::cl_float fast_length(s::cl_float4 p) { + return __fast_length(p); +} // fast_normalize -s::cl_float fast_normalize(s::cl_float p) { +__SYCL_EXPORT s::cl_float fast_normalize(s::cl_float p) { if (p == 0.0f) return p; s::cl_float r = std::sqrt(FMul(p, p)); return p / r; } -s::cl_float2 fast_normalize(s::cl_float2 p) { return __fast_normalize(p); } -s::cl_float3 fast_normalize(s::cl_float3 p) { return __fast_normalize(p); } -s::cl_float4 fast_normalize(s::cl_float4 p) { return __fast_normalize(p); } +__SYCL_EXPORT s::cl_float2 fast_normalize(s::cl_float2 p) { + return __fast_normalize(p); +} +__SYCL_EXPORT s::cl_float3 fast_normalize(s::cl_float3 p) { + return __fast_normalize(p); +} +__SYCL_EXPORT s::cl_float4 fast_normalize(s::cl_float4 p) { + return __fast_normalize(p); +} // fast_distance -s::cl_float fast_distance(s::cl_float p0, s::cl_float p1) { +__SYCL_EXPORT s::cl_float fast_distance(s::cl_float p0, s::cl_float p1) { return fast_length(p0 - p1); } -s::cl_float fast_distance(s::cl_float2 p0, s::cl_float2 p1) { +__SYCL_EXPORT s::cl_float fast_distance(s::cl_float2 p0, s::cl_float2 p1) { return fast_length(p0 - p1); } -s::cl_float fast_distance(s::cl_float3 p0, s::cl_float3 p1) { +__SYCL_EXPORT s::cl_float fast_distance(s::cl_float3 p0, s::cl_float3 p1) { return fast_length(p0 - p1); } -s::cl_float fast_distance(s::cl_float4 p0, s::cl_float4 p1) { +__SYCL_EXPORT s::cl_float fast_distance(s::cl_float4 p0, s::cl_float4 p1) { return fast_length(p0 - p1); } diff --git a/sycl/source/detail/builtins_helper.hpp b/sycl/source/detail/builtins_helper.hpp index 948182a65475a..0de1b33c19f95 100644 --- a/sycl/source/detail/builtins_helper.hpp +++ b/sycl/source/detail/builtins_helper.hpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// +#include #include #include #include @@ -15,110 +16,97 @@ #define __NOEXC /*noexcept*/ #define __MAKE_1V(Fun, Call, N, Ret, Arg1) \ - Ret##N Fun __NOEXC(Arg1##N x) { \ + __SYCL_EXPORT Ret##N Fun __NOEXC(Arg1##N x) { \ Ret##N r; \ detail::helper().run_1v( \ - r, [](Arg1 x) { \ - return cl::__host_std::Call(x); }, x); \ + r, [](Arg1 x) { return cl::__host_std::Call(x); }, x); \ return r; \ } #define __MAKE_1V_2V(Fun, Call, N, Ret, Arg1, Arg2) \ - Ret##N Fun __NOEXC(Arg1##N x, Arg2##N y) { \ + __SYCL_EXPORT Ret##N Fun __NOEXC(Arg1##N x, Arg2##N y) { \ Ret##N r; \ detail::helper().run_1v_2v( \ - r, [](Arg1 x, Arg2 y) { \ - return cl::__host_std::Call(x, y); }, x, y); \ + r, [](Arg1 x, Arg2 y) { return cl::__host_std::Call(x, y); }, x, y); \ return r; \ } #define __MAKE_1V_2V_RS(Fun, Call, N, Ret, Arg1, Arg2) \ - Ret Fun __NOEXC(Arg1##N x, Arg2##N y) { \ + __SYCL_EXPORT Ret Fun __NOEXC(Arg1##N x, Arg2##N y) { \ Ret r = Ret(); \ detail::helper().run_1v_2v_rs( \ r, \ - [](Ret &r, Arg1 x, Arg2 y) { \ - return cl::__host_std::Call(r, x, y); \ - }, \ + [](Ret &r, Arg1 x, Arg2 y) { return cl::__host_std::Call(r, x, y); }, \ x, y); \ return r; \ } #define __MAKE_1V_RS(Fun, Call, N, Ret, Arg1) \ - Ret Fun __NOEXC(Arg1##N x) { \ + __SYCL_EXPORT Ret Fun __NOEXC(Arg1##N x) { \ Ret r = Ret(); \ detail::helper().run_1v_rs( \ - r, [](Ret &r, Arg1 x) { \ - return cl::__host_std::Call(r, x); }, x); \ + r, [](Ret &r, Arg1 x) { return cl::__host_std::Call(r, x); }, x); \ return r; \ } #define __MAKE_1V_2V_3V(Fun, Call, N, Ret, Arg1, Arg2, Arg3) \ - Ret##N Fun __NOEXC(Arg1##N x, Arg2##N y, Arg3##N z) { \ + __SYCL_EXPORT Ret##N Fun __NOEXC(Arg1##N x, Arg2##N y, Arg3##N z) { \ Ret##N r; \ detail::helper().run_1v_2v_3v( \ r, \ - [](Arg1 x, Arg2 y, Arg3 z) { \ - return cl::__host_std::Call(x, y, z); \ - }, \ + [](Arg1 x, Arg2 y, Arg3 z) { return cl::__host_std::Call(x, y, z); }, \ x, y, z); \ return r; \ } #define __MAKE_1V_2S_3S(Fun, N, Ret, Arg1, Arg2, Arg3) \ - Ret##N Fun __NOEXC(Arg1##N x, Arg2 y, Arg3 z) { \ + __SYCL_EXPORT Ret##N Fun __NOEXC(Arg1##N x, Arg2 y, Arg3 z) { \ Ret##N r; \ detail::helper().run_1v_2s_3s( \ r, \ - [](Arg1 x, Arg2 y, Arg3 z) { \ - return cl::__host_std::Fun(x, y, z); \ - }, \ + [](Arg1 x, Arg2 y, Arg3 z) { return cl::__host_std::Fun(x, y, z); }, \ x, y, z); \ return r; \ } #define __MAKE_1V_2S(Fun, N, Ret, Arg1, Arg2) \ - Ret##N Fun __NOEXC(Arg1##N x, Arg2 y) { \ + __SYCL_EXPORT Ret##N Fun __NOEXC(Arg1##N x, Arg2 y) { \ Ret##N r; \ detail::helper().run_1v_2s( \ - r, [](Arg1 x, Arg2 y) { return cl::__host_std::Fun(x, y); }, \ - x, y); \ + r, [](Arg1 x, Arg2 y) { return cl::__host_std::Fun(x, y); }, x, y); \ return r; \ } #define __MAKE_SR_1V_AND(Fun, Call, N, Ret, Arg1) \ - Ret Fun __NOEXC(Arg1##N x) { \ + __SYCL_EXPORT Ret Fun __NOEXC(Arg1##N x) { \ Ret r; \ detail::helper().run_1v_sr_and( \ - r, [](Arg1 x) { return cl::__host_std::Call(x); }, x); \ + r, [](Arg1 x) { return cl::__host_std::Call(x); }, x); \ return r; \ } #define __MAKE_SR_1V_OR(Fun, Call, N, Ret, Arg1) \ - Ret Fun __NOEXC(Arg1##N x) { \ + __SYCL_EXPORT Ret Fun __NOEXC(Arg1##N x) { \ Ret r; \ detail::helper().run_1v_sr_or( \ - r, [](Arg1 x) { return cl::__host_std::Call(x); }, x); \ + r, [](Arg1 x) { return cl::__host_std::Call(x); }, x); \ return r; \ } #define __MAKE_1V_2P(Fun, N, Ret, Arg1, Arg2) \ - Ret##N Fun __NOEXC(Arg1##N x, Arg2##N *y) { \ + __SYCL_EXPORT Ret##N Fun __NOEXC(Arg1##N x, Arg2##N *y) { \ Ret##N r; \ detail::helper().run_1v_2p( \ - r, [](Arg1 x, Arg2 *y) { \ - return cl::__host_std::Fun(x, y); }, x, y); \ + r, [](Arg1 x, Arg2 *y) { return cl::__host_std::Fun(x, y); }, x, y); \ return r; \ } #define __MAKE_1V_2V_3P(Fun, N, Ret, Arg1, Arg2, Arg3) \ - Ret##N Fun __NOEXC(Arg1##N x, Arg2##N y, Arg3##N *z) { \ + __SYCL_EXPORT Ret##N Fun __NOEXC(Arg1##N x, Arg2##N y, Arg3##N *z) { \ Ret##N r; \ detail::helper().run_1v_2v_3p( \ r, \ - [](Arg1 x, Arg2 y, Arg3 *z) { \ - return cl::__host_std::Fun(x, y, z); \ - }, \ + [](Arg1 x, Arg2 y, Arg3 *z) { return cl::__host_std::Fun(x, y, z); }, \ x, y, z); \ return r; \ } @@ -127,8 +115,9 @@ #define MAKE_1V_FUNC(Fun, Call, Ret, Arg1) \ __MAKE_1V(Fun, Call, 2, Ret, Arg1) \ - __MAKE_1V(Fun, Call, 3, Ret, Arg1) __MAKE_1V(Fun, Call, 4, Ret, Arg1) \ - __MAKE_1V(Fun, Call, 8, Ret, Arg1) __MAKE_1V(Fun, Call, 16, Ret, Arg1) + __MAKE_1V(Fun, Call, 3, Ret, Arg1) \ + __MAKE_1V(Fun, Call, 4, Ret, Arg1) \ + __MAKE_1V(Fun, Call, 8, Ret, Arg1) __MAKE_1V(Fun, Call, 16, Ret, Arg1) #define MAKE_1V_2V(Fun, Ret, Arg1, Arg2) \ MAKE_1V_2V_FUNC(Fun, Fun, Ret, Arg1, Arg2) @@ -159,13 +148,15 @@ MAKE_1V_2V_3V_FUNC(FunSc, FunV, Ret, Arg1, Arg2, Arg3) #define MAKE_SC_3ARG(Fun, Ret, Arg1, Arg2, Arg3) \ - Ret Fun __NOEXC(Arg1 x, Arg2 y, Arg3 z) { return (Ret)__##Fun(x, y, z); } + __SYCL_EXPORT Ret Fun __NOEXC(Arg1 x, Arg2 y, Arg3 z) { \ + return (Ret)__##Fun(x, y, z); \ + } #define MAKE_1V_2S(Fun, Ret, Arg1, Arg2) \ __MAKE_1V_2S(Fun, 2, Ret, Arg1, Arg2) \ - __MAKE_1V_2S(Fun, 3, Ret, Arg1, Arg2) __MAKE_1V_2S(Fun, 4, Ret, Arg1, Arg2) \ - __MAKE_1V_2S(Fun, 8, Ret, Arg1, Arg2) \ - __MAKE_1V_2S(Fun, 16, Ret, Arg1, Arg2) + __MAKE_1V_2S(Fun, 3, Ret, Arg1, Arg2) \ + __MAKE_1V_2S(Fun, 4, Ret, Arg1, Arg2) \ + __MAKE_1V_2S(Fun, 8, Ret, Arg1, Arg2) __MAKE_1V_2S(Fun, 16, Ret, Arg1, Arg2) #define MAKE_1V_2S_3S(Fun, Ret, Arg1, Arg2, Arg3) \ __MAKE_1V_2S_3S(Fun, 2, Ret, Arg1, Arg2, Arg3) \ @@ -190,9 +181,9 @@ #define MAKE_1V_2P(Fun, Ret, Arg1, Arg2) \ __MAKE_1V_2P(Fun, 2, Ret, Arg1, Arg2) \ - __MAKE_1V_2P(Fun, 3, Ret, Arg1, Arg2) __MAKE_1V_2P(Fun, 4, Ret, Arg1, Arg2) \ - __MAKE_1V_2P(Fun, 8, Ret, Arg1, Arg2) \ - __MAKE_1V_2P(Fun, 16, Ret, Arg1, Arg2) + __MAKE_1V_2P(Fun, 3, Ret, Arg1, Arg2) \ + __MAKE_1V_2P(Fun, 4, Ret, Arg1, Arg2) \ + __MAKE_1V_2P(Fun, 8, Ret, Arg1, Arg2) __MAKE_1V_2P(Fun, 16, Ret, Arg1, Arg2) #define MAKE_GEO_1V_2V_RS(Fun, Call, Ret, Arg1, Arg2) \ __MAKE_1V_2V_RS(Fun, Call, 2, Ret, Arg1, Arg2) \ diff --git a/sycl/source/detail/builtins_integer.cpp b/sycl/source/detail/builtins_integer.cpp index 739d429524d96..15c9a63243e3c 100644 --- a/sycl/source/detail/builtins_integer.cpp +++ b/sycl/source/detail/builtins_integer.cpp @@ -10,6 +10,7 @@ // in SYCL SPEC section - 4.13.4 Integer functions. #include "builtins_helper.hpp" +#include #include #include @@ -237,36 +238,36 @@ template inline T __mul24(T x, T y) { return (x * y); } // --------------- 4.13.4 Integer functions. Host implementations -------------- // u_abs -s::cl_uchar u_abs(s::cl_uchar x) __NOEXC { return x; } -s::cl_ushort u_abs(s::cl_ushort x) __NOEXC { return x; } -s::cl_uint u_abs(s::cl_uint x) __NOEXC { return x; } -s::cl_ulong u_abs(s::cl_ulong x) __NOEXC { return x; } +__SYCL_EXPORT s::cl_uchar u_abs(s::cl_uchar x) __NOEXC { return x; } +__SYCL_EXPORT s::cl_ushort u_abs(s::cl_ushort x) __NOEXC { return x; } +__SYCL_EXPORT s::cl_uint u_abs(s::cl_uint x) __NOEXC { return x; } +__SYCL_EXPORT s::cl_ulong u_abs(s::cl_ulong x) __NOEXC { return x; } MAKE_1V(u_abs, s::cl_uchar, s::cl_uchar) MAKE_1V(u_abs, s::cl_ushort, s::cl_ushort) MAKE_1V(u_abs, s::cl_uint, s::cl_uint) MAKE_1V(u_abs, s::cl_ulong, s::cl_ulong) // s_abs -s::cl_uchar s_abs(s::cl_char x) __NOEXC { return std::abs(x); } -s::cl_ushort s_abs(s::cl_short x) __NOEXC { return std::abs(x); } -s::cl_uint s_abs(s::cl_int x) __NOEXC { return std::abs(x); } -s::cl_ulong s_abs(s::cl_long x) __NOEXC { return std::abs(x); } +__SYCL_EXPORT s::cl_uchar s_abs(s::cl_char x) __NOEXC { return std::abs(x); } +__SYCL_EXPORT s::cl_ushort s_abs(s::cl_short x) __NOEXC { return std::abs(x); } +__SYCL_EXPORT s::cl_uint s_abs(s::cl_int x) __NOEXC { return std::abs(x); } +__SYCL_EXPORT s::cl_ulong s_abs(s::cl_long x) __NOEXC { return std::abs(x); } MAKE_1V(s_abs, s::cl_uchar, s::cl_char) MAKE_1V(s_abs, s::cl_ushort, s::cl_short) MAKE_1V(s_abs, s::cl_uint, s::cl_int) MAKE_1V(s_abs, s::cl_ulong, s::cl_long) // u_abs_diff -s::cl_uchar u_abs_diff(s::cl_uchar x, s::cl_uchar y) __NOEXC { +__SYCL_EXPORT s::cl_uchar u_abs_diff(s::cl_uchar x, s::cl_uchar y) __NOEXC { return __abs_diff(x, y); } -s::cl_ushort u_abs_diff(s::cl_ushort x, s::cl_ushort y) __NOEXC { +__SYCL_EXPORT s::cl_ushort u_abs_diff(s::cl_ushort x, s::cl_ushort y) __NOEXC { return __abs_diff(x, y); } -s::cl_uint u_abs_diff(s::cl_uint x, s::cl_uint y) __NOEXC { +__SYCL_EXPORT s::cl_uint u_abs_diff(s::cl_uint x, s::cl_uint y) __NOEXC { return __abs_diff(x, y); } -s::cl_ulong u_abs_diff(s::cl_ulong x, s::cl_ulong y) __NOEXC { +__SYCL_EXPORT s::cl_ulong u_abs_diff(s::cl_ulong x, s::cl_ulong y) __NOEXC { return __abs_diff(x, y); } @@ -276,16 +277,16 @@ MAKE_1V_2V(u_abs_diff, s::cl_uint, s::cl_uint, s::cl_uint) MAKE_1V_2V(u_abs_diff, s::cl_ulong, s::cl_ulong, s::cl_ulong) // s_abs_diff -s::cl_uchar s_abs_diff(s::cl_char x, s::cl_char y) __NOEXC { +__SYCL_EXPORT s::cl_uchar s_abs_diff(s::cl_char x, s::cl_char y) __NOEXC { return __abs_diff(x, y); } -s::cl_ushort s_abs_diff(s::cl_short x, s::cl_short y) __NOEXC { +__SYCL_EXPORT s::cl_ushort s_abs_diff(s::cl_short x, s::cl_short y) __NOEXC { return __abs_diff(x, y); } -s::cl_uint s_abs_diff(s::cl_int x, s::cl_int y) __NOEXC { +__SYCL_EXPORT s::cl_uint s_abs_diff(s::cl_int x, s::cl_int y) __NOEXC { return __abs_diff(x, y); } -s::cl_ulong s_abs_diff(s::cl_long x, s::cl_long y) __NOEXC { +__SYCL_EXPORT s::cl_ulong s_abs_diff(s::cl_long x, s::cl_long y) __NOEXC { return __abs_diff(x, y); } MAKE_1V_2V(s_abs_diff, s::cl_uchar, s::cl_char, s::cl_char) @@ -294,16 +295,16 @@ MAKE_1V_2V(s_abs_diff, s::cl_uint, s::cl_int, s::cl_int) MAKE_1V_2V(s_abs_diff, s::cl_ulong, s::cl_long, s::cl_long) // u_add_sat -s::cl_uchar u_add_sat(s::cl_uchar x, s::cl_uchar y) __NOEXC { +__SYCL_EXPORT s::cl_uchar u_add_sat(s::cl_uchar x, s::cl_uchar y) __NOEXC { return __u_add_sat(x, y); } -s::cl_ushort u_add_sat(s::cl_ushort x, s::cl_ushort y) __NOEXC { +__SYCL_EXPORT s::cl_ushort u_add_sat(s::cl_ushort x, s::cl_ushort y) __NOEXC { return __u_add_sat(x, y); } -s::cl_uint u_add_sat(s::cl_uint x, s::cl_uint y) __NOEXC { +__SYCL_EXPORT s::cl_uint u_add_sat(s::cl_uint x, s::cl_uint y) __NOEXC { return __u_add_sat(x, y); } -s::cl_ulong u_add_sat(s::cl_ulong x, s::cl_ulong y) __NOEXC { +__SYCL_EXPORT s::cl_ulong u_add_sat(s::cl_ulong x, s::cl_ulong y) __NOEXC { return __u_add_sat(x, y); } MAKE_1V_2V(u_add_sat, s::cl_uchar, s::cl_uchar, s::cl_uchar) @@ -312,16 +313,16 @@ MAKE_1V_2V(u_add_sat, s::cl_uint, s::cl_uint, s::cl_uint) MAKE_1V_2V(u_add_sat, s::cl_ulong, s::cl_ulong, s::cl_ulong) // s_add_sat -s::cl_char s_add_sat(s::cl_char x, s::cl_char y) __NOEXC { +__SYCL_EXPORT s::cl_char s_add_sat(s::cl_char x, s::cl_char y) __NOEXC { return __s_add_sat(x, y); } -s::cl_short s_add_sat(s::cl_short x, s::cl_short y) __NOEXC { +__SYCL_EXPORT s::cl_short s_add_sat(s::cl_short x, s::cl_short y) __NOEXC { return __s_add_sat(x, y); } -s::cl_int s_add_sat(s::cl_int x, s::cl_int y) __NOEXC { +__SYCL_EXPORT s::cl_int s_add_sat(s::cl_int x, s::cl_int y) __NOEXC { return __s_add_sat(x, y); } -s::cl_long s_add_sat(s::cl_long x, s::cl_long y) __NOEXC { +__SYCL_EXPORT s::cl_long s_add_sat(s::cl_long x, s::cl_long y) __NOEXC { return __s_add_sat(x, y); } MAKE_1V_2V(s_add_sat, s::cl_char, s::cl_char, s::cl_char) @@ -330,14 +331,16 @@ MAKE_1V_2V(s_add_sat, s::cl_int, s::cl_int, s::cl_int) MAKE_1V_2V(s_add_sat, s::cl_long, s::cl_long, s::cl_long) // u_hadd -s::cl_uchar u_hadd(s::cl_uchar x, s::cl_uchar y) __NOEXC { +__SYCL_EXPORT s::cl_uchar u_hadd(s::cl_uchar x, s::cl_uchar y) __NOEXC { return __hadd(x, y); } -s::cl_ushort u_hadd(s::cl_ushort x, s::cl_ushort y) __NOEXC { +__SYCL_EXPORT s::cl_ushort u_hadd(s::cl_ushort x, s::cl_ushort y) __NOEXC { return __hadd(x, y); } -s::cl_uint u_hadd(s::cl_uint x, s::cl_uint y) __NOEXC { return __hadd(x, y); } -s::cl_ulong u_hadd(s::cl_ulong x, s::cl_ulong y) __NOEXC { +__SYCL_EXPORT s::cl_uint u_hadd(s::cl_uint x, s::cl_uint y) __NOEXC { + return __hadd(x, y); +} +__SYCL_EXPORT s::cl_ulong u_hadd(s::cl_ulong x, s::cl_ulong y) __NOEXC { return __hadd(x, y); } MAKE_1V_2V(u_hadd, s::cl_uchar, s::cl_uchar, s::cl_uchar) @@ -346,26 +349,34 @@ MAKE_1V_2V(u_hadd, s::cl_uint, s::cl_uint, s::cl_uint) MAKE_1V_2V(u_hadd, s::cl_ulong, s::cl_ulong, s::cl_ulong) // s_hadd -s::cl_char s_hadd(s::cl_char x, s::cl_char y) __NOEXC { return __hadd(x, y); } -s::cl_short s_hadd(s::cl_short x, s::cl_short y) __NOEXC { +__SYCL_EXPORT s::cl_char s_hadd(s::cl_char x, s::cl_char y) __NOEXC { + return __hadd(x, y); +} +__SYCL_EXPORT s::cl_short s_hadd(s::cl_short x, s::cl_short y) __NOEXC { + return __hadd(x, y); +} +__SYCL_EXPORT s::cl_int s_hadd(s::cl_int x, s::cl_int y) __NOEXC { + return __hadd(x, y); +} +__SYCL_EXPORT s::cl_long s_hadd(s::cl_long x, s::cl_long y) __NOEXC { return __hadd(x, y); } -s::cl_int s_hadd(s::cl_int x, s::cl_int y) __NOEXC { return __hadd(x, y); } -s::cl_long s_hadd(s::cl_long x, s::cl_long y) __NOEXC { return __hadd(x, y); } MAKE_1V_2V(s_hadd, s::cl_char, s::cl_char, s::cl_char) MAKE_1V_2V(s_hadd, s::cl_short, s::cl_short, s::cl_short) MAKE_1V_2V(s_hadd, s::cl_int, s::cl_int, s::cl_int) MAKE_1V_2V(s_hadd, s::cl_long, s::cl_long, s::cl_long) // u_rhadd -s::cl_uchar u_rhadd(s::cl_uchar x, s::cl_uchar y) __NOEXC { +__SYCL_EXPORT s::cl_uchar u_rhadd(s::cl_uchar x, s::cl_uchar y) __NOEXC { + return __rhadd(x, y); +} +__SYCL_EXPORT s::cl_ushort u_rhadd(s::cl_ushort x, s::cl_ushort y) __NOEXC { return __rhadd(x, y); } -s::cl_ushort u_rhadd(s::cl_ushort x, s::cl_ushort y) __NOEXC { +__SYCL_EXPORT s::cl_uint u_rhadd(s::cl_uint x, s::cl_uint y) __NOEXC { return __rhadd(x, y); } -s::cl_uint u_rhadd(s::cl_uint x, s::cl_uint y) __NOEXC { return __rhadd(x, y); } -s::cl_ulong u_rhadd(s::cl_ulong x, s::cl_ulong y) __NOEXC { +__SYCL_EXPORT s::cl_ulong u_rhadd(s::cl_ulong x, s::cl_ulong y) __NOEXC { return __rhadd(x, y); } MAKE_1V_2V(u_rhadd, s::cl_uchar, s::cl_uchar, s::cl_uchar) @@ -374,31 +385,38 @@ MAKE_1V_2V(u_rhadd, s::cl_uint, s::cl_uint, s::cl_uint) MAKE_1V_2V(u_rhadd, s::cl_ulong, s::cl_ulong, s::cl_ulong) // s_rhadd -s::cl_char s_rhadd(s::cl_char x, s::cl_char y) __NOEXC { return __rhadd(x, y); } -s::cl_short s_rhadd(s::cl_short x, s::cl_short y) __NOEXC { +__SYCL_EXPORT s::cl_char s_rhadd(s::cl_char x, s::cl_char y) __NOEXC { + return __rhadd(x, y); +} +__SYCL_EXPORT s::cl_short s_rhadd(s::cl_short x, s::cl_short y) __NOEXC { + return __rhadd(x, y); +} +__SYCL_EXPORT s::cl_int s_rhadd(s::cl_int x, s::cl_int y) __NOEXC { + return __rhadd(x, y); +} +__SYCL_EXPORT s::cl_long s_rhadd(s::cl_long x, s::cl_long y) __NOEXC { return __rhadd(x, y); } -s::cl_int s_rhadd(s::cl_int x, s::cl_int y) __NOEXC { return __rhadd(x, y); } -s::cl_long s_rhadd(s::cl_long x, s::cl_long y) __NOEXC { return __rhadd(x, y); } MAKE_1V_2V(s_rhadd, s::cl_char, s::cl_char, s::cl_char) MAKE_1V_2V(s_rhadd, s::cl_short, s::cl_short, s::cl_short) MAKE_1V_2V(s_rhadd, s::cl_int, s::cl_int, s::cl_int) MAKE_1V_2V(s_rhadd, s::cl_long, s::cl_long, s::cl_long) // u_clamp -s::cl_uchar u_clamp(s::cl_uchar x, s::cl_uchar minval, - s::cl_uchar maxval) __NOEXC { +__SYCL_EXPORT s::cl_uchar u_clamp(s::cl_uchar x, s::cl_uchar minval, + s::cl_uchar maxval) __NOEXC { return __clamp(x, minval, maxval); } -s::cl_ushort u_clamp(s::cl_ushort x, s::cl_ushort minval, - s::cl_ushort maxval) __NOEXC { +__SYCL_EXPORT s::cl_ushort u_clamp(s::cl_ushort x, s::cl_ushort minval, + s::cl_ushort maxval) __NOEXC { return __clamp(x, minval, maxval); } -s::cl_uint u_clamp(s::cl_uint x, s::cl_uint minval, s::cl_uint maxval) __NOEXC { +__SYCL_EXPORT s::cl_uint u_clamp(s::cl_uint x, s::cl_uint minval, + s::cl_uint maxval) __NOEXC { return __clamp(x, minval, maxval); } -s::cl_ulong u_clamp(s::cl_ulong x, s::cl_ulong minval, - s::cl_ulong maxval) __NOEXC { +__SYCL_EXPORT s::cl_ulong u_clamp(s::cl_ulong x, s::cl_ulong minval, + s::cl_ulong maxval) __NOEXC { return __clamp(x, minval, maxval); } MAKE_1V_2V_3V(u_clamp, s::cl_uchar, s::cl_uchar, s::cl_uchar, s::cl_uchar) @@ -411,17 +429,20 @@ MAKE_1V_2S_3S(u_clamp, s::cl_uint, s::cl_uint, s::cl_uint, s::cl_uint) MAKE_1V_2S_3S(u_clamp, s::cl_ulong, s::cl_ulong, s::cl_ulong, s::cl_ulong) // s_clamp -s::cl_char s_clamp(s::cl_char x, s::cl_char minval, s::cl_char maxval) __NOEXC { +__SYCL_EXPORT s::cl_char s_clamp(s::cl_char x, s::cl_char minval, + s::cl_char maxval) __NOEXC { return __clamp(x, minval, maxval); } -s::cl_short s_clamp(s::cl_short x, s::cl_short minval, - s::cl_short maxval) __NOEXC { +__SYCL_EXPORT s::cl_short s_clamp(s::cl_short x, s::cl_short minval, + s::cl_short maxval) __NOEXC { return __clamp(x, minval, maxval); } -s::cl_int s_clamp(s::cl_int x, s::cl_int minval, s::cl_int maxval) __NOEXC { +__SYCL_EXPORT s::cl_int s_clamp(s::cl_int x, s::cl_int minval, + s::cl_int maxval) __NOEXC { return __clamp(x, minval, maxval); } -s::cl_long s_clamp(s::cl_long x, s::cl_long minval, s::cl_long maxval) __NOEXC { +__SYCL_EXPORT s::cl_long s_clamp(s::cl_long x, s::cl_long minval, + s::cl_long maxval) __NOEXC { return __clamp(x, minval, maxval); } MAKE_1V_2V_3V(s_clamp, s::cl_char, s::cl_char, s::cl_char, s::cl_char) @@ -434,14 +455,14 @@ MAKE_1V_2S_3S(s_clamp, s::cl_int, s::cl_int, s::cl_int, s::cl_int) MAKE_1V_2S_3S(s_clamp, s::cl_long, s::cl_long, s::cl_long, s::cl_long) // clz -s::cl_uchar clz(s::cl_uchar x) __NOEXC { return __clz(x); } -s::cl_char clz(s::cl_char x) __NOEXC { return __clz(x); } -s::cl_ushort clz(s::cl_ushort x) __NOEXC { return __clz(x); } -s::cl_short clz(s::cl_short x) __NOEXC { return __clz(x); } -s::cl_uint clz(s::cl_uint x) __NOEXC { return __clz(x); } -s::cl_int clz(s::cl_int x) __NOEXC { return __clz(x); } -s::cl_ulong clz(s::cl_ulong x) __NOEXC { return __clz(x); } -s::cl_long clz(s::cl_long x) __NOEXC { return __clz(x); } +__SYCL_EXPORT s::cl_uchar clz(s::cl_uchar x) __NOEXC { return __clz(x); } +__SYCL_EXPORT s::cl_char clz(s::cl_char x) __NOEXC { return __clz(x); } +__SYCL_EXPORT s::cl_ushort clz(s::cl_ushort x) __NOEXC { return __clz(x); } +__SYCL_EXPORT s::cl_short clz(s::cl_short x) __NOEXC { return __clz(x); } +__SYCL_EXPORT s::cl_uint clz(s::cl_uint x) __NOEXC { return __clz(x); } +__SYCL_EXPORT s::cl_int clz(s::cl_int x) __NOEXC { return __clz(x); } +__SYCL_EXPORT s::cl_ulong clz(s::cl_ulong x) __NOEXC { return __clz(x); } +__SYCL_EXPORT s::cl_long clz(s::cl_long x) __NOEXC { return __clz(x); } MAKE_1V(clz, s::cl_uchar, s::cl_uchar) MAKE_1V(clz, s::cl_char, s::cl_char) MAKE_1V(clz, s::cl_ushort, s::cl_ushort) @@ -452,14 +473,14 @@ MAKE_1V(clz, s::cl_ulong, s::cl_ulong) MAKE_1V(clz, s::cl_long, s::cl_long) // ctz -s::cl_uchar ctz(s::cl_uchar x) __NOEXC { return __ctz(x); } -s::cl_char ctz(s::cl_char x) __NOEXC { return __ctz(x); } -s::cl_ushort ctz(s::cl_ushort x) __NOEXC { return __ctz(x); } -s::cl_short ctz(s::cl_short x) __NOEXC { return __ctz(x); } -s::cl_uint ctz(s::cl_uint x) __NOEXC { return __ctz(x); } -s::cl_int ctz(s::cl_int x) __NOEXC { return __ctz(x); } -s::cl_ulong ctz(s::cl_ulong x) __NOEXC { return __ctz(x); } -s::cl_long ctz(s::cl_long x) __NOEXC { return __ctz(x); } +__SYCL_EXPORT s::cl_uchar ctz(s::cl_uchar x) __NOEXC { return __ctz(x); } +__SYCL_EXPORT s::cl_char ctz(s::cl_char x) __NOEXC { return __ctz(x); } +__SYCL_EXPORT s::cl_ushort ctz(s::cl_ushort x) __NOEXC { return __ctz(x); } +__SYCL_EXPORT s::cl_short ctz(s::cl_short x) __NOEXC { return __ctz(x); } +__SYCL_EXPORT s::cl_uint ctz(s::cl_uint x) __NOEXC { return __ctz(x); } +__SYCL_EXPORT s::cl_int ctz(s::cl_int x) __NOEXC { return __ctz(x); } +__SYCL_EXPORT s::cl_ulong ctz(s::cl_ulong x) __NOEXC { return __ctz(x); } +__SYCL_EXPORT s::cl_long ctz(s::cl_long x) __NOEXC { return __ctz(x); } MAKE_1V(ctz, s::cl_uchar, s::cl_uchar) MAKE_1V(ctz, s::cl_char, s::cl_char) MAKE_1V(ctz, s::cl_ushort, s::cl_ushort) @@ -470,10 +491,16 @@ MAKE_1V(ctz, s::cl_ulong, s::cl_ulong) MAKE_1V(ctz, s::cl_long, s::cl_long) // s_mul_hi -s::cl_char s_mul_hi(s::cl_char a, s::cl_char b) { return __mul_hi(a, b); } -s::cl_short s_mul_hi(s::cl_short a, s::cl_short b) { return __mul_hi(a, b); } -s::cl_int s_mul_hi(s::cl_int a, s::cl_int b) { return __mul_hi(a, b); } -s::cl_long s_mul_hi(s::cl_long x, s::cl_long y) __NOEXC { +__SYCL_EXPORT s::cl_char s_mul_hi(s::cl_char a, s::cl_char b) { + return __mul_hi(a, b); +} +__SYCL_EXPORT s::cl_short s_mul_hi(s::cl_short a, s::cl_short b) { + return __mul_hi(a, b); +} +__SYCL_EXPORT s::cl_int s_mul_hi(s::cl_int a, s::cl_int b) { + return __mul_hi(a, b); +} +__SYCL_EXPORT s::cl_long s_mul_hi(s::cl_long x, s::cl_long y) __NOEXC { return __s_long_mul_hi(x, y); } MAKE_1V_2V(s_mul_hi, s::cl_char, s::cl_char, s::cl_char) @@ -482,10 +509,16 @@ MAKE_1V_2V(s_mul_hi, s::cl_int, s::cl_int, s::cl_int) MAKE_1V_2V(s_mul_hi, s::cl_long, s::cl_long, s::cl_long) // u_mul_hi -s::cl_uchar u_mul_hi(s::cl_uchar a, s::cl_uchar b) { return __mul_hi(a, b); } -s::cl_ushort u_mul_hi(s::cl_ushort a, s::cl_ushort b) { return __mul_hi(a, b); } -s::cl_uint u_mul_hi(s::cl_uint a, s::cl_uint b) { return __mul_hi(a, b); } -s::cl_ulong u_mul_hi(s::cl_ulong x, s::cl_ulong y) __NOEXC { +__SYCL_EXPORT s::cl_uchar u_mul_hi(s::cl_uchar a, s::cl_uchar b) { + return __mul_hi(a, b); +} +__SYCL_EXPORT s::cl_ushort u_mul_hi(s::cl_ushort a, s::cl_ushort b) { + return __mul_hi(a, b); +} +__SYCL_EXPORT s::cl_uint u_mul_hi(s::cl_uint a, s::cl_uint b) { + return __mul_hi(a, b); +} +__SYCL_EXPORT s::cl_ulong u_mul_hi(s::cl_ulong x, s::cl_ulong y) __NOEXC { return __u_long_mul_hi(x, y); } MAKE_1V_2V(u_mul_hi, s::cl_uchar, s::cl_uchar, s::cl_uchar) @@ -494,19 +527,20 @@ MAKE_1V_2V(u_mul_hi, s::cl_uint, s::cl_uint, s::cl_uint) MAKE_1V_2V(u_mul_hi, s::cl_ulong, s::cl_ulong, s::cl_ulong) // s_mad_hi -s::cl_char s_mad_hi(s::cl_char x, s::cl_char minval, - s::cl_char maxval) __NOEXC { +__SYCL_EXPORT s::cl_char s_mad_hi(s::cl_char x, s::cl_char minval, + s::cl_char maxval) __NOEXC { return __mad_hi(x, minval, maxval); } -s::cl_short s_mad_hi(s::cl_short x, s::cl_short minval, - s::cl_short maxval) __NOEXC { +__SYCL_EXPORT s::cl_short s_mad_hi(s::cl_short x, s::cl_short minval, + s::cl_short maxval) __NOEXC { return __mad_hi(x, minval, maxval); } -s::cl_int s_mad_hi(s::cl_int x, s::cl_int minval, s::cl_int maxval) __NOEXC { +__SYCL_EXPORT s::cl_int s_mad_hi(s::cl_int x, s::cl_int minval, + s::cl_int maxval) __NOEXC { return __mad_hi(x, minval, maxval); } -s::cl_long s_mad_hi(s::cl_long x, s::cl_long minval, - s::cl_long maxval) __NOEXC { +__SYCL_EXPORT s::cl_long s_mad_hi(s::cl_long x, s::cl_long minval, + s::cl_long maxval) __NOEXC { return __s_long_mad_hi(x, minval, maxval); } MAKE_1V_2V_3V(s_mad_hi, s::cl_char, s::cl_char, s::cl_char, s::cl_char) @@ -515,20 +549,20 @@ MAKE_1V_2V_3V(s_mad_hi, s::cl_int, s::cl_int, s::cl_int, s::cl_int) MAKE_1V_2V_3V(s_mad_hi, s::cl_long, s::cl_long, s::cl_long, s::cl_long) // u_mad_hi -s::cl_uchar u_mad_hi(s::cl_uchar x, s::cl_uchar minval, - s::cl_uchar maxval) __NOEXC { +__SYCL_EXPORT s::cl_uchar u_mad_hi(s::cl_uchar x, s::cl_uchar minval, + s::cl_uchar maxval) __NOEXC { return __mad_hi(x, minval, maxval); } -s::cl_ushort u_mad_hi(s::cl_ushort x, s::cl_ushort minval, - s::cl_ushort maxval) __NOEXC { +__SYCL_EXPORT s::cl_ushort u_mad_hi(s::cl_ushort x, s::cl_ushort minval, + s::cl_ushort maxval) __NOEXC { return __mad_hi(x, minval, maxval); } -s::cl_uint u_mad_hi(s::cl_uint x, s::cl_uint minval, - s::cl_uint maxval) __NOEXC { +__SYCL_EXPORT s::cl_uint u_mad_hi(s::cl_uint x, s::cl_uint minval, + s::cl_uint maxval) __NOEXC { return __mad_hi(x, minval, maxval); } -s::cl_ulong u_mad_hi(s::cl_ulong x, s::cl_ulong minval, - s::cl_ulong maxval) __NOEXC { +__SYCL_EXPORT s::cl_ulong u_mad_hi(s::cl_ulong x, s::cl_ulong minval, + s::cl_ulong maxval) __NOEXC { return __u_long_mad_hi(x, minval, maxval); } MAKE_1V_2V_3V(u_mad_hi, s::cl_uchar, s::cl_uchar, s::cl_uchar, s::cl_uchar) @@ -537,16 +571,20 @@ MAKE_1V_2V_3V(u_mad_hi, s::cl_uint, s::cl_uint, s::cl_uint, s::cl_uint) MAKE_1V_2V_3V(u_mad_hi, s::cl_ulong, s::cl_ulong, s::cl_ulong, s::cl_ulong) // s_mad_sat -s::cl_char s_mad_sat(s::cl_char a, s::cl_char b, s::cl_char c) __NOEXC { +__SYCL_EXPORT s::cl_char s_mad_sat(s::cl_char a, s::cl_char b, + s::cl_char c) __NOEXC { return __s_mad_sat(a, b, c); } -s::cl_short s_mad_sat(s::cl_short a, s::cl_short b, s::cl_short c) __NOEXC { +__SYCL_EXPORT s::cl_short s_mad_sat(s::cl_short a, s::cl_short b, + s::cl_short c) __NOEXC { return __s_mad_sat(a, b, c); } -s::cl_int s_mad_sat(s::cl_int a, s::cl_int b, s::cl_int c) __NOEXC { +__SYCL_EXPORT s::cl_int s_mad_sat(s::cl_int a, s::cl_int b, + s::cl_int c) __NOEXC { return __s_mad_sat(a, b, c); } -s::cl_long s_mad_sat(s::cl_long a, s::cl_long b, s::cl_long c) __NOEXC { +__SYCL_EXPORT s::cl_long s_mad_sat(s::cl_long a, s::cl_long b, + s::cl_long c) __NOEXC { return __s_long_mad_sat(a, b, c); } MAKE_1V_2V_3V(s_mad_sat, s::cl_char, s::cl_char, s::cl_char, s::cl_char) @@ -555,16 +593,20 @@ MAKE_1V_2V_3V(s_mad_sat, s::cl_int, s::cl_int, s::cl_int, s::cl_int) MAKE_1V_2V_3V(s_mad_sat, s::cl_long, s::cl_long, s::cl_long, s::cl_long) // u_mad_sat -s::cl_uchar u_mad_sat(s::cl_uchar a, s::cl_uchar b, s::cl_uchar c) __NOEXC { +__SYCL_EXPORT s::cl_uchar u_mad_sat(s::cl_uchar a, s::cl_uchar b, + s::cl_uchar c) __NOEXC { return __u_mad_sat(a, b, c); } -s::cl_ushort u_mad_sat(s::cl_ushort a, s::cl_ushort b, s::cl_ushort c) __NOEXC { +__SYCL_EXPORT s::cl_ushort u_mad_sat(s::cl_ushort a, s::cl_ushort b, + s::cl_ushort c) __NOEXC { return __u_mad_sat(a, b, c); } -s::cl_uint u_mad_sat(s::cl_uint a, s::cl_uint b, s::cl_uint c) __NOEXC { +__SYCL_EXPORT s::cl_uint u_mad_sat(s::cl_uint a, s::cl_uint b, + s::cl_uint c) __NOEXC { return __u_mad_sat(a, b, c); } -s::cl_ulong u_mad_sat(s::cl_ulong a, s::cl_ulong b, s::cl_ulong c) __NOEXC { +__SYCL_EXPORT s::cl_ulong u_mad_sat(s::cl_ulong a, s::cl_ulong b, + s::cl_ulong c) __NOEXC { return __u_long_mad_sat(a, b, c); } MAKE_1V_2V_3V(u_mad_sat, s::cl_uchar, s::cl_uchar, s::cl_uchar, s::cl_uchar) @@ -573,12 +615,18 @@ MAKE_1V_2V_3V(u_mad_sat, s::cl_uint, s::cl_uint, s::cl_uint, s::cl_uint) MAKE_1V_2V_3V(u_mad_sat, s::cl_ulong, s::cl_ulong, s::cl_ulong, s::cl_ulong) // s_max -s::cl_char s_max(s::cl_char x, s::cl_char y) __NOEXC { return std::max(x, y); } -s::cl_short s_max(s::cl_short x, s::cl_short y) __NOEXC { +__SYCL_EXPORT s::cl_char s_max(s::cl_char x, s::cl_char y) __NOEXC { + return std::max(x, y); +} +__SYCL_EXPORT s::cl_short s_max(s::cl_short x, s::cl_short y) __NOEXC { + return std::max(x, y); +} +__SYCL_EXPORT s::cl_int s_max(s::cl_int x, s::cl_int y) __NOEXC { + return std::max(x, y); +} +__SYCL_EXPORT s::cl_long s_max(s::cl_long x, s::cl_long y) __NOEXC { return std::max(x, y); } -s::cl_int s_max(s::cl_int x, s::cl_int y) __NOEXC { return std::max(x, y); } -s::cl_long s_max(s::cl_long x, s::cl_long y) __NOEXC { return std::max(x, y); } MAKE_1V_2V(s_max, s::cl_char, s::cl_char, s::cl_char) MAKE_1V_2V(s_max, s::cl_short, s::cl_short, s::cl_short) MAKE_1V_2V(s_max, s::cl_int, s::cl_int, s::cl_int) @@ -589,14 +637,16 @@ MAKE_1V_2S(s_max, s::cl_int, s::cl_int, s::cl_int) MAKE_1V_2S(s_max, s::cl_long, s::cl_long, s::cl_long) // u_max -s::cl_uchar u_max(s::cl_uchar x, s::cl_uchar y) __NOEXC { +__SYCL_EXPORT s::cl_uchar u_max(s::cl_uchar x, s::cl_uchar y) __NOEXC { return std::max(x, y); } -s::cl_ushort u_max(s::cl_ushort x, s::cl_ushort y) __NOEXC { +__SYCL_EXPORT s::cl_ushort u_max(s::cl_ushort x, s::cl_ushort y) __NOEXC { return std::max(x, y); } -s::cl_uint u_max(s::cl_uint x, s::cl_uint y) __NOEXC { return std::max(x, y); } -s::cl_ulong u_max(s::cl_ulong x, s::cl_ulong y) __NOEXC { +__SYCL_EXPORT s::cl_uint u_max(s::cl_uint x, s::cl_uint y) __NOEXC { + return std::max(x, y); +} +__SYCL_EXPORT s::cl_ulong u_max(s::cl_ulong x, s::cl_ulong y) __NOEXC { return std::max(x, y); } MAKE_1V_2V(u_max, s::cl_uchar, s::cl_uchar, s::cl_uchar) @@ -609,12 +659,18 @@ MAKE_1V_2S(u_max, s::cl_uint, s::cl_uint, s::cl_uint) MAKE_1V_2S(u_max, s::cl_ulong, s::cl_ulong, s::cl_ulong) // s_min -s::cl_char s_min(s::cl_char x, s::cl_char y) __NOEXC { return std::min(x, y); } -s::cl_short s_min(s::cl_short x, s::cl_short y) __NOEXC { +__SYCL_EXPORT s::cl_char s_min(s::cl_char x, s::cl_char y) __NOEXC { + return std::min(x, y); +} +__SYCL_EXPORT s::cl_short s_min(s::cl_short x, s::cl_short y) __NOEXC { + return std::min(x, y); +} +__SYCL_EXPORT s::cl_int s_min(s::cl_int x, s::cl_int y) __NOEXC { + return std::min(x, y); +} +__SYCL_EXPORT s::cl_long s_min(s::cl_long x, s::cl_long y) __NOEXC { return std::min(x, y); } -s::cl_int s_min(s::cl_int x, s::cl_int y) __NOEXC { return std::min(x, y); } -s::cl_long s_min(s::cl_long x, s::cl_long y) __NOEXC { return std::min(x, y); } MAKE_1V_2V(s_min, s::cl_char, s::cl_char, s::cl_char) MAKE_1V_2V(s_min, s::cl_short, s::cl_short, s::cl_short) MAKE_1V_2V(s_min, s::cl_int, s::cl_int, s::cl_int) @@ -625,14 +681,16 @@ MAKE_1V_2S(s_min, s::cl_int, s::cl_int, s::cl_int) MAKE_1V_2S(s_min, s::cl_long, s::cl_long, s::cl_long) // u_min -s::cl_uchar u_min(s::cl_uchar x, s::cl_uchar y) __NOEXC { +__SYCL_EXPORT s::cl_uchar u_min(s::cl_uchar x, s::cl_uchar y) __NOEXC { + return std::min(x, y); +} +__SYCL_EXPORT s::cl_ushort u_min(s::cl_ushort x, s::cl_ushort y) __NOEXC { return std::min(x, y); } -s::cl_ushort u_min(s::cl_ushort x, s::cl_ushort y) __NOEXC { +__SYCL_EXPORT s::cl_uint u_min(s::cl_uint x, s::cl_uint y) __NOEXC { return std::min(x, y); } -s::cl_uint u_min(s::cl_uint x, s::cl_uint y) __NOEXC { return std::min(x, y); } -s::cl_ulong u_min(s::cl_ulong x, s::cl_ulong y) __NOEXC { +__SYCL_EXPORT s::cl_ulong u_min(s::cl_ulong x, s::cl_ulong y) __NOEXC { return std::min(x, y); } MAKE_1V_2V(u_min, s::cl_uchar, s::cl_uchar, s::cl_uchar) @@ -645,22 +703,30 @@ MAKE_1V_2S(u_min, s::cl_uint, s::cl_uint, s::cl_uint) MAKE_1V_2S(u_min, s::cl_ulong, s::cl_ulong, s::cl_ulong) // rotate -s::cl_uchar rotate(s::cl_uchar x, s::cl_uchar y) __NOEXC { +__SYCL_EXPORT s::cl_uchar rotate(s::cl_uchar x, s::cl_uchar y) __NOEXC { + return __rotate(x, y); +} +__SYCL_EXPORT s::cl_ushort rotate(s::cl_ushort x, s::cl_ushort y) __NOEXC { + return __rotate(x, y); +} +__SYCL_EXPORT s::cl_uint rotate(s::cl_uint x, s::cl_uint y) __NOEXC { return __rotate(x, y); } -s::cl_ushort rotate(s::cl_ushort x, s::cl_ushort y) __NOEXC { +__SYCL_EXPORT s::cl_ulong rotate(s::cl_ulong x, s::cl_ulong y) __NOEXC { return __rotate(x, y); } -s::cl_uint rotate(s::cl_uint x, s::cl_uint y) __NOEXC { return __rotate(x, y); } -s::cl_ulong rotate(s::cl_ulong x, s::cl_ulong y) __NOEXC { +__SYCL_EXPORT s::cl_char rotate(s::cl_char x, s::cl_char y) __NOEXC { return __rotate(x, y); } -s::cl_char rotate(s::cl_char x, s::cl_char y) __NOEXC { return __rotate(x, y); } -s::cl_short rotate(s::cl_short x, s::cl_short y) __NOEXC { +__SYCL_EXPORT s::cl_short rotate(s::cl_short x, s::cl_short y) __NOEXC { + return __rotate(x, y); +} +__SYCL_EXPORT s::cl_int rotate(s::cl_int x, s::cl_int y) __NOEXC { + return __rotate(x, y); +} +__SYCL_EXPORT s::cl_long rotate(s::cl_long x, s::cl_long y) __NOEXC { return __rotate(x, y); } -s::cl_int rotate(s::cl_int x, s::cl_int y) __NOEXC { return __rotate(x, y); } -s::cl_long rotate(s::cl_long x, s::cl_long y) __NOEXC { return __rotate(x, y); } MAKE_1V_2V(rotate, s::cl_uchar, s::cl_uchar, s::cl_uchar) MAKE_1V_2V(rotate, s::cl_ushort, s::cl_ushort, s::cl_ushort) MAKE_1V_2V(rotate, s::cl_uint, s::cl_uint, s::cl_uint) @@ -671,16 +737,16 @@ MAKE_1V_2V(rotate, s::cl_int, s::cl_int, s::cl_int) MAKE_1V_2V(rotate, s::cl_long, s::cl_long, s::cl_long) // u_sub_sat -s::cl_uchar u_sub_sat(s::cl_uchar x, s::cl_uchar y) __NOEXC { +__SYCL_EXPORT s::cl_uchar u_sub_sat(s::cl_uchar x, s::cl_uchar y) __NOEXC { return __u_sub_sat(x, y); } -s::cl_ushort u_sub_sat(s::cl_ushort x, s::cl_ushort y) __NOEXC { +__SYCL_EXPORT s::cl_ushort u_sub_sat(s::cl_ushort x, s::cl_ushort y) __NOEXC { return __u_sub_sat(x, y); } -s::cl_uint u_sub_sat(s::cl_uint x, s::cl_uint y) __NOEXC { +__SYCL_EXPORT s::cl_uint u_sub_sat(s::cl_uint x, s::cl_uint y) __NOEXC { return __u_sub_sat(x, y); } -s::cl_ulong u_sub_sat(s::cl_ulong x, s::cl_ulong y) __NOEXC { +__SYCL_EXPORT s::cl_ulong u_sub_sat(s::cl_ulong x, s::cl_ulong y) __NOEXC { return __u_sub_sat(x, y); } MAKE_1V_2V(u_sub_sat, s::cl_uchar, s::cl_uchar, s::cl_uchar) @@ -689,16 +755,16 @@ MAKE_1V_2V(u_sub_sat, s::cl_uint, s::cl_uint, s::cl_uint) MAKE_1V_2V(u_sub_sat, s::cl_ulong, s::cl_ulong, s::cl_ulong) // s_sub_sat -s::cl_char s_sub_sat(s::cl_char x, s::cl_char y) __NOEXC { +__SYCL_EXPORT s::cl_char s_sub_sat(s::cl_char x, s::cl_char y) __NOEXC { return __s_sub_sat(x, y); } -s::cl_short s_sub_sat(s::cl_short x, s::cl_short y) __NOEXC { +__SYCL_EXPORT s::cl_short s_sub_sat(s::cl_short x, s::cl_short y) __NOEXC { return __s_sub_sat(x, y); } -s::cl_int s_sub_sat(s::cl_int x, s::cl_int y) __NOEXC { +__SYCL_EXPORT s::cl_int s_sub_sat(s::cl_int x, s::cl_int y) __NOEXC { return __s_sub_sat(x, y); } -s::cl_long s_sub_sat(s::cl_long x, s::cl_long y) __NOEXC { +__SYCL_EXPORT s::cl_long s_sub_sat(s::cl_long x, s::cl_long y) __NOEXC { return __s_sub_sat(x, y); } MAKE_1V_2V(s_sub_sat, s::cl_char, s::cl_char, s::cl_char) @@ -707,26 +773,26 @@ MAKE_1V_2V(s_sub_sat, s::cl_int, s::cl_int, s::cl_int) MAKE_1V_2V(s_sub_sat, s::cl_long, s::cl_long, s::cl_long) // u_upsample -s::cl_ushort u_upsample(s::cl_uchar x, s::cl_uchar y) __NOEXC { +__SYCL_EXPORT s::cl_ushort u_upsample(s::cl_uchar x, s::cl_uchar y) __NOEXC { return __upsample(x, y); } -s::cl_uint u_upsample(s::cl_ushort x, s::cl_ushort y) __NOEXC { +__SYCL_EXPORT s::cl_uint u_upsample(s::cl_ushort x, s::cl_ushort y) __NOEXC { return __upsample(x, y); } -s::cl_ulong u_upsample(s::cl_uint x, s::cl_uint y) __NOEXC { +__SYCL_EXPORT s::cl_ulong u_upsample(s::cl_uint x, s::cl_uint y) __NOEXC { return __upsample(x, y); } MAKE_1V_2V(u_upsample, s::cl_ushort, s::cl_uchar, s::cl_uchar) MAKE_1V_2V(u_upsample, s::cl_uint, s::cl_ushort, s::cl_ushort) MAKE_1V_2V(u_upsample, s::cl_ulong, s::cl_uint, s::cl_uint) -s::cl_short s_upsample(s::cl_char x, s::cl_uchar y) __NOEXC { +__SYCL_EXPORT s::cl_short s_upsample(s::cl_char x, s::cl_uchar y) __NOEXC { return __upsample(x, y); } -s::cl_int s_upsample(s::cl_short x, s::cl_ushort y) __NOEXC { +__SYCL_EXPORT s::cl_int s_upsample(s::cl_short x, s::cl_ushort y) __NOEXC { return __upsample(x, y); } -s::cl_long s_upsample(s::cl_int x, s::cl_uint y) __NOEXC { +__SYCL_EXPORT s::cl_long s_upsample(s::cl_int x, s::cl_uint y) __NOEXC { return __upsample(x, y); } MAKE_1V_2V(s_upsample, s::cl_short, s::cl_char, s::cl_uchar) @@ -734,42 +800,61 @@ MAKE_1V_2V(s_upsample, s::cl_int, s::cl_short, s::cl_ushort) MAKE_1V_2V(s_upsample, s::cl_long, s::cl_int, s::cl_uint) // popcount -s::cl_uchar popcount(s::cl_uchar x) __NOEXC { return __popcount(x); } -s::cl_ushort popcount(s::cl_ushort x) __NOEXC { return __popcount(x); } -s::cl_uint popcount(s::cl_uint x) __NOEXC { return __popcount(x); } -s::cl_ulong popcount(s::cl_ulong x) __NOEXC { return __popcount(x); } +__SYCL_EXPORT s::cl_uchar popcount(s::cl_uchar x) __NOEXC { + return __popcount(x); +} +__SYCL_EXPORT s::cl_ushort popcount(s::cl_ushort x) __NOEXC { + return __popcount(x); +} +__SYCL_EXPORT s::cl_uint popcount(s::cl_uint x) __NOEXC { + return __popcount(x); +} +__SYCL_EXPORT s::cl_ulong popcount(s::cl_ulong x) __NOEXC { + return __popcount(x); +} MAKE_1V(popcount, s::cl_uchar, s::cl_uchar) MAKE_1V(popcount, s::cl_ushort, s::cl_ushort) MAKE_1V(popcount, s::cl_uint, s::cl_uint) MAKE_1V(popcount, s::cl_ulong, s::cl_ulong) -s::cl_char popcount(s::cl_char x) __NOEXC { return __popcount(x); } -s::cl_short popcount(s::cl_short x) __NOEXC { return __popcount(x); } -s::cl_int popcount(s::cl_int x) __NOEXC { return __popcount(x); } -s::cl_long popcount(s::cl_long x) __NOEXC { return __popcount(x); } +__SYCL_EXPORT s::cl_char popcount(s::cl_char x) __NOEXC { + return __popcount(x); +} +__SYCL_EXPORT s::cl_short popcount(s::cl_short x) __NOEXC { + return __popcount(x); +} +__SYCL_EXPORT s::cl_int popcount(s::cl_int x) __NOEXC { return __popcount(x); } +__SYCL_EXPORT s::cl_long popcount(s::cl_long x) __NOEXC { + return __popcount(x); +} MAKE_1V(popcount, s::cl_char, s::cl_char) MAKE_1V(popcount, s::cl_short, s::cl_short) MAKE_1V(popcount, s::cl_int, s::cl_int) MAKE_1V(popcount, s::cl_long, s::cl_long) // u_mad24 -s::cl_uint u_mad24(s::cl_uint x, s::cl_uint y, s::cl_uint z) __NOEXC { +__SYCL_EXPORT s::cl_uint u_mad24(s::cl_uint x, s::cl_uint y, + s::cl_uint z) __NOEXC { return __mad24(x, y, z); } MAKE_1V_2V_3V(u_mad24, s::cl_uint, s::cl_uint, s::cl_uint, s::cl_uint) // s_mad24 -s::cl_int s_mad24(s::cl_int x, s::cl_int y, s::cl_int z) __NOEXC { +__SYCL_EXPORT s::cl_int s_mad24(s::cl_int x, s::cl_int y, s::cl_int z) __NOEXC { return __mad24(x, y, z); } MAKE_1V_2V_3V(s_mad24, s::cl_int, s::cl_int, s::cl_int, s::cl_int) // u_mul24 -s::cl_uint u_mul24(s::cl_uint x, s::cl_uint y) __NOEXC { return __mul24(x, y); } +__SYCL_EXPORT s::cl_uint u_mul24(s::cl_uint x, s::cl_uint y) __NOEXC { + return __mul24(x, y); +} MAKE_1V_2V(u_mul24, s::cl_uint, s::cl_uint, s::cl_uint) // s_mul24 -s::cl_int s_mul24(s::cl_int x, s::cl_int y) __NOEXC { return __mul24(x, y); } +__SYCL_EXPORT s::cl_int s_mul24(s::cl_int x, s::cl_int y) __NOEXC { + return __mul24(x, y); +} MAKE_1V_2V(s_mul24, s::cl_int, s::cl_int, s::cl_int) } // namespace __host_std diff --git a/sycl/source/detail/builtins_math.cpp b/sycl/source/detail/builtins_math.cpp index a71c13034b5cb..b84a08e842aec 100644 --- a/sycl/source/detail/builtins_math.cpp +++ b/sycl/source/detail/builtins_math.cpp @@ -16,6 +16,7 @@ #define _USE_MATH_DEFINES #include "builtins_helper.hpp" +#include #include @@ -91,69 +92,77 @@ template inline T __tanpi(T x) { return std::tan(M_PI * x); } // ----------------- 4.13.3 Math functions. Host implementations --------------- // acos -s::cl_float acos(s::cl_float x) __NOEXC { return std::acos(x); } -s::cl_double acos(s::cl_double x) __NOEXC { return std::acos(x); } -s::cl_half acos(s::cl_half x) __NOEXC { return std::acos(x); } +__SYCL_EXPORT s::cl_float acos(s::cl_float x) __NOEXC { return std::acos(x); } +__SYCL_EXPORT s::cl_double acos(s::cl_double x) __NOEXC { return std::acos(x); } +__SYCL_EXPORT s::cl_half acos(s::cl_half x) __NOEXC { return std::acos(x); } MAKE_1V(acos, s::cl_float, s::cl_float) MAKE_1V(acos, s::cl_double, s::cl_double) MAKE_1V(acos, s::cl_half, s::cl_half) // acosh -s::cl_float acosh(s::cl_float x) __NOEXC { return std::acosh(x); } -s::cl_double acosh(s::cl_double x) __NOEXC { return std::acosh(x); } -s::cl_half acosh(s::cl_half x) __NOEXC { return std::acosh(x); } +__SYCL_EXPORT s::cl_float acosh(s::cl_float x) __NOEXC { return std::acosh(x); } +__SYCL_EXPORT s::cl_double acosh(s::cl_double x) __NOEXC { + return std::acosh(x); +} +__SYCL_EXPORT s::cl_half acosh(s::cl_half x) __NOEXC { return std::acosh(x); } MAKE_1V(acosh, s::cl_float, s::cl_float) MAKE_1V(acosh, s::cl_double, s::cl_double) MAKE_1V(acosh, s::cl_half, s::cl_half) // acospi -s::cl_float acospi(s::cl_float x) __NOEXC { return __acospi(x); } -s::cl_double acospi(s::cl_double x) __NOEXC { return __acospi(x); } -s::cl_half acospi(s::cl_half x) __NOEXC { return __acospi(x); } +__SYCL_EXPORT s::cl_float acospi(s::cl_float x) __NOEXC { return __acospi(x); } +__SYCL_EXPORT s::cl_double acospi(s::cl_double x) __NOEXC { + return __acospi(x); +} +__SYCL_EXPORT s::cl_half acospi(s::cl_half x) __NOEXC { return __acospi(x); } MAKE_1V(acospi, s::cl_float, s::cl_float) MAKE_1V(acospi, s::cl_double, s::cl_double) MAKE_1V(acospi, s::cl_half, s::cl_half) // asin -s::cl_float asin(s::cl_float x) __NOEXC { return std::asin(x); } -s::cl_double asin(s::cl_double x) __NOEXC { return std::asin(x); } -s::cl_half asin(s::cl_half x) __NOEXC { return std::asin(x); } +__SYCL_EXPORT s::cl_float asin(s::cl_float x) __NOEXC { return std::asin(x); } +__SYCL_EXPORT s::cl_double asin(s::cl_double x) __NOEXC { return std::asin(x); } +__SYCL_EXPORT s::cl_half asin(s::cl_half x) __NOEXC { return std::asin(x); } MAKE_1V(asin, s::cl_float, s::cl_float) MAKE_1V(asin, s::cl_double, s::cl_double) MAKE_1V(asin, s::cl_half, s::cl_half) // asinh -s::cl_float asinh(s::cl_float x) __NOEXC { return std::asinh(x); } -s::cl_double asinh(s::cl_double x) __NOEXC { return std::asinh(x); } -s::cl_half asinh(s::cl_half x) __NOEXC { return std::asinh(x); } +__SYCL_EXPORT s::cl_float asinh(s::cl_float x) __NOEXC { return std::asinh(x); } +__SYCL_EXPORT s::cl_double asinh(s::cl_double x) __NOEXC { + return std::asinh(x); +} +__SYCL_EXPORT s::cl_half asinh(s::cl_half x) __NOEXC { return std::asinh(x); } MAKE_1V(asinh, s::cl_float, s::cl_float) MAKE_1V(asinh, s::cl_double, s::cl_double) MAKE_1V(asinh, s::cl_half, s::cl_half) // asinpi -s::cl_float asinpi(s::cl_float x) __NOEXC { return __asinpi(x); } -s::cl_double asinpi(s::cl_double x) __NOEXC { return __asinpi(x); } -s::cl_half asinpi(s::cl_half x) __NOEXC { return __asinpi(x); } +__SYCL_EXPORT s::cl_float asinpi(s::cl_float x) __NOEXC { return __asinpi(x); } +__SYCL_EXPORT s::cl_double asinpi(s::cl_double x) __NOEXC { + return __asinpi(x); +} +__SYCL_EXPORT s::cl_half asinpi(s::cl_half x) __NOEXC { return __asinpi(x); } MAKE_1V(asinpi, s::cl_float, s::cl_float) MAKE_1V(asinpi, s::cl_double, s::cl_double) MAKE_1V(asinpi, s::cl_half, s::cl_half) // atan -s::cl_float atan(s::cl_float x) __NOEXC { return std::atan(x); } -s::cl_double atan(s::cl_double x) __NOEXC { return std::atan(x); } -s::cl_half atan(s::cl_half x) __NOEXC { return std::atan(x); } +__SYCL_EXPORT s::cl_float atan(s::cl_float x) __NOEXC { return std::atan(x); } +__SYCL_EXPORT s::cl_double atan(s::cl_double x) __NOEXC { return std::atan(x); } +__SYCL_EXPORT s::cl_half atan(s::cl_half x) __NOEXC { return std::atan(x); } MAKE_1V(atan, s::cl_float, s::cl_float) MAKE_1V(atan, s::cl_double, s::cl_double) MAKE_1V(atan, s::cl_half, s::cl_half) // atan2 -s::cl_float atan2(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float atan2(s::cl_float x, s::cl_float y) __NOEXC { return std::atan2(x, y); } -s::cl_double atan2(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_double atan2(s::cl_double x, s::cl_double y) __NOEXC { return std::atan2(x, y); } -s::cl_half atan2(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_half atan2(s::cl_half x, s::cl_half y) __NOEXC { return std::atan2(x, y); } MAKE_1V_2V(atan2, s::cl_float, s::cl_float, s::cl_float) @@ -161,29 +170,33 @@ MAKE_1V_2V(atan2, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(atan2, s::cl_half, s::cl_half, s::cl_half) // atanh -s::cl_float atanh(s::cl_float x) __NOEXC { return std::atanh(x); } -s::cl_double atanh(s::cl_double x) __NOEXC { return std::atanh(x); } -s::cl_half atanh(s::cl_half x) __NOEXC { return std::atanh(x); } +__SYCL_EXPORT s::cl_float atanh(s::cl_float x) __NOEXC { return std::atanh(x); } +__SYCL_EXPORT s::cl_double atanh(s::cl_double x) __NOEXC { + return std::atanh(x); +} +__SYCL_EXPORT s::cl_half atanh(s::cl_half x) __NOEXC { return std::atanh(x); } MAKE_1V(atanh, s::cl_float, s::cl_float) MAKE_1V(atanh, s::cl_double, s::cl_double) MAKE_1V(atanh, s::cl_half, s::cl_half) // atanpi -s::cl_float atanpi(s::cl_float x) __NOEXC { return __atanpi(x); } -s::cl_double atanpi(s::cl_double x) __NOEXC { return __atanpi(x); } -s::cl_half atanpi(s::cl_half x) __NOEXC { return __atanpi(x); } +__SYCL_EXPORT s::cl_float atanpi(s::cl_float x) __NOEXC { return __atanpi(x); } +__SYCL_EXPORT s::cl_double atanpi(s::cl_double x) __NOEXC { + return __atanpi(x); +} +__SYCL_EXPORT s::cl_half atanpi(s::cl_half x) __NOEXC { return __atanpi(x); } MAKE_1V(atanpi, s::cl_float, s::cl_float) MAKE_1V(atanpi, s::cl_double, s::cl_double) MAKE_1V(atanpi, s::cl_half, s::cl_half) // atan2pi -s::cl_float atan2pi(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float atan2pi(s::cl_float x, s::cl_float y) __NOEXC { return __atan2pi(x, y); } -s::cl_double atan2pi(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_double atan2pi(s::cl_double x, s::cl_double y) __NOEXC { return __atan2pi(x, y); } -s::cl_half atan2pi(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_half atan2pi(s::cl_half x, s::cl_half y) __NOEXC { return __atan2pi(x, y); } MAKE_1V_2V(atan2pi, s::cl_float, s::cl_float, s::cl_float) @@ -191,29 +204,29 @@ MAKE_1V_2V(atan2pi, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(atan2pi, s::cl_half, s::cl_half, s::cl_half) // cbrt -s::cl_float cbrt(s::cl_float x) __NOEXC { return std::cbrt(x); } -s::cl_double cbrt(s::cl_double x) __NOEXC { return std::cbrt(x); } -s::cl_half cbrt(s::cl_half x) __NOEXC { return std::cbrt(x); } +__SYCL_EXPORT s::cl_float cbrt(s::cl_float x) __NOEXC { return std::cbrt(x); } +__SYCL_EXPORT s::cl_double cbrt(s::cl_double x) __NOEXC { return std::cbrt(x); } +__SYCL_EXPORT s::cl_half cbrt(s::cl_half x) __NOEXC { return std::cbrt(x); } MAKE_1V(cbrt, s::cl_float, s::cl_float) MAKE_1V(cbrt, s::cl_double, s::cl_double) MAKE_1V(cbrt, s::cl_half, s::cl_half) // ceil -s::cl_float ceil(s::cl_float x) __NOEXC { return std::ceil(x); } -s::cl_double ceil(s::cl_double x) __NOEXC { return std::ceil(x); } -s::cl_half ceil(s::cl_half x) __NOEXC { return std::ceil(x); } +__SYCL_EXPORT s::cl_float ceil(s::cl_float x) __NOEXC { return std::ceil(x); } +__SYCL_EXPORT s::cl_double ceil(s::cl_double x) __NOEXC { return std::ceil(x); } +__SYCL_EXPORT s::cl_half ceil(s::cl_half x) __NOEXC { return std::ceil(x); } MAKE_1V(ceil, s::cl_float, s::cl_float) MAKE_1V(ceil, s::cl_double, s::cl_double) MAKE_1V(ceil, s::cl_half, s::cl_half) // copysign -s::cl_float copysign(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float copysign(s::cl_float x, s::cl_float y) __NOEXC { return std::copysign(x, y); } -s::cl_double copysign(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_double copysign(s::cl_double x, s::cl_double y) __NOEXC { return std::copysign(x, y); } -s::cl_half copysign(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_half copysign(s::cl_half x, s::cl_half y) __NOEXC { return std::copysign(x, y); } MAKE_1V_2V(copysign, s::cl_float, s::cl_float, s::cl_float) @@ -221,113 +234,125 @@ MAKE_1V_2V(copysign, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(copysign, s::cl_half, s::cl_half, s::cl_half) // cos -s::cl_float cos(s::cl_float x) __NOEXC { return std::cos(x); } -s::cl_double cos(s::cl_double x) __NOEXC { return std::cos(x); } -s::cl_half cos(s::cl_half x) __NOEXC { return std::cos(x); } +__SYCL_EXPORT s::cl_float cos(s::cl_float x) __NOEXC { return std::cos(x); } +__SYCL_EXPORT s::cl_double cos(s::cl_double x) __NOEXC { return std::cos(x); } +__SYCL_EXPORT s::cl_half cos(s::cl_half x) __NOEXC { return std::cos(x); } MAKE_1V(cos, s::cl_float, s::cl_float) MAKE_1V(cos, s::cl_double, s::cl_double) MAKE_1V(cos, s::cl_half, s::cl_half) // cosh -s::cl_float cosh(s::cl_float x) __NOEXC { return std::cosh(x); } -s::cl_double cosh(s::cl_double x) __NOEXC { return std::cosh(x); } -s::cl_half cosh(s::cl_half x) __NOEXC { return std::cosh(x); } +__SYCL_EXPORT s::cl_float cosh(s::cl_float x) __NOEXC { return std::cosh(x); } +__SYCL_EXPORT s::cl_double cosh(s::cl_double x) __NOEXC { return std::cosh(x); } +__SYCL_EXPORT s::cl_half cosh(s::cl_half x) __NOEXC { return std::cosh(x); } MAKE_1V(cosh, s::cl_float, s::cl_float) MAKE_1V(cosh, s::cl_double, s::cl_double) MAKE_1V(cosh, s::cl_half, s::cl_half) // cospi -s::cl_float cospi(s::cl_float x) __NOEXC { return __cospi(x); } -s::cl_double cospi(s::cl_double x) __NOEXC { return __cospi(x); } -s::cl_half cospi(s::cl_half x) __NOEXC { return __cospi(x); } +__SYCL_EXPORT s::cl_float cospi(s::cl_float x) __NOEXC { return __cospi(x); } +__SYCL_EXPORT s::cl_double cospi(s::cl_double x) __NOEXC { return __cospi(x); } +__SYCL_EXPORT s::cl_half cospi(s::cl_half x) __NOEXC { return __cospi(x); } MAKE_1V(cospi, s::cl_float, s::cl_float) MAKE_1V(cospi, s::cl_double, s::cl_double) MAKE_1V(cospi, s::cl_half, s::cl_half) // erfc -s::cl_float erfc(s::cl_float x) __NOEXC { return std::erfc(x); } -s::cl_double erfc(s::cl_double x) __NOEXC { return std::erfc(x); } -s::cl_half erfc(s::cl_half x) __NOEXC { return std::erfc(x); } +__SYCL_EXPORT s::cl_float erfc(s::cl_float x) __NOEXC { return std::erfc(x); } +__SYCL_EXPORT s::cl_double erfc(s::cl_double x) __NOEXC { return std::erfc(x); } +__SYCL_EXPORT s::cl_half erfc(s::cl_half x) __NOEXC { return std::erfc(x); } MAKE_1V(erfc, s::cl_float, s::cl_float) MAKE_1V(erfc, s::cl_double, s::cl_double) MAKE_1V(erfc, s::cl_half, s::cl_half) // erf -s::cl_float erf(s::cl_float x) __NOEXC { return std::erf(x); } -s::cl_double erf(s::cl_double x) __NOEXC { return std::erf(x); } -s::cl_half erf(s::cl_half x) __NOEXC { return std::erf(x); } +__SYCL_EXPORT s::cl_float erf(s::cl_float x) __NOEXC { return std::erf(x); } +__SYCL_EXPORT s::cl_double erf(s::cl_double x) __NOEXC { return std::erf(x); } +__SYCL_EXPORT s::cl_half erf(s::cl_half x) __NOEXC { return std::erf(x); } MAKE_1V(erf, s::cl_float, s::cl_float) MAKE_1V(erf, s::cl_double, s::cl_double) MAKE_1V(erf, s::cl_half, s::cl_half) // exp -s::cl_float exp(s::cl_float x) __NOEXC { return std::exp(x); } -s::cl_double exp(s::cl_double x) __NOEXC { return std::exp(x); } -s::cl_half exp(s::cl_half x) __NOEXC { return std::exp(x); } +__SYCL_EXPORT s::cl_float exp(s::cl_float x) __NOEXC { return std::exp(x); } +__SYCL_EXPORT s::cl_double exp(s::cl_double x) __NOEXC { return std::exp(x); } +__SYCL_EXPORT s::cl_half exp(s::cl_half x) __NOEXC { return std::exp(x); } MAKE_1V(exp, s::cl_float, s::cl_float) MAKE_1V(exp, s::cl_double, s::cl_double) MAKE_1V(exp, s::cl_half, s::cl_half) // exp2 -s::cl_float exp2(s::cl_float x) __NOEXC { return std::exp2(x); } -s::cl_double exp2(s::cl_double x) __NOEXC { return std::exp2(x); } -s::cl_half exp2(s::cl_half x) __NOEXC { return std::exp2(x); } +__SYCL_EXPORT s::cl_float exp2(s::cl_float x) __NOEXC { return std::exp2(x); } +__SYCL_EXPORT s::cl_double exp2(s::cl_double x) __NOEXC { return std::exp2(x); } +__SYCL_EXPORT s::cl_half exp2(s::cl_half x) __NOEXC { return std::exp2(x); } MAKE_1V(exp2, s::cl_float, s::cl_float) MAKE_1V(exp2, s::cl_double, s::cl_double) MAKE_1V(exp2, s::cl_half, s::cl_half) // exp10 -s::cl_float exp10(s::cl_float x) __NOEXC { return std::pow(10, x); } -s::cl_double exp10(s::cl_double x) __NOEXC { return std::pow(10, x); } -s::cl_half exp10(s::cl_half x) __NOEXC { return std::pow(10, x); } +__SYCL_EXPORT s::cl_float exp10(s::cl_float x) __NOEXC { + return std::pow(10, x); +} +__SYCL_EXPORT s::cl_double exp10(s::cl_double x) __NOEXC { + return std::pow(10, x); +} +__SYCL_EXPORT s::cl_half exp10(s::cl_half x) __NOEXC { return std::pow(10, x); } MAKE_1V(exp10, s::cl_float, s::cl_float) MAKE_1V(exp10, s::cl_double, s::cl_double) MAKE_1V(exp10, s::cl_half, s::cl_half) // expm1 -s::cl_float expm1(s::cl_float x) __NOEXC { return std::expm1(x); } -s::cl_double expm1(s::cl_double x) __NOEXC { return std::expm1(x); } -s::cl_half expm1(s::cl_half x) __NOEXC { return std::expm1(x); } +__SYCL_EXPORT s::cl_float expm1(s::cl_float x) __NOEXC { return std::expm1(x); } +__SYCL_EXPORT s::cl_double expm1(s::cl_double x) __NOEXC { + return std::expm1(x); +} +__SYCL_EXPORT s::cl_half expm1(s::cl_half x) __NOEXC { return std::expm1(x); } MAKE_1V(expm1, s::cl_float, s::cl_float) MAKE_1V(expm1, s::cl_double, s::cl_double) MAKE_1V(expm1, s::cl_half, s::cl_half) // fabs -s::cl_float fabs(s::cl_float x) __NOEXC { return std::fabs(x); } -s::cl_double fabs(s::cl_double x) __NOEXC { return std::fabs(x); } -s::cl_half fabs(s::cl_half x) __NOEXC { return std::fabs(x); } +__SYCL_EXPORT s::cl_float fabs(s::cl_float x) __NOEXC { return std::fabs(x); } +__SYCL_EXPORT s::cl_double fabs(s::cl_double x) __NOEXC { return std::fabs(x); } +__SYCL_EXPORT s::cl_half fabs(s::cl_half x) __NOEXC { return std::fabs(x); } MAKE_1V(fabs, s::cl_float, s::cl_float) MAKE_1V(fabs, s::cl_double, s::cl_double) MAKE_1V(fabs, s::cl_half, s::cl_half) // fdim -s::cl_float fdim(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float fdim(s::cl_float x, s::cl_float y) __NOEXC { return std::fdim(x, y); } -s::cl_double fdim(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_double fdim(s::cl_double x, s::cl_double y) __NOEXC { + return std::fdim(x, y); +} +__SYCL_EXPORT s::cl_half fdim(s::cl_half x, s::cl_half y) __NOEXC { return std::fdim(x, y); } -s::cl_half fdim(s::cl_half x, s::cl_half y) __NOEXC { return std::fdim(x, y); } MAKE_1V_2V(fdim, s::cl_float, s::cl_float, s::cl_float) MAKE_1V_2V(fdim, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(fdim, s::cl_half, s::cl_half, s::cl_half) // floor -s::cl_float floor(s::cl_float x) __NOEXC { return std::floor(x); } -s::cl_double floor(s::cl_double x) __NOEXC { return std::floor(x); } -s::cl_half floor(s::cl_half x) __NOEXC { return std::floor(x); } +__SYCL_EXPORT s::cl_float floor(s::cl_float x) __NOEXC { return std::floor(x); } +__SYCL_EXPORT s::cl_double floor(s::cl_double x) __NOEXC { + return std::floor(x); +} +__SYCL_EXPORT s::cl_half floor(s::cl_half x) __NOEXC { return std::floor(x); } MAKE_1V(floor, s::cl_float, s::cl_float) MAKE_1V(floor, s::cl_double, s::cl_double) MAKE_1V(floor, s::cl_half, s::cl_half) // fma -s::cl_float fma(s::cl_float a, s::cl_float b, s::cl_float c) __NOEXC { +__SYCL_EXPORT s::cl_float fma(s::cl_float a, s::cl_float b, + s::cl_float c) __NOEXC { return std::fma(a, b, c); } -s::cl_double fma(s::cl_double a, s::cl_double b, s::cl_double c) __NOEXC { +__SYCL_EXPORT s::cl_double fma(s::cl_double a, s::cl_double b, + s::cl_double c) __NOEXC { return std::fma(a, b, c); } -s::cl_half fma(s::cl_half a, s::cl_half b, s::cl_half c) __NOEXC { +__SYCL_EXPORT s::cl_half fma(s::cl_half a, s::cl_half b, s::cl_half c) __NOEXC { return std::fma(a, b, c); } MAKE_1V_2V_3V(fma, s::cl_float, s::cl_float, s::cl_float, s::cl_float) @@ -335,49 +360,55 @@ MAKE_1V_2V_3V(fma, s::cl_double, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V_3V(fma, s::cl_half, s::cl_half, s::cl_half, s::cl_half) // fmax -s::cl_float fmax(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float fmax(s::cl_float x, s::cl_float y) __NOEXC { return std::fmax(x, y); } -s::cl_double fmax(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_double fmax(s::cl_double x, s::cl_double y) __NOEXC { + return std::fmax(x, y); +} +__SYCL_EXPORT s::cl_half fmax(s::cl_half x, s::cl_half y) __NOEXC { return std::fmax(x, y); } -s::cl_half fmax(s::cl_half x, s::cl_half y) __NOEXC { return std::fmax(x, y); } MAKE_1V_2V(fmax, s::cl_float, s::cl_float, s::cl_float) MAKE_1V_2V(fmax, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(fmax, s::cl_half, s::cl_half, s::cl_half) // fmin -s::cl_float fmin(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float fmin(s::cl_float x, s::cl_float y) __NOEXC { return std::fmin(x, y); } -s::cl_double fmin(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_double fmin(s::cl_double x, s::cl_double y) __NOEXC { + return std::fmin(x, y); +} +__SYCL_EXPORT s::cl_half fmin(s::cl_half x, s::cl_half y) __NOEXC { return std::fmin(x, y); } -s::cl_half fmin(s::cl_half x, s::cl_half y) __NOEXC { return std::fmin(x, y); } MAKE_1V_2V(fmin, s::cl_float, s::cl_float, s::cl_float) MAKE_1V_2V(fmin, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(fmin, s::cl_half, s::cl_half, s::cl_half) // fmod -s::cl_float fmod(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float fmod(s::cl_float x, s::cl_float y) __NOEXC { + return std::fmod(x, y); +} +__SYCL_EXPORT s::cl_double fmod(s::cl_double x, s::cl_double y) __NOEXC { return std::fmod(x, y); } -s::cl_double fmod(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_half fmod(s::cl_half x, s::cl_half y) __NOEXC { return std::fmod(x, y); } -s::cl_half fmod(s::cl_half x, s::cl_half y) __NOEXC { return std::fmod(x, y); } MAKE_1V_2V(fmod, s::cl_float, s::cl_float, s::cl_float) MAKE_1V_2V(fmod, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(fmod, s::cl_half, s::cl_half, s::cl_half) // nextafter -s::cl_float nextafter(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float nextafter(s::cl_float x, s::cl_float y) __NOEXC { return std::nextafter(x, y); } -s::cl_double nextafter(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_double nextafter(s::cl_double x, s::cl_double y) __NOEXC { return std::nextafter(x, y); } -s::cl_half nextafter(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_half nextafter(s::cl_half x, s::cl_half y) __NOEXC { return std::nextafter(x, y); } MAKE_1V_2V(nextafter, s::cl_float, s::cl_float, s::cl_float) @@ -385,13 +416,13 @@ MAKE_1V_2V(nextafter, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(nextafter, s::cl_half, s::cl_half, s::cl_half) // fract -s::cl_float fract(s::cl_float x, s::cl_float *iptr) __NOEXC { +__SYCL_EXPORT s::cl_float fract(s::cl_float x, s::cl_float *iptr) __NOEXC { return __fract(x, iptr); } -s::cl_double fract(s::cl_double x, s::cl_double *iptr) __NOEXC { +__SYCL_EXPORT s::cl_double fract(s::cl_double x, s::cl_double *iptr) __NOEXC { return __fract(x, iptr); } -s::cl_half fract(s::cl_half x, s::cl_half *iptr) __NOEXC { +__SYCL_EXPORT s::cl_half fract(s::cl_half x, s::cl_half *iptr) __NOEXC { return __fract(x, iptr); } MAKE_1V_2P(fract, s::cl_float, s::cl_float, s::cl_float) @@ -399,13 +430,13 @@ MAKE_1V_2P(fract, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2P(fract, s::cl_half, s::cl_half, s::cl_half) // frexp -s::cl_float frexp(s::cl_float x, s::cl_int *exp) __NOEXC { +__SYCL_EXPORT s::cl_float frexp(s::cl_float x, s::cl_int *exp) __NOEXC { return std::frexp(x, exp); } -s::cl_double frexp(s::cl_double x, s::cl_int *exp) __NOEXC { +__SYCL_EXPORT s::cl_double frexp(s::cl_double x, s::cl_int *exp) __NOEXC { return std::frexp(x, exp); } -s::cl_half frexp(s::cl_half x, s::cl_int *exp) __NOEXC { +__SYCL_EXPORT s::cl_half frexp(s::cl_half x, s::cl_int *exp) __NOEXC { return std::frexp(x, exp); } MAKE_1V_2P(frexp, s::cl_float, s::cl_float, s::cl_int) @@ -413,13 +444,13 @@ MAKE_1V_2P(frexp, s::cl_double, s::cl_double, s::cl_int) MAKE_1V_2P(frexp, s::cl_half, s::cl_half, s::cl_int) // hypot -s::cl_float hypot(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float hypot(s::cl_float x, s::cl_float y) __NOEXC { return std::hypot(x, y); } -s::cl_double hypot(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_double hypot(s::cl_double x, s::cl_double y) __NOEXC { return std::hypot(x, y); } -s::cl_half hypot(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_half hypot(s::cl_half x, s::cl_half y) __NOEXC { return std::hypot(x, y); } MAKE_1V_2V(hypot, s::cl_float, s::cl_float, s::cl_float) @@ -427,41 +458,47 @@ MAKE_1V_2V(hypot, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(hypot, s::cl_half, s::cl_half, s::cl_half) // ilogb -s::cl_int ilogb(s::cl_float x) __NOEXC { return std::ilogb(x); } -s::cl_int ilogb(s::cl_double x) __NOEXC { return std::ilogb(x); } -s::cl_int ilogb(s::cl_half x) __NOEXC { return std::ilogb(x); } +__SYCL_EXPORT s::cl_int ilogb(s::cl_float x) __NOEXC { return std::ilogb(x); } +__SYCL_EXPORT s::cl_int ilogb(s::cl_double x) __NOEXC { return std::ilogb(x); } +__SYCL_EXPORT s::cl_int ilogb(s::cl_half x) __NOEXC { return std::ilogb(x); } MAKE_1V(ilogb, s::cl_int, s::cl_float) MAKE_1V(ilogb, s::cl_int, s::cl_double) MAKE_1V(ilogb, s::cl_int, s::cl_half) // ldexp -s::cl_float ldexp(s::cl_float x, s::cl_int k) __NOEXC { +__SYCL_EXPORT s::cl_float ldexp(s::cl_float x, s::cl_int k) __NOEXC { + return std::ldexp(x, k); +} +__SYCL_EXPORT s::cl_double ldexp(s::cl_double x, s::cl_int k) __NOEXC { return std::ldexp(x, k); } -s::cl_double ldexp(s::cl_double x, s::cl_int k) __NOEXC { +__SYCL_EXPORT s::cl_half ldexp(s::cl_half x, s::cl_int k) __NOEXC { return std::ldexp(x, k); } -s::cl_half ldexp(s::cl_half x, s::cl_int k) __NOEXC { return std::ldexp(x, k); } MAKE_1V_2V(ldexp, s::cl_float, s::cl_float, s::cl_int) MAKE_1V_2V(ldexp, s::cl_double, s::cl_double, s::cl_int) MAKE_1V_2V(ldexp, s::cl_half, s::cl_half, s::cl_int) // lgamma -s::cl_float lgamma(s::cl_float x) __NOEXC { return std::lgamma(x); } -s::cl_double lgamma(s::cl_double x) __NOEXC { return std::lgamma(x); } -s::cl_half lgamma(s::cl_half x) __NOEXC { return std::lgamma(x); } +__SYCL_EXPORT s::cl_float lgamma(s::cl_float x) __NOEXC { + return std::lgamma(x); +} +__SYCL_EXPORT s::cl_double lgamma(s::cl_double x) __NOEXC { + return std::lgamma(x); +} +__SYCL_EXPORT s::cl_half lgamma(s::cl_half x) __NOEXC { return std::lgamma(x); } MAKE_1V(lgamma, s::cl_float, s::cl_float) MAKE_1V(lgamma, s::cl_double, s::cl_double) MAKE_1V(lgamma, s::cl_half, s::cl_half) // lgamma_r -s::cl_float lgamma_r(s::cl_float x, s::cl_int *signp) __NOEXC { +__SYCL_EXPORT s::cl_float lgamma_r(s::cl_float x, s::cl_int *signp) __NOEXC { return __lgamma_r(x, signp); } -s::cl_double lgamma_r(s::cl_double x, s::cl_int *signp) __NOEXC { +__SYCL_EXPORT s::cl_double lgamma_r(s::cl_double x, s::cl_int *signp) __NOEXC { return __lgamma_r(x, signp); } -s::cl_half lgamma_r(s::cl_half x, s::cl_int *signp) __NOEXC { +__SYCL_EXPORT s::cl_half lgamma_r(s::cl_half x, s::cl_int *signp) __NOEXC { return __lgamma_r(x, signp); } MAKE_1V_2P(lgamma_r, s::cl_float, s::cl_float, s::cl_int) @@ -469,53 +506,59 @@ MAKE_1V_2P(lgamma_r, s::cl_double, s::cl_double, s::cl_int) MAKE_1V_2P(lgamma_r, s::cl_half, s::cl_half, s::cl_int) // log -s::cl_float log(s::cl_float x) __NOEXC { return std::log(x); } -s::cl_double log(s::cl_double x) __NOEXC { return std::log(x); } -s::cl_half log(s::cl_half x) __NOEXC { return std::log(x); } +__SYCL_EXPORT s::cl_float log(s::cl_float x) __NOEXC { return std::log(x); } +__SYCL_EXPORT s::cl_double log(s::cl_double x) __NOEXC { return std::log(x); } +__SYCL_EXPORT s::cl_half log(s::cl_half x) __NOEXC { return std::log(x); } MAKE_1V(log, s::cl_float, s::cl_float) MAKE_1V(log, s::cl_double, s::cl_double) MAKE_1V(log, s::cl_half, s::cl_half) // log2 -s::cl_float log2(s::cl_float x) __NOEXC { return std::log2(x); } -s::cl_double log2(s::cl_double x) __NOEXC { return std::log2(x); } -s::cl_half log2(s::cl_half x) __NOEXC { return std::log2(x); } +__SYCL_EXPORT s::cl_float log2(s::cl_float x) __NOEXC { return std::log2(x); } +__SYCL_EXPORT s::cl_double log2(s::cl_double x) __NOEXC { return std::log2(x); } +__SYCL_EXPORT s::cl_half log2(s::cl_half x) __NOEXC { return std::log2(x); } MAKE_1V(log2, s::cl_float, s::cl_float) MAKE_1V(log2, s::cl_double, s::cl_double) MAKE_1V(log2, s::cl_half, s::cl_half) // log10 -s::cl_float log10(s::cl_float x) __NOEXC { return std::log10(x); } -s::cl_double log10(s::cl_double x) __NOEXC { return std::log10(x); } -s::cl_half log10(s::cl_half x) __NOEXC { return std::log10(x); } +__SYCL_EXPORT s::cl_float log10(s::cl_float x) __NOEXC { return std::log10(x); } +__SYCL_EXPORT s::cl_double log10(s::cl_double x) __NOEXC { + return std::log10(x); +} +__SYCL_EXPORT s::cl_half log10(s::cl_half x) __NOEXC { return std::log10(x); } MAKE_1V(log10, s::cl_float, s::cl_float) MAKE_1V(log10, s::cl_double, s::cl_double) MAKE_1V(log10, s::cl_half, s::cl_half) // log1p -s::cl_float log1p(s::cl_float x) __NOEXC { return std::log1p(x); } -s::cl_double log1p(s::cl_double x) __NOEXC { return std::log1p(x); } -s::cl_half log1p(s::cl_half x) __NOEXC { return std::log1p(x); } +__SYCL_EXPORT s::cl_float log1p(s::cl_float x) __NOEXC { return std::log1p(x); } +__SYCL_EXPORT s::cl_double log1p(s::cl_double x) __NOEXC { + return std::log1p(x); +} +__SYCL_EXPORT s::cl_half log1p(s::cl_half x) __NOEXC { return std::log1p(x); } MAKE_1V(log1p, s::cl_float, s::cl_float) MAKE_1V(log1p, s::cl_double, s::cl_double) MAKE_1V(log1p, s::cl_half, s::cl_half) // logb -s::cl_float logb(s::cl_float x) __NOEXC { return std::logb(x); } -s::cl_double logb(s::cl_double x) __NOEXC { return std::logb(x); } -s::cl_half logb(s::cl_half x) __NOEXC { return std::logb(x); } +__SYCL_EXPORT s::cl_float logb(s::cl_float x) __NOEXC { return std::logb(x); } +__SYCL_EXPORT s::cl_double logb(s::cl_double x) __NOEXC { return std::logb(x); } +__SYCL_EXPORT s::cl_half logb(s::cl_half x) __NOEXC { return std::logb(x); } MAKE_1V(logb, s::cl_float, s::cl_float) MAKE_1V(logb, s::cl_double, s::cl_double) MAKE_1V(logb, s::cl_half, s::cl_half) // mad -s::cl_float mad(s::cl_float a, s::cl_float b, s::cl_float c) __NOEXC { +__SYCL_EXPORT s::cl_float mad(s::cl_float a, s::cl_float b, + s::cl_float c) __NOEXC { return __mad(a, b, c); } -s::cl_double mad(s::cl_double a, s::cl_double b, s::cl_double c) __NOEXC { +__SYCL_EXPORT s::cl_double mad(s::cl_double a, s::cl_double b, + s::cl_double c) __NOEXC { return __mad(a, b, c); } -s::cl_half mad(s::cl_half a, s::cl_half b, s::cl_half c) __NOEXC { +__SYCL_EXPORT s::cl_half mad(s::cl_half a, s::cl_half b, s::cl_half c) __NOEXC { return __mad(a, b, c); } MAKE_1V_2V_3V(mad, s::cl_float, s::cl_float, s::cl_float, s::cl_float) @@ -523,37 +566,41 @@ MAKE_1V_2V_3V(mad, s::cl_double, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V_3V(mad, s::cl_half, s::cl_half, s::cl_half, s::cl_half) // maxmag -s::cl_float maxmag(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float maxmag(s::cl_float x, s::cl_float y) __NOEXC { return __maxmag(x, y); } -s::cl_double maxmag(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_double maxmag(s::cl_double x, s::cl_double y) __NOEXC { + return __maxmag(x, y); +} +__SYCL_EXPORT s::cl_half maxmag(s::cl_half x, s::cl_half y) __NOEXC { return __maxmag(x, y); } -s::cl_half maxmag(s::cl_half x, s::cl_half y) __NOEXC { return __maxmag(x, y); } MAKE_1V_2V(maxmag, s::cl_float, s::cl_float, s::cl_float) MAKE_1V_2V(maxmag, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(maxmag, s::cl_half, s::cl_half, s::cl_half) // minmag -s::cl_float minmag(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float minmag(s::cl_float x, s::cl_float y) __NOEXC { return __minmag(x, y); } -s::cl_double minmag(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_double minmag(s::cl_double x, s::cl_double y) __NOEXC { + return __minmag(x, y); +} +__SYCL_EXPORT s::cl_half minmag(s::cl_half x, s::cl_half y) __NOEXC { return __minmag(x, y); } -s::cl_half minmag(s::cl_half x, s::cl_half y) __NOEXC { return __minmag(x, y); } MAKE_1V_2V(minmag, s::cl_float, s::cl_float, s::cl_float) MAKE_1V_2V(minmag, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(minmag, s::cl_half, s::cl_half, s::cl_half) // modf -s::cl_float modf(s::cl_float x, s::cl_float *iptr) __NOEXC { +__SYCL_EXPORT s::cl_float modf(s::cl_float x, s::cl_float *iptr) __NOEXC { return std::modf(x, iptr); } -s::cl_double modf(s::cl_double x, s::cl_double *iptr) __NOEXC { +__SYCL_EXPORT s::cl_double modf(s::cl_double x, s::cl_double *iptr) __NOEXC { return std::modf(x, iptr); } -s::cl_half modf(s::cl_half x, s::cl_half *iptr) __NOEXC { +__SYCL_EXPORT s::cl_half modf(s::cl_half x, s::cl_half *iptr) __NOEXC { return std::modf(x, reinterpret_cast(iptr)); } MAKE_1V_2P(modf, s::cl_float, s::cl_float, s::cl_float) @@ -561,13 +608,13 @@ MAKE_1V_2P(modf, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2P(modf, s::cl_half, s::cl_half, s::cl_half) // nan -s::cl_float nan(s::cl_uint nancode) __NOEXC { +__SYCL_EXPORT s::cl_float nan(s::cl_uint nancode) __NOEXC { return d::quiet_NaN(); } -s::cl_double nan(s::cl_ulong nancode) __NOEXC { +__SYCL_EXPORT s::cl_double nan(s::cl_ulong nancode) __NOEXC { return d::quiet_NaN(); } -s::cl_half nan(s::cl_ushort nancode) __NOEXC { +__SYCL_EXPORT s::cl_half nan(s::cl_ushort nancode) __NOEXC { return s::cl_half(d::quiet_NaN()); } MAKE_1V(nan, s::cl_float, s::cl_uint) @@ -575,43 +622,55 @@ MAKE_1V(nan, s::cl_double, s::cl_ulong) MAKE_1V(nan, s::cl_half, s::cl_ushort) // pow -s::cl_float pow(s::cl_float x, s::cl_float y) __NOEXC { return std::pow(x, y); } -s::cl_double pow(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_float pow(s::cl_float x, s::cl_float y) __NOEXC { + return std::pow(x, y); +} +__SYCL_EXPORT s::cl_double pow(s::cl_double x, s::cl_double y) __NOEXC { + return std::pow(x, y); +} +__SYCL_EXPORT s::cl_half pow(s::cl_half x, s::cl_half y) __NOEXC { return std::pow(x, y); } -s::cl_half pow(s::cl_half x, s::cl_half y) __NOEXC { return std::pow(x, y); } MAKE_1V_2V(pow, s::cl_float, s::cl_float, s::cl_float) MAKE_1V_2V(pow, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(pow, s::cl_half, s::cl_half, s::cl_half) // pown -s::cl_float pown(s::cl_float x, s::cl_int y) __NOEXC { return std::pow(x, y); } -s::cl_double pown(s::cl_double x, s::cl_int y) __NOEXC { +__SYCL_EXPORT s::cl_float pown(s::cl_float x, s::cl_int y) __NOEXC { + return std::pow(x, y); +} +__SYCL_EXPORT s::cl_double pown(s::cl_double x, s::cl_int y) __NOEXC { + return std::pow(x, y); +} +__SYCL_EXPORT s::cl_half pown(s::cl_half x, s::cl_int y) __NOEXC { return std::pow(x, y); } -s::cl_half pown(s::cl_half x, s::cl_int y) __NOEXC { return std::pow(x, y); } MAKE_1V_2V(pown, s::cl_float, s::cl_float, s::cl_int) MAKE_1V_2V(pown, s::cl_double, s::cl_double, s::cl_int) MAKE_1V_2V(pown, s::cl_half, s::cl_half, s::cl_int) // powr -s::cl_float powr(s::cl_float x, s::cl_float y) __NOEXC { return __powr(x, y); } -s::cl_double powr(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_float powr(s::cl_float x, s::cl_float y) __NOEXC { + return __powr(x, y); +} +__SYCL_EXPORT s::cl_double powr(s::cl_double x, s::cl_double y) __NOEXC { + return __powr(x, y); +} +__SYCL_EXPORT s::cl_half powr(s::cl_half x, s::cl_half y) __NOEXC { return __powr(x, y); } -s::cl_half powr(s::cl_half x, s::cl_half y) __NOEXC { return __powr(x, y); } MAKE_1V_2V(powr, s::cl_float, s::cl_float, s::cl_float) MAKE_1V_2V(powr, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(powr, s::cl_half, s::cl_half, s::cl_half) // remainder -s::cl_float remainder(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float remainder(s::cl_float x, s::cl_float y) __NOEXC { return std::remainder(x, y); } -s::cl_double remainder(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_double remainder(s::cl_double x, s::cl_double y) __NOEXC { return std::remainder(x, y); } -s::cl_half remainder(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_half remainder(s::cl_half x, s::cl_half y) __NOEXC { return std::remainder(x, y); } MAKE_1V_2V(remainder, s::cl_float, s::cl_float, s::cl_float) @@ -619,13 +678,16 @@ MAKE_1V_2V(remainder, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2V(remainder, s::cl_half, s::cl_half, s::cl_half) // remquo -s::cl_float remquo(s::cl_float x, s::cl_float y, s::cl_int *quo) __NOEXC { +__SYCL_EXPORT s::cl_float remquo(s::cl_float x, s::cl_float y, + s::cl_int *quo) __NOEXC { return std::remquo(x, y, quo); } -s::cl_double remquo(s::cl_double x, s::cl_double y, s::cl_int *quo) __NOEXC { +__SYCL_EXPORT s::cl_double remquo(s::cl_double x, s::cl_double y, + s::cl_int *quo) __NOEXC { return std::remquo(x, y, quo); } -s::cl_half remquo(s::cl_half x, s::cl_half y, s::cl_int *quo) __NOEXC { +__SYCL_EXPORT s::cl_half remquo(s::cl_half x, s::cl_half y, + s::cl_int *quo) __NOEXC { return std::remquo(x, y, quo); } MAKE_1V_2V_3P(remquo, s::cl_float, s::cl_float, s::cl_float, s::cl_int) @@ -633,55 +695,62 @@ MAKE_1V_2V_3P(remquo, s::cl_double, s::cl_double, s::cl_double, s::cl_int) MAKE_1V_2V_3P(remquo, s::cl_half, s::cl_half, s::cl_half, s::cl_int) // rint -s::cl_float rint(s::cl_float x) __NOEXC { return std::rint(x); } -s::cl_double rint(s::cl_double x) __NOEXC { return std::rint(x); } -s::cl_half rint(s::cl_half x) __NOEXC { return std::rint(x); } +__SYCL_EXPORT s::cl_float rint(s::cl_float x) __NOEXC { return std::rint(x); } +__SYCL_EXPORT s::cl_double rint(s::cl_double x) __NOEXC { return std::rint(x); } +__SYCL_EXPORT s::cl_half rint(s::cl_half x) __NOEXC { return std::rint(x); } MAKE_1V(rint, s::cl_float, s::cl_float) MAKE_1V(rint, s::cl_double, s::cl_double) MAKE_1V(rint, s::cl_half, s::cl_half) // rootn -s::cl_float rootn(s::cl_float x, s::cl_int y) __NOEXC { return __rootn(x, y); } -s::cl_double rootn(s::cl_double x, s::cl_int y) __NOEXC { +__SYCL_EXPORT s::cl_float rootn(s::cl_float x, s::cl_int y) __NOEXC { + return __rootn(x, y); +} +__SYCL_EXPORT s::cl_double rootn(s::cl_double x, s::cl_int y) __NOEXC { + return __rootn(x, y); +} +__SYCL_EXPORT s::cl_half rootn(s::cl_half x, s::cl_int y) __NOEXC { return __rootn(x, y); } -s::cl_half rootn(s::cl_half x, s::cl_int y) __NOEXC { return __rootn(x, y); } MAKE_1V_2V(rootn, s::cl_float, s::cl_float, s::cl_int) MAKE_1V_2V(rootn, s::cl_double, s::cl_double, s::cl_int) MAKE_1V_2V(rootn, s::cl_half, s::cl_half, s::cl_int) // round -s::cl_float round(s::cl_float x) __NOEXC { return std::round(x); } -s::cl_double round(s::cl_double x) __NOEXC { return std::round(x); } -s::cl_half round(s::cl_half x) __NOEXC { return std::round(x); } +__SYCL_EXPORT s::cl_float round(s::cl_float x) __NOEXC { return std::round(x); } +__SYCL_EXPORT s::cl_double round(s::cl_double x) __NOEXC { + return std::round(x); +} +__SYCL_EXPORT s::cl_half round(s::cl_half x) __NOEXC { return std::round(x); } MAKE_1V(round, s::cl_float, s::cl_float) MAKE_1V(round, s::cl_double, s::cl_double) MAKE_1V(round, s::cl_half, s::cl_half) // rsqrt -s::cl_float rsqrt(s::cl_float x) __NOEXC { return __rsqrt(x); } -s::cl_double rsqrt(s::cl_double x) __NOEXC { return __rsqrt(x); } -s::cl_half rsqrt(s::cl_half x) __NOEXC { return __rsqrt(x); } +__SYCL_EXPORT s::cl_float rsqrt(s::cl_float x) __NOEXC { return __rsqrt(x); } +__SYCL_EXPORT s::cl_double rsqrt(s::cl_double x) __NOEXC { return __rsqrt(x); } +__SYCL_EXPORT s::cl_half rsqrt(s::cl_half x) __NOEXC { return __rsqrt(x); } MAKE_1V(rsqrt, s::cl_float, s::cl_float) MAKE_1V(rsqrt, s::cl_double, s::cl_double) MAKE_1V(rsqrt, s::cl_half, s::cl_half) // sin -s::cl_float sin(s::cl_float x) __NOEXC { return std::sin(x); } -s::cl_double sin(s::cl_double x) __NOEXC { return std::sin(x); } -s::cl_half sin(s::cl_half x) __NOEXC { return std::sin(x); } +__SYCL_EXPORT s::cl_float sin(s::cl_float x) __NOEXC { return std::sin(x); } +__SYCL_EXPORT s::cl_double sin(s::cl_double x) __NOEXC { return std::sin(x); } +__SYCL_EXPORT s::cl_half sin(s::cl_half x) __NOEXC { return std::sin(x); } MAKE_1V(sin, s::cl_float, s::cl_float) MAKE_1V(sin, s::cl_double, s::cl_double) MAKE_1V(sin, s::cl_half, s::cl_half) // sincos -s::cl_float sincos(s::cl_float x, s::cl_float *cosval) __NOEXC { +__SYCL_EXPORT s::cl_float sincos(s::cl_float x, s::cl_float *cosval) __NOEXC { return __sincos(x, cosval); } -s::cl_double sincos(s::cl_double x, s::cl_double *cosval) __NOEXC { +__SYCL_EXPORT s::cl_double sincos(s::cl_double x, + s::cl_double *cosval) __NOEXC { return __sincos(x, cosval); } -s::cl_half sincos(s::cl_half x, s::cl_half *cosval) __NOEXC { +__SYCL_EXPORT s::cl_half sincos(s::cl_half x, s::cl_half *cosval) __NOEXC { return __sincos(x, cosval); } MAKE_1V_2P(sincos, s::cl_float, s::cl_float, s::cl_float) @@ -689,185 +758,239 @@ MAKE_1V_2P(sincos, s::cl_double, s::cl_double, s::cl_double) MAKE_1V_2P(sincos, s::cl_half, s::cl_half, s::cl_half) // sinh -s::cl_float sinh(s::cl_float x) __NOEXC { return std::sinh(x); } -s::cl_double sinh(s::cl_double x) __NOEXC { return std::sinh(x); } -s::cl_half sinh(s::cl_half x) __NOEXC { return std::sinh(x); } +__SYCL_EXPORT s::cl_float sinh(s::cl_float x) __NOEXC { return std::sinh(x); } +__SYCL_EXPORT s::cl_double sinh(s::cl_double x) __NOEXC { return std::sinh(x); } +__SYCL_EXPORT s::cl_half sinh(s::cl_half x) __NOEXC { return std::sinh(x); } MAKE_1V(sinh, s::cl_float, s::cl_float) MAKE_1V(sinh, s::cl_double, s::cl_double) MAKE_1V(sinh, s::cl_half, s::cl_half) // sinpi -s::cl_float sinpi(s::cl_float x) __NOEXC { return __sinpi(x); } -s::cl_double sinpi(s::cl_double x) __NOEXC { return __sinpi(x); } -s::cl_half sinpi(s::cl_half x) __NOEXC { return __sinpi(x); } +__SYCL_EXPORT s::cl_float sinpi(s::cl_float x) __NOEXC { return __sinpi(x); } +__SYCL_EXPORT s::cl_double sinpi(s::cl_double x) __NOEXC { return __sinpi(x); } +__SYCL_EXPORT s::cl_half sinpi(s::cl_half x) __NOEXC { return __sinpi(x); } MAKE_1V(sinpi, s::cl_float, s::cl_float) MAKE_1V(sinpi, s::cl_double, s::cl_double) MAKE_1V(sinpi, s::cl_half, s::cl_half) // sqrt -s::cl_float sqrt(s::cl_float x) __NOEXC { return std::sqrt(x); } -s::cl_double sqrt(s::cl_double x) __NOEXC { return std::sqrt(x); } -s::cl_half sqrt(s::cl_half x) __NOEXC { return std::sqrt(x); } +__SYCL_EXPORT s::cl_float sqrt(s::cl_float x) __NOEXC { return std::sqrt(x); } +__SYCL_EXPORT s::cl_double sqrt(s::cl_double x) __NOEXC { return std::sqrt(x); } +__SYCL_EXPORT s::cl_half sqrt(s::cl_half x) __NOEXC { return std::sqrt(x); } MAKE_1V(sqrt, s::cl_float, s::cl_float) MAKE_1V(sqrt, s::cl_double, s::cl_double) MAKE_1V(sqrt, s::cl_half, s::cl_half) // tan -s::cl_float tan(s::cl_float x) __NOEXC { return std::tan(x); } -s::cl_double tan(s::cl_double x) __NOEXC { return std::tan(x); } -s::cl_half tan(s::cl_half x) __NOEXC { return std::tan(x); } +__SYCL_EXPORT s::cl_float tan(s::cl_float x) __NOEXC { return std::tan(x); } +__SYCL_EXPORT s::cl_double tan(s::cl_double x) __NOEXC { return std::tan(x); } +__SYCL_EXPORT s::cl_half tan(s::cl_half x) __NOEXC { return std::tan(x); } MAKE_1V(tan, s::cl_float, s::cl_float) MAKE_1V(tan, s::cl_double, s::cl_double) MAKE_1V(tan, s::cl_half, s::cl_half) // tanh -s::cl_float tanh(s::cl_float x) __NOEXC { return std::tanh(x); } -s::cl_double tanh(s::cl_double x) __NOEXC { return std::tanh(x); } -s::cl_half tanh(s::cl_half x) __NOEXC { return std::tanh(x); } +__SYCL_EXPORT s::cl_float tanh(s::cl_float x) __NOEXC { return std::tanh(x); } +__SYCL_EXPORT s::cl_double tanh(s::cl_double x) __NOEXC { return std::tanh(x); } +__SYCL_EXPORT s::cl_half tanh(s::cl_half x) __NOEXC { return std::tanh(x); } MAKE_1V(tanh, s::cl_float, s::cl_float) MAKE_1V(tanh, s::cl_double, s::cl_double) MAKE_1V(tanh, s::cl_half, s::cl_half) // tanpi -s::cl_float tanpi(s::cl_float x) __NOEXC { return __tanpi(x); } -s::cl_double tanpi(s::cl_double x) __NOEXC { return __tanpi(x); } -s::cl_half tanpi(s::cl_half x) __NOEXC { return __tanpi(x); } +__SYCL_EXPORT s::cl_float tanpi(s::cl_float x) __NOEXC { return __tanpi(x); } +__SYCL_EXPORT s::cl_double tanpi(s::cl_double x) __NOEXC { return __tanpi(x); } +__SYCL_EXPORT s::cl_half tanpi(s::cl_half x) __NOEXC { return __tanpi(x); } MAKE_1V(tanpi, s::cl_float, s::cl_float) MAKE_1V(tanpi, s::cl_double, s::cl_double) MAKE_1V(tanpi, s::cl_half, s::cl_half) // tgamma -s::cl_float tgamma(s::cl_float x) __NOEXC { return std::tgamma(x); } -s::cl_double tgamma(s::cl_double x) __NOEXC { return std::tgamma(x); } -s::cl_half tgamma(s::cl_half x) __NOEXC { return std::tgamma(x); } +__SYCL_EXPORT s::cl_float tgamma(s::cl_float x) __NOEXC { + return std::tgamma(x); +} +__SYCL_EXPORT s::cl_double tgamma(s::cl_double x) __NOEXC { + return std::tgamma(x); +} +__SYCL_EXPORT s::cl_half tgamma(s::cl_half x) __NOEXC { return std::tgamma(x); } MAKE_1V(tgamma, s::cl_float, s::cl_float) MAKE_1V(tgamma, s::cl_double, s::cl_double) MAKE_1V(tgamma, s::cl_half, s::cl_half) // trunc -s::cl_float trunc(s::cl_float x) __NOEXC { return std::trunc(x); } -s::cl_double trunc(s::cl_double x) __NOEXC { return std::trunc(x); } -s::cl_half trunc(s::cl_half x) __NOEXC { return std::trunc(x); } +__SYCL_EXPORT s::cl_float trunc(s::cl_float x) __NOEXC { return std::trunc(x); } +__SYCL_EXPORT s::cl_double trunc(s::cl_double x) __NOEXC { + return std::trunc(x); +} +__SYCL_EXPORT s::cl_half trunc(s::cl_half x) __NOEXC { return std::trunc(x); } MAKE_1V(trunc, s::cl_float, s::cl_float) MAKE_1V(trunc, s::cl_double, s::cl_double) MAKE_1V(trunc, s::cl_half, s::cl_half) // --------------- 4.13.3 Native Math functions. Host implementations. --------- // native_cos -s::cl_float native_cos(s::cl_float x) __NOEXC { return std::cos(x); } +__SYCL_EXPORT s::cl_float native_cos(s::cl_float x) __NOEXC { + return std::cos(x); +} MAKE_1V(native_cos, s::cl_float, s::cl_float) // native_divide -s::cl_float native_divide(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float native_divide(s::cl_float x, s::cl_float y) __NOEXC { return x / y; } MAKE_1V_2V(native_divide, s::cl_float, s::cl_float, s::cl_float) // native_exp -s::cl_float native_exp(s::cl_float x) __NOEXC { return std::exp(x); } +__SYCL_EXPORT s::cl_float native_exp(s::cl_float x) __NOEXC { + return std::exp(x); +} MAKE_1V(native_exp, s::cl_float, s::cl_float) // native_exp2 -s::cl_float native_exp2(s::cl_float x) __NOEXC { return std::exp2(x); } +__SYCL_EXPORT s::cl_float native_exp2(s::cl_float x) __NOEXC { + return std::exp2(x); +} MAKE_1V(native_exp2, s::cl_float, s::cl_float) // native_exp10 -s::cl_float native_exp10(s::cl_float x) __NOEXC { return std::pow(10, x); } +__SYCL_EXPORT s::cl_float native_exp10(s::cl_float x) __NOEXC { + return std::pow(10, x); +} MAKE_1V(native_exp10, s::cl_float, s::cl_float) // native_log -s::cl_float native_log(s::cl_float x) __NOEXC { return std::log(x); } +__SYCL_EXPORT s::cl_float native_log(s::cl_float x) __NOEXC { + return std::log(x); +} MAKE_1V(native_log, s::cl_float, s::cl_float) // native_log2 -s::cl_float native_log2(s::cl_float x) __NOEXC { return std::log2(x); } +__SYCL_EXPORT s::cl_float native_log2(s::cl_float x) __NOEXC { + return std::log2(x); +} MAKE_1V(native_log2, s::cl_float, s::cl_float) // native_log10 -s::cl_float native_log10(s::cl_float x) __NOEXC { return std::log10(x); } +__SYCL_EXPORT s::cl_float native_log10(s::cl_float x) __NOEXC { + return std::log10(x); +} MAKE_1V(native_log10, s::cl_float, s::cl_float) // native_powr -s::cl_float native_powr(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float native_powr(s::cl_float x, s::cl_float y) __NOEXC { return (x >= 0 ? std::pow(x, y) : x); } MAKE_1V_2V(native_powr, s::cl_float, s::cl_float, s::cl_float) // native_recip -s::cl_float native_recip(s::cl_float x) __NOEXC { return 1.0 / x; } +__SYCL_EXPORT s::cl_float native_recip(s::cl_float x) __NOEXC { + return 1.0 / x; +} MAKE_1V(native_recip, s::cl_float, s::cl_float) // native_rsqrt -s::cl_float native_rsqrt(s::cl_float x) __NOEXC { return 1.0 / std::sqrt(x); } +__SYCL_EXPORT s::cl_float native_rsqrt(s::cl_float x) __NOEXC { + return 1.0 / std::sqrt(x); +} MAKE_1V(native_rsqrt, s::cl_float, s::cl_float) // native_sin -s::cl_float native_sin(s::cl_float x) __NOEXC { return std::sin(x); } +__SYCL_EXPORT s::cl_float native_sin(s::cl_float x) __NOEXC { + return std::sin(x); +} MAKE_1V(native_sin, s::cl_float, s::cl_float) // native_sqrt -s::cl_float native_sqrt(s::cl_float x) __NOEXC { return std::sqrt(x); } +__SYCL_EXPORT s::cl_float native_sqrt(s::cl_float x) __NOEXC { + return std::sqrt(x); +} MAKE_1V(native_sqrt, s::cl_float, s::cl_float) // native_tan -s::cl_float native_tan(s::cl_float x) __NOEXC { return std::tan(x); } +__SYCL_EXPORT s::cl_float native_tan(s::cl_float x) __NOEXC { + return std::tan(x); +} MAKE_1V(native_tan, s::cl_float, s::cl_float) // ---------- 4.13.3 Half Precision Math functions. Host implementations. ------ // half_cos -s::cl_float half_cos(s::cl_float x) __NOEXC { return std::cos(x); } +__SYCL_EXPORT s::cl_float half_cos(s::cl_float x) __NOEXC { + return std::cos(x); +} MAKE_1V(half_cos, s::cl_float, s::cl_float) // half_divide -s::cl_float half_divide(s::cl_float x, s::cl_float y) __NOEXC { return x / y; } +__SYCL_EXPORT s::cl_float half_divide(s::cl_float x, s::cl_float y) __NOEXC { + return x / y; +} MAKE_1V_2V(half_divide, s::cl_float, s::cl_float, s::cl_float) // half_exp -s::cl_float half_exp(s::cl_float x) __NOEXC { return std::exp(x); } +__SYCL_EXPORT s::cl_float half_exp(s::cl_float x) __NOEXC { + return std::exp(x); +} MAKE_1V(half_exp, s::cl_float, s::cl_float) // half_exp2 -s::cl_float half_exp2(s::cl_float x) __NOEXC { return std::exp2(x); } +__SYCL_EXPORT s::cl_float half_exp2(s::cl_float x) __NOEXC { + return std::exp2(x); +} MAKE_1V(half_exp2, s::cl_float, s::cl_float) // half_exp10 -s::cl_float half_exp10(s::cl_float x) __NOEXC { return std::pow(10, x); } +__SYCL_EXPORT s::cl_float half_exp10(s::cl_float x) __NOEXC { + return std::pow(10, x); +} MAKE_1V(half_exp10, s::cl_float, s::cl_float) // half_log -s::cl_float half_log(s::cl_float x) __NOEXC { return std::log(x); } +__SYCL_EXPORT s::cl_float half_log(s::cl_float x) __NOEXC { + return std::log(x); +} MAKE_1V(half_log, s::cl_float, s::cl_float) // half_log2 -s::cl_float half_log2(s::cl_float x) __NOEXC { return std::log2(x); } +__SYCL_EXPORT s::cl_float half_log2(s::cl_float x) __NOEXC { + return std::log2(x); +} MAKE_1V(half_log2, s::cl_float, s::cl_float) // half_log10 -s::cl_float half_log10(s::cl_float x) __NOEXC { return std::log10(x); } +__SYCL_EXPORT s::cl_float half_log10(s::cl_float x) __NOEXC { + return std::log10(x); +} MAKE_1V(half_log10, s::cl_float, s::cl_float) // half_powr -s::cl_float half_powr(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_float half_powr(s::cl_float x, s::cl_float y) __NOEXC { return (x >= 0 ? std::pow(x, y) : x); } MAKE_1V_2V(half_powr, s::cl_float, s::cl_float, s::cl_float) // half_recip -s::cl_float half_recip(s::cl_float x) __NOEXC { return 1.0 / x; } +__SYCL_EXPORT s::cl_float half_recip(s::cl_float x) __NOEXC { return 1.0 / x; } MAKE_1V(half_recip, s::cl_float, s::cl_float) // half_rsqrt -s::cl_float half_rsqrt(s::cl_float x) __NOEXC { return 1.0 / std::sqrt(x); } +__SYCL_EXPORT s::cl_float half_rsqrt(s::cl_float x) __NOEXC { + return 1.0 / std::sqrt(x); +} MAKE_1V(half_rsqrt, s::cl_float, s::cl_float) // half_sin -s::cl_float half_sin(s::cl_float x) __NOEXC { return std::sin(x); } +__SYCL_EXPORT s::cl_float half_sin(s::cl_float x) __NOEXC { + return std::sin(x); +} MAKE_1V(half_sin, s::cl_float, s::cl_float) // half_sqrt -s::cl_float half_sqrt(s::cl_float x) __NOEXC { return std::sqrt(x); } +__SYCL_EXPORT s::cl_float half_sqrt(s::cl_float x) __NOEXC { + return std::sqrt(x); +} MAKE_1V(half_sqrt, s::cl_float, s::cl_float) // half_tan -s::cl_float half_tan(s::cl_float x) __NOEXC { return std::tan(x); } +__SYCL_EXPORT s::cl_float half_tan(s::cl_float x) __NOEXC { + return std::tan(x); +} MAKE_1V(half_tan, s::cl_float, s::cl_float) } // namespace __host_std diff --git a/sycl/source/detail/builtins_relational.cpp b/sycl/source/detail/builtins_relational.cpp index 289b40ecbf237..80611e26d7f00 100644 --- a/sycl/source/detail/builtins_relational.cpp +++ b/sycl/source/detail/builtins_relational.cpp @@ -132,13 +132,13 @@ template inline T2 __vselect(T2 a, T2 b, T c) { // ---------- 4.13.7 Relational functions. Host implementations. --------------- // FOrdEqual-isequal -s::cl_int FOrdEqual(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdEqual(s::cl_float x, s::cl_float y) __NOEXC { return __sFOrdEqual(x, y); } -s::cl_int FOrdEqual(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdEqual(s::cl_double x, s::cl_double y) __NOEXC { return __sFOrdEqual(x, y); } -s::cl_int FOrdEqual(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdEqual(s::cl_half x, s::cl_half y) __NOEXC { return __sFOrdEqual(x, y); } MAKE_1V_2V_FUNC(FOrdEqual, __vFOrdEqual, s::cl_int, s::cl_float, s::cl_float) @@ -146,13 +146,13 @@ MAKE_1V_2V_FUNC(FOrdEqual, __vFOrdEqual, s::cl_long, s::cl_double, s::cl_double) MAKE_1V_2V_FUNC(FOrdEqual, __vFOrdEqual, s::cl_short, s::cl_half, s::cl_half) // FUnordNotEqual-isnotequal -s::cl_int FUnordNotEqual(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_int FUnordNotEqual(s::cl_float x, s::cl_float y) __NOEXC { return __sFUnordNotEqual(x, y); } -s::cl_int FUnordNotEqual(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_int FUnordNotEqual(s::cl_double x, s::cl_double y) __NOEXC { return __sFUnordNotEqual(x, y); } -s::cl_int FUnordNotEqual(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_int FUnordNotEqual(s::cl_half x, s::cl_half y) __NOEXC { return __sFUnordNotEqual(x, y); } MAKE_1V_2V_FUNC(FUnordNotEqual, __vFUnordNotEqual, s::cl_int, s::cl_float, @@ -163,13 +163,14 @@ MAKE_1V_2V_FUNC(FUnordNotEqual, __vFUnordNotEqual, s::cl_short, s::cl_half, s::cl_half) // (FOrdGreaterThan) // isgreater -s::cl_int FOrdGreaterThan(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdGreaterThan(s::cl_float x, s::cl_float y) __NOEXC { return __sFOrdGreaterThan(x, y); } -s::cl_int FOrdGreaterThan(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdGreaterThan(s::cl_double x, + s::cl_double y) __NOEXC { return __sFOrdGreaterThan(x, y); } -s::cl_int FOrdGreaterThan(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdGreaterThan(s::cl_half x, s::cl_half y) __NOEXC { return __sFOrdGreaterThan(x, y); } MAKE_1V_2V_FUNC(FOrdGreaterThan, __vFOrdGreaterThan, s::cl_int, s::cl_float, @@ -180,13 +181,16 @@ MAKE_1V_2V_FUNC(FOrdGreaterThan, __vFOrdGreaterThan, s::cl_short, s::cl_half, s::cl_half) // (FOrdGreaterThanEqual) // isgreaterequal -s::cl_int FOrdGreaterThanEqual(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdGreaterThanEqual(s::cl_float x, + s::cl_float y) __NOEXC { return __sFOrdGreaterThanEqual(x, y); } -s::cl_int FOrdGreaterThanEqual(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdGreaterThanEqual(s::cl_double x, + s::cl_double y) __NOEXC { return __sFOrdGreaterThanEqual(x, y); } -s::cl_int FOrdGreaterThanEqual(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdGreaterThanEqual(s::cl_half x, + s::cl_half y) __NOEXC { return __sFOrdGreaterThanEqual(x, y); } MAKE_1V_2V_FUNC(FOrdGreaterThanEqual, __vFOrdGreaterThanEqual, s::cl_int, @@ -197,18 +201,23 @@ MAKE_1V_2V_FUNC(FOrdGreaterThanEqual, __vFOrdGreaterThanEqual, s::cl_short, s::cl_half, s::cl_half) // (FOrdLessThan) // isless -s::cl_int FOrdLessThan(s::cl_float x, s::cl_float y) __NOEXC { return (x < y); } -s::cl_int FOrdLessThan(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdLessThan(s::cl_float x, s::cl_float y) __NOEXC { return (x < y); } -s::cl_int __vFOrdLessThan(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdLessThan(s::cl_double x, s::cl_double y) __NOEXC { + return (x < y); +} +__SYCL_EXPORT s::cl_int __vFOrdLessThan(s::cl_float x, s::cl_float y) __NOEXC { return -(x < y); } -s::cl_long __vFOrdLessThan(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_long __vFOrdLessThan(s::cl_double x, + s::cl_double y) __NOEXC { return -(x < y); } -s::cl_int FOrdLessThan(s::cl_half x, s::cl_half y) __NOEXC { return (x < y); } -s::cl_short __vFOrdLessThan(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdLessThan(s::cl_half x, s::cl_half y) __NOEXC { + return (x < y); +} +__SYCL_EXPORT s::cl_short __vFOrdLessThan(s::cl_half x, s::cl_half y) __NOEXC { return -(x < y); } MAKE_1V_2V_FUNC(FOrdLessThan, __vFOrdLessThan, s::cl_int, s::cl_float, @@ -219,13 +228,15 @@ MAKE_1V_2V_FUNC(FOrdLessThan, __vFOrdLessThan, s::cl_short, s::cl_half, s::cl_half) // (FOrdLessThanEqual) // islessequal -s::cl_int FOrdLessThanEqual(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdLessThanEqual(s::cl_float x, + s::cl_float y) __NOEXC { return __sFOrdLessThanEqual(x, y); } -s::cl_int FOrdLessThanEqual(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdLessThanEqual(s::cl_double x, + s::cl_double y) __NOEXC { return __sFOrdLessThanEqual(x, y); } -s::cl_int FOrdLessThanEqual(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_int FOrdLessThanEqual(s::cl_half x, s::cl_half y) __NOEXC { return __sFOrdLessThanEqual(x, y); } MAKE_1V_2V_FUNC(FOrdLessThanEqual, __vFOrdLessThanEqual, s::cl_int, s::cl_float, @@ -236,13 +247,13 @@ MAKE_1V_2V_FUNC(FOrdLessThanEqual, __vFOrdLessThanEqual, s::cl_short, s::cl_half, s::cl_half) // (LessOrGreater) // islessgreater -s::cl_int LessOrGreater(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_int LessOrGreater(s::cl_float x, s::cl_float y) __NOEXC { return __sLessOrGreater(x, y); } -s::cl_int LessOrGreater(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_int LessOrGreater(s::cl_double x, s::cl_double y) __NOEXC { return __sLessOrGreater(x, y); } -s::cl_int LessOrGreater(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_int LessOrGreater(s::cl_half x, s::cl_half y) __NOEXC { return __sLessOrGreater(x, y); } MAKE_1V_2V_FUNC(LessOrGreater, __vLessOrGreater, s::cl_int, s::cl_float, @@ -253,18 +264,22 @@ MAKE_1V_2V_FUNC(LessOrGreater, __vLessOrGreater, s::cl_short, s::cl_half, s::cl_half) // (IsFinite) // isfinite -s::cl_int IsFinite(s::cl_float x) __NOEXC { return std::isfinite(x); } -s::cl_int IsFinite(s::cl_double x) __NOEXC { return std::isfinite(x); } -s::cl_int __vIsFinite(s::cl_float x) __NOEXC { +__SYCL_EXPORT s::cl_int IsFinite(s::cl_float x) __NOEXC { + return std::isfinite(x); +} +__SYCL_EXPORT s::cl_int IsFinite(s::cl_double x) __NOEXC { + return std::isfinite(x); +} +__SYCL_EXPORT s::cl_int __vIsFinite(s::cl_float x) __NOEXC { return -static_cast(std::isfinite(x)); } -s::cl_long __vIsFinite(s::cl_double x) __NOEXC { +__SYCL_EXPORT s::cl_long __vIsFinite(s::cl_double x) __NOEXC { return -static_cast(std::isfinite(x)); } -s::cl_int IsFinite(s::cl_half x) __NOEXC { +__SYCL_EXPORT s::cl_int IsFinite(s::cl_half x) __NOEXC { return std::isfinite(d::cast_if_host_half(x)); } -s::cl_short __vIsFinite(s::cl_half x) __NOEXC { +__SYCL_EXPORT s::cl_short __vIsFinite(s::cl_half x) __NOEXC { return -static_cast(std::isfinite(d::cast_if_host_half(x))); } MAKE_1V_FUNC(IsFinite, __vIsFinite, s::cl_int, s::cl_float) @@ -272,18 +287,18 @@ MAKE_1V_FUNC(IsFinite, __vIsFinite, s::cl_long, s::cl_double) MAKE_1V_FUNC(IsFinite, __vIsFinite, s::cl_short, s::cl_half) // (IsInf) // isinf -s::cl_int IsInf(s::cl_float x) __NOEXC { return std::isinf(x); } -s::cl_int IsInf(s::cl_double x) __NOEXC { return std::isinf(x); } -s::cl_int __vIsInf(s::cl_float x) __NOEXC { +__SYCL_EXPORT s::cl_int IsInf(s::cl_float x) __NOEXC { return std::isinf(x); } +__SYCL_EXPORT s::cl_int IsInf(s::cl_double x) __NOEXC { return std::isinf(x); } +__SYCL_EXPORT s::cl_int __vIsInf(s::cl_float x) __NOEXC { return -static_cast(std::isinf(x)); } -s::cl_long __vIsInf(s::cl_double x) __NOEXC { +__SYCL_EXPORT s::cl_long __vIsInf(s::cl_double x) __NOEXC { return -static_cast(std::isinf(x)); } -s::cl_int IsInf(s::cl_half x) __NOEXC { +__SYCL_EXPORT s::cl_int IsInf(s::cl_half x) __NOEXC { return std::isinf(d::cast_if_host_half(x)); } -s::cl_short __vIsInf(s::cl_half x) __NOEXC { +__SYCL_EXPORT s::cl_short __vIsInf(s::cl_half x) __NOEXC { return -static_cast(std::isinf(d::cast_if_host_half(x))); } MAKE_1V_FUNC(IsInf, __vIsInf, s::cl_int, s::cl_float) @@ -291,19 +306,19 @@ MAKE_1V_FUNC(IsInf, __vIsInf, s::cl_long, s::cl_double) MAKE_1V_FUNC(IsInf, __vIsInf, s::cl_short, s::cl_half) // (IsNan) // isnan -s::cl_int IsNan(s::cl_float x) __NOEXC { return std::isnan(x); } -s::cl_int IsNan(s::cl_double x) __NOEXC { return std::isnan(x); } -s::cl_int __vIsNan(s::cl_float x) __NOEXC { +__SYCL_EXPORT s::cl_int IsNan(s::cl_float x) __NOEXC { return std::isnan(x); } +__SYCL_EXPORT s::cl_int IsNan(s::cl_double x) __NOEXC { return std::isnan(x); } +__SYCL_EXPORT s::cl_int __vIsNan(s::cl_float x) __NOEXC { return -static_cast(std::isnan(x)); } -s::cl_long __vIsNan(s::cl_double x) __NOEXC { +__SYCL_EXPORT s::cl_long __vIsNan(s::cl_double x) __NOEXC { return -static_cast(std::isnan(x)); } -s::cl_int IsNan(s::cl_half x) __NOEXC { +__SYCL_EXPORT s::cl_int IsNan(s::cl_half x) __NOEXC { return std::isnan(d::cast_if_host_half(x)); } -s::cl_short __vIsNan(s::cl_half x) __NOEXC { +__SYCL_EXPORT s::cl_short __vIsNan(s::cl_half x) __NOEXC { return -static_cast(std::isnan(d::cast_if_host_half(x))); } MAKE_1V_FUNC(IsNan, __vIsNan, s::cl_int, s::cl_float) @@ -311,18 +326,22 @@ MAKE_1V_FUNC(IsNan, __vIsNan, s::cl_long, s::cl_double) MAKE_1V_FUNC(IsNan, __vIsNan, s::cl_short, s::cl_half) // (IsNormal) // isnormal -s::cl_int IsNormal(s::cl_float x) __NOEXC { return std::isnormal(x); } -s::cl_int IsNormal(s::cl_double x) __NOEXC { return std::isnormal(x); } -s::cl_int __vIsNormal(s::cl_float x) __NOEXC { +__SYCL_EXPORT s::cl_int IsNormal(s::cl_float x) __NOEXC { + return std::isnormal(x); +} +__SYCL_EXPORT s::cl_int IsNormal(s::cl_double x) __NOEXC { + return std::isnormal(x); +} +__SYCL_EXPORT s::cl_int __vIsNormal(s::cl_float x) __NOEXC { return -static_cast(std::isnormal(x)); } -s::cl_long __vIsNormal(s::cl_double x) __NOEXC { +__SYCL_EXPORT s::cl_long __vIsNormal(s::cl_double x) __NOEXC { return -static_cast(std::isnormal(x)); } -s::cl_int IsNormal(s::cl_half x) __NOEXC { +__SYCL_EXPORT s::cl_int IsNormal(s::cl_half x) __NOEXC { return std::isnormal(d::cast_if_host_half(x)); } -s::cl_short __vIsNormal(s::cl_half x) __NOEXC { +__SYCL_EXPORT s::cl_short __vIsNormal(s::cl_half x) __NOEXC { return -static_cast(std::isnormal(d::cast_if_host_half(x))); } MAKE_1V_FUNC(IsNormal, __vIsNormal, s::cl_int, s::cl_float) @@ -330,13 +349,13 @@ MAKE_1V_FUNC(IsNormal, __vIsNormal, s::cl_long, s::cl_double) MAKE_1V_FUNC(IsNormal, __vIsNormal, s::cl_short, s::cl_half) // (Ordered) // isordered -s::cl_int Ordered(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_int Ordered(s::cl_float x, s::cl_float y) __NOEXC { return __vOrdered(x, y); } -s::cl_int Ordered(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_int Ordered(s::cl_double x, s::cl_double y) __NOEXC { return __vOrdered(x, y); } -s::cl_int Ordered(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_int Ordered(s::cl_half x, s::cl_half y) __NOEXC { return __vOrdered(x, y); } MAKE_1V_2V_FUNC(Ordered, __vOrdered, s::cl_int, s::cl_float, s::cl_float) @@ -344,13 +363,13 @@ MAKE_1V_2V_FUNC(Ordered, __vOrdered, s::cl_long, s::cl_double, s::cl_double) MAKE_1V_2V_FUNC(Ordered, __vOrdered, s::cl_short, s::cl_half, s::cl_half) // (Unordered) // isunordered -s::cl_int Unordered(s::cl_float x, s::cl_float y) __NOEXC { +__SYCL_EXPORT s::cl_int Unordered(s::cl_float x, s::cl_float y) __NOEXC { return __sUnordered(x, y); } -s::cl_int Unordered(s::cl_double x, s::cl_double y) __NOEXC { +__SYCL_EXPORT s::cl_int Unordered(s::cl_double x, s::cl_double y) __NOEXC { return __sUnordered(x, y); } -s::cl_int Unordered(s::cl_half x, s::cl_half y) __NOEXC { +__SYCL_EXPORT s::cl_int Unordered(s::cl_half x, s::cl_half y) __NOEXC { return __sUnordered(x, y); } MAKE_1V_2V_FUNC(Unordered, __vUnordered, s::cl_int, s::cl_float, s::cl_float) @@ -358,18 +377,22 @@ MAKE_1V_2V_FUNC(Unordered, __vUnordered, s::cl_long, s::cl_double, s::cl_double) MAKE_1V_2V_FUNC(Unordered, __vUnordered, s::cl_short, s::cl_half, s::cl_half) // (SignBitSet) // signbit -s::cl_int SignBitSet(s::cl_float x) __NOEXC { return std::signbit(x); } -s::cl_int SignBitSet(s::cl_double x) __NOEXC { return std::signbit(x); } -s::cl_int __vSignBitSet(s::cl_float x) __NOEXC { +__SYCL_EXPORT s::cl_int SignBitSet(s::cl_float x) __NOEXC { + return std::signbit(x); +} +__SYCL_EXPORT s::cl_int SignBitSet(s::cl_double x) __NOEXC { + return std::signbit(x); +} +__SYCL_EXPORT s::cl_int __vSignBitSet(s::cl_float x) __NOEXC { return -static_cast(std::signbit(x)); } -s::cl_long __vSignBitSet(s::cl_double x) __NOEXC { +__SYCL_EXPORT s::cl_long __vSignBitSet(s::cl_double x) __NOEXC { return -static_cast(std::signbit(x)); } -s::cl_int SignBitSet(s::cl_half x) __NOEXC { +__SYCL_EXPORT s::cl_int SignBitSet(s::cl_half x) __NOEXC { return std::signbit(d::cast_if_host_half(x)); } -s::cl_short __vSignBitSet(s::cl_half x) __NOEXC { +__SYCL_EXPORT s::cl_short __vSignBitSet(s::cl_half x) __NOEXC { return -static_cast(std::signbit(d::cast_if_host_half(x))); } MAKE_1V_FUNC(SignBitSet, __vSignBitSet, s::cl_int, s::cl_float) diff --git a/sycl/source/detail/force_device.cpp b/sycl/source/detail/force_device.cpp index a88b129d86f8d..052d6bbd1fb25 100644 --- a/sycl/source/detail/force_device.cpp +++ b/sycl/source/detail/force_device.cpp @@ -6,9 +6,9 @@ // //===----------------------------------------------------------------------===// -#include #include #include +#include #include #include diff --git a/sycl/include/CL/sycl/detail/force_device.hpp b/sycl/source/detail/force_device.hpp similarity index 100% rename from sycl/include/CL/sycl/detail/force_device.hpp rename to sycl/source/detail/force_device.hpp diff --git a/sycl/source/detail/platform_impl.hpp b/sycl/source/detail/platform_impl.hpp index 8b9fd17e84ba1..80a62b56af96d 100644 --- a/sycl/source/detail/platform_impl.hpp +++ b/sycl/source/detail/platform_impl.hpp @@ -8,12 +8,12 @@ #pragma once #include -#include #include #include #include -#include +#include #include +#include __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp index cdfd14df7a245..a05eb4e4412ff 100644 --- a/sycl/source/detail/program_manager/program_manager.hpp +++ b/sycl/source/detail/program_manager/program_manager.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include #include @@ -23,12 +24,12 @@ /// Executed as a part of current module's (.exe, .dll) static initialization. /// Registers device executable images with the runtime. -extern "C" void __sycl_register_lib(pi_device_binaries desc); +extern "C" __SYCL_EXPORT void __sycl_register_lib(pi_device_binaries desc); /// Executed as a part of current module's (.exe, .dll) static /// de-initialization. /// Unregisters device executable images with the runtime. -extern "C" void __sycl_unregister_lib(pi_device_binaries desc); +extern "C" __SYCL_EXPORT void __sycl_unregister_lib(pi_device_binaries desc); // +++ } diff --git a/sycl/source/device.cpp b/sycl/source/device.cpp index 9db8763602cf9..8d28be7ac855b 100644 --- a/sycl/source/device.cpp +++ b/sycl/source/device.cpp @@ -6,11 +6,12 @@ // //===----------------------------------------------------------------------===// -#include +#include #include #include #include #include +#include __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { @@ -78,7 +79,8 @@ template vector_class device::create_sub_devices(size_t ComputeUnits) const { return impl->create_sub_devices(ComputeUnits); } -template vector_class + +template __SYCL_EXPORT vector_class device::create_sub_devices( size_t ComputeUnits) const; @@ -87,7 +89,8 @@ vector_class device::create_sub_devices(const vector_class &Counts) const { return impl->create_sub_devices(Counts); } -template vector_class + +template __SYCL_EXPORT vector_class device::create_sub_devices( const vector_class &Counts) const; @@ -96,7 +99,8 @@ vector_class device::create_sub_devices( info::partition_affinity_domain AffinityDomain) const { return impl->create_sub_devices(AffinityDomain); } -template vector_class device::create_sub_devices< + +template __SYCL_EXPORT vector_class device::create_sub_devices< info::partition_property::partition_by_affinity_domain>( info::partition_affinity_domain AffinityDomain) const; @@ -111,7 +115,8 @@ device::get_info() const { } #define PARAM_TRAITS_SPEC(param_type, param, ret_type) \ - template ret_type device::get_info() const; + template __SYCL_EXPORT ret_type device::get_info() \ + const; #include diff --git a/sycl/source/device_selector.cpp b/sycl/source/device_selector.cpp index 9ca9e3d14a37a..6eb1a32d13471 100644 --- a/sycl/source/device_selector.cpp +++ b/sycl/source/device_selector.cpp @@ -6,11 +6,11 @@ // //===----------------------------------------------------------------------===// -#include #include #include #include #include +#include // 4.6.1 Device selection class __SYCL_INLINE_NAMESPACE(cl) { diff --git a/sycl/source/event.cpp b/sycl/source/event.cpp index bf0b0780edca9..900a6a06a4af9 100644 --- a/sycl/source/event.cpp +++ b/sycl/source/event.cpp @@ -64,20 +64,22 @@ event::event(shared_ptr_class event_impl) : impl(event_impl) {} #define PARAM_TRAITS_SPEC(param_type, param, ret_type) \ - template <> ret_type event::get_info() const { \ - return impl->get_info(); \ - } + template <> \ + __SYCL_EXPORT ret_type event::get_info() const { \ + return impl->get_info(); \ + } #include #undef PARAM_TRAITS_SPEC #define PARAM_TRAITS_SPEC(param_type, param, ret_type) \ - template <> \ - ret_type event::get_profiling_info() const { \ - impl->wait(impl); \ - return impl->get_profiling_info(); \ - } + template <> \ + __SYCL_EXPORT ret_type event::get_profiling_info() \ + const { \ + impl->wait(impl); \ + return impl->get_profiling_info(); \ + } #include diff --git a/sycl/source/kernel.cpp b/sycl/source/kernel.cpp index 2b4b2631d9e9f..b22be2a49852b 100644 --- a/sycl/source/kernel.cpp +++ b/sycl/source/kernel.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// +#include #include #include #include @@ -16,7 +17,7 @@ namespace sycl { kernel::kernel(cl_kernel ClKernel, const context &SyclContext) : impl(std::make_shared( detail::pi::cast(ClKernel), - detail::getSyclObjImpl(SyclContext))) {} + detail::getSyclObjImpl(SyclContext))) {} cl_kernel kernel::get() const { return impl->get(); } @@ -37,7 +38,8 @@ kernel::get_info() const { } #define PARAM_TRAITS_SPEC(param_type, param, ret_type) \ - template ret_type kernel::get_info() const; + template __SYCL_EXPORT ret_type kernel::get_info() \ + const; #include @@ -50,8 +52,8 @@ kernel::get_work_group_info(const device &dev) const { } #define PARAM_TRAITS_SPEC(param_type, param, ret_type) \ - template ret_type kernel::get_work_group_info( \ - const device &) const; + template __SYCL_EXPORT ret_type \ + kernel::get_work_group_info(const device &) const; #include @@ -73,11 +75,12 @@ kernel::get_sub_group_info( } #define PARAM_TRAITS_SPEC(param_type, param, ret_type) \ - template ret_type kernel::get_sub_group_info( \ - const device &) const; + template __SYCL_EXPORT ret_type \ + kernel::get_sub_group_info(const device &) const; #define PARAM_TRAITS_SPEC_WITH_INPUT(param_type, param, ret_type, in_type) \ - template ret_type kernel::get_sub_group_info( \ - const device &, in_type) const; + template __SYCL_EXPORT ret_type \ + kernel::get_sub_group_info(const device &, in_type) \ + const; #include diff --git a/sycl/source/platform.cpp b/sycl/source/platform.cpp index 077a95531df66..25e2e8bdd9292 100644 --- a/sycl/source/platform.cpp +++ b/sycl/source/platform.cpp @@ -6,11 +6,11 @@ // //===----------------------------------------------------------------------===// -#include #include #include #include #include +#include #include __SYCL_INLINE_NAMESPACE(cl) { @@ -50,7 +50,8 @@ platform::get_info() const { } #define PARAM_TRAITS_SPEC(param_type, param, ret_type) \ - template ret_type platform::get_info() const; + template __SYCL_EXPORT ret_type \ + platform::get_info() const; #include diff --git a/sycl/source/program.cpp b/sycl/source/program.cpp index 54fa78f8a926c..4e24ee9088929 100644 --- a/sycl/source/program.cpp +++ b/sycl/source/program.cpp @@ -86,7 +86,8 @@ program::get_info() const { } #define PARAM_TRAITS_SPEC(param_type, param, ret_type) \ - template ret_type program::get_info() const; + template __SYCL_EXPORT ret_type program::get_info() \ + const; #include diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp index 7cf8e32733b80..1ca3605eb494e 100644 --- a/sycl/source/queue.cpp +++ b/sycl/source/queue.cpp @@ -128,7 +128,8 @@ queue::get_info() const { } #define PARAM_TRAITS_SPEC(param_type, param, ret_type) \ - template ret_type queue::get_info() const; + template __SYCL_EXPORT ret_type queue::get_info() \ + const; #include @@ -142,8 +143,9 @@ template propertyT queue::get_property() const { return impl->get_property(); } -template bool queue::has_property() const; -template property::queue::enable_profiling +template __SYCL_EXPORT bool +queue::has_property() const; +template __SYCL_EXPORT property::queue::enable_profiling queue::get_property() const; bool queue::is_in_order() const { diff --git a/sycl/source/spirv_ops.cpp b/sycl/source/spirv_ops.cpp index 9dc66ede1bb71..3128f6164b7df 100644 --- a/sycl/source/spirv_ops.cpp +++ b/sycl/source/spirv_ops.cpp @@ -15,17 +15,20 @@ // This operation is NOP on HOST as all operations there are blocking and // by the moment this function was called, the operations generating // the __ocl_event_t objects had already been finished. -void __spirv_GroupWaitEvents(__spv::Scope Execution, uint32_t NumEvents, - __ocl_event_t * WaitEvents) noexcept { +__SYCL_EXPORT void __spirv_GroupWaitEvents(__spv::Scope Execution, + uint32_t NumEvents, + __ocl_event_t *WaitEvents) noexcept { } -void __spirv_ControlBarrier(__spv::Scope Execution, __spv::Scope Memory, - uint32_t Semantics) noexcept { +__SYCL_EXPORT void __spirv_ControlBarrier(__spv::Scope Execution, + __spv::Scope Memory, + uint32_t Semantics) noexcept { std::cerr << "Barrier is not supported on the host device yet.\n"; abort(); } -void __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept { +__SYCL_EXPORT void __spirv_MemoryBarrier(__spv::Scope Memory, + uint32_t Semantics) noexcept { // 1. The 'Memory' parameter is ignored on HOST because there is no memory // separation to global and local there. // 2. The 'Semantics' parameter is ignored because there is no need @@ -33,6 +36,7 @@ void __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept { atomic_thread_fence(std::memory_order_seq_cst); } -void __spirv_ocl_prefetch(const char *Ptr, size_t NumBytes) noexcept { +__SYCL_EXPORT void __spirv_ocl_prefetch(const char *Ptr, + size_t NumBytes) noexcept { cl::sycl::detail::PlatformUtil::prefetch(Ptr, NumBytes); } diff --git a/sycl/test/Unit/lit.cfg.py b/sycl/test/Unit/lit.cfg.py index 58d72af9b2627..e3aecaaef5524 100644 --- a/sycl/test/Unit/lit.cfg.py +++ b/sycl/test/Unit/lit.cfg.py @@ -30,6 +30,15 @@ if 'TEMP' in os.environ: config.environment['TEMP'] = os.environ['TEMP'] +if 'SYCL_CONFIG_FILE_NAME' in os.environ: + config.environment['SYCL_CONFIG_FILE_NAME'] = os.environ['SYCL_CONFIG_FILE_NAME'] +else: + # Since SYCL RT can be now statically linked into the unit test binary, + # dynamic library location resolution mechanisms can be incorrect for such + # tests. Provide the runtime with non-existing configuration file name to + # force it load the default configuration. + config.environment['SYCL_CONFIG_FILE_NAME'] = "null.cfg" + # Propagate path to symbolizer for ASan/MSan. for symbolizer in ['ASAN_SYMBOLIZER_PATH', 'MSAN_SYMBOLIZER_PATH']: if symbolizer in os.environ: diff --git a/sycl/test/kernel-and-program/cache.cpp b/sycl/test/kernel-and-program/cache.cpp index 8b5638ac95723..ed38034dd686e 100644 --- a/sycl/test/kernel-and-program/cache.cpp +++ b/sycl/test/kernel-and-program/cache.cpp @@ -1,4 +1,6 @@ // REQUIRES: opencl +// TODO rewrite as unit test +// XFAIL: * // RUN: %clangxx -fsycl -I %sycl_source_dir %s -o %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out diff --git a/sycl/test/program_manager/program_manager.cpp b/sycl/test/program_manager/program_manager.cpp index 64c3a967d5b35..7ef97c9f1325a 100644 --- a/sycl/test/program_manager/program_manager.cpp +++ b/sycl/test/program_manager/program_manager.cpp @@ -2,6 +2,8 @@ // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out +// TODO rewrite as unit test +// XFAIL: * //==--- program_manager.cpp - SYCL program manager test --------------------==// // diff --git a/sycl/unittests/CMakeLists.txt b/sycl/unittests/CMakeLists.txt index 27a9199b958aa..543f17b8d0687 100644 --- a/sycl/unittests/CMakeLists.txt +++ b/sycl/unittests/CMakeLists.txt @@ -1,26 +1,38 @@ add_custom_target(SYCLUnitTests) set_target_properties(SYCLUnitTests PROPERTIES FOLDER "SYCL tests") -# add_sycl_unittest(test_dirname file1.cpp, file2.cpp ...) +# add_sycl_unittest(test_dirname SHARED|OBJECT file1.cpp, file2.cpp ...) # # Will compile the list of files together and link against SYCL. # Produces a binary names `basename(test_dirname)`. -function(add_sycl_unittest test_dirname) +function(add_sycl_unittest test_dirname link_variant) # Enable exception handling for these unit tests set(LLVM_REQUIRES_EH 1) if (MSVC AND CMAKE_BUILD_TYPE MATCHES "Debug") - set(sycl_lib "sycld") + set(sycl_obj_target "sycld_object") + set(sycl_so_target "sycld") else() - set(sycl_lib "sycl") + set(sycl_obj_target "sycl_object") + set(sycl_so_target "sycl") endif() - add_unittest(SYCLUnitTests ${test_dirname} ${ARGN}) + if ("${link_variant}" MATCHES "SHARED") + set(SYCL_LINK_LIBS ${sycl_so_target}) + add_unittest(SYCLUnitTests ${test_dirname} ${ARGN}) + else() + add_unittest(SYCLUnitTests ${test_dirname} + $ ${ARGN}) + target_compile_definitions(${test_dirname} PRIVATE __SYCL_BUILD_SYCL_DLL) + + get_target_property(SYCL_LINK_LIBS ${sycl_so_target} LINK_LIBRARIES) + endif() + target_link_libraries(${test_dirname} PRIVATE - ${sycl_lib} LLVMTestingSupport OpenCL-Headers + ${SYCL_LINK_LIBS} ) target_include_directories(${test_dirname} PRIVATE SYSTEM ${sycl_inc_dir} diff --git a/sycl/unittests/misc/CMakeLists.txt b/sycl/unittests/misc/CMakeLists.txt index 94247c6eec6b5..6587f9edea163 100644 --- a/sycl/unittests/misc/CMakeLists.txt +++ b/sycl/unittests/misc/CMakeLists.txt @@ -1,5 +1,5 @@ set(sycl_lib_dir $) add_definitions(-DSYCL_LIB_DIR="${sycl_lib_dir}") -add_sycl_unittest(MiscTests +add_sycl_unittest(MiscTests SHARED OsUtils.cpp ) diff --git a/sycl/unittests/pi/CMakeLists.txt b/sycl/unittests/pi/CMakeLists.txt index c6ec05f37eb5b..97c46aafa6d08 100644 --- a/sycl/unittests/pi/CMakeLists.txt +++ b/sycl/unittests/pi/CMakeLists.txt @@ -2,7 +2,7 @@ set(CMAKE_CXX_EXTENSIONS OFF) # Enable exception handling for these unit tests set(LLVM_REQUIRES_EH 1) -add_sycl_unittest(PiTests +add_sycl_unittest(PiTests OBJECT EnqueueMemTest.cpp PlatformTest.cpp EventTest.cpp diff --git a/sycl/unittests/pi/cuda/CMakeLists.txt b/sycl/unittests/pi/cuda/CMakeLists.txt index 0d68616bc5d5d..dc23715b1c8d1 100644 --- a/sycl/unittests/pi/cuda/CMakeLists.txt +++ b/sycl/unittests/pi/cuda/CMakeLists.txt @@ -1,5 +1,5 @@ set(LLVM_REQUIRES_EH 1) -add_sycl_unittest(PiCudaTests +add_sycl_unittest(PiCudaTests OBJECT test_base_objects.cpp test_commands.cpp test_device.cpp @@ -12,7 +12,6 @@ add_sycl_unittest(PiCudaTests add_dependencies(PiCudaTests sycl) target_link_libraries(PiCudaTests PRIVATE - sycl LLVMTestingSupport OpenCL-Headers) diff --git a/sycl/unittests/scheduler/CMakeLists.txt b/sycl/unittests/scheduler/CMakeLists.txt index 3a3edfa68c05c..7756129f804c6 100644 --- a/sycl/unittests/scheduler/CMakeLists.txt +++ b/sycl/unittests/scheduler/CMakeLists.txt @@ -1,19 +1,8 @@ -set(clang $) - -set(__cxx_comp_backup ${CMAKE_CXX_COMPILER}) -set(__c_comp_backup ${CMAKE_C_COMPILER}) - -set(CMAKE_C_COMPILER ${clang}) -set(CMAKE_CXX_COMPILER ${clang}) - -add_sycl_unittest(SchedulerTests +add_sycl_unittest(SchedulerTests OBJECT BlockedCommands.cpp FailedCommands.cpp FinishedCmdCleanup.cpp LeafLimit.cpp MemObjCommandCleanup.cpp utils.cpp -) - -set(CMAKE_C_COMPILER ${__c_comp_backup}) -set(CMAKE_CXX_COMPILER ${__cxx_comp_backup}) +) \ No newline at end of file diff --git a/sycl/unittests/thread_safety/CMakeLists.txt b/sycl/unittests/thread_safety/CMakeLists.txt index a0360d95ffdba..78dc6f2190178 100644 --- a/sycl/unittests/thread_safety/CMakeLists.txt +++ b/sycl/unittests/thread_safety/CMakeLists.txt @@ -1,3 +1,3 @@ -add_sycl_unittest(ThreadSafetyTests +add_sycl_unittest(ThreadSafetyTests OBJECT HostAccessorDeadLock.cpp )