DPC++ Runtime
Runtime libraries for oneAPI DPC++
builtins.hpp
Go to the documentation of this file.
1 //==------ builtins.hpp - Non-standard SYCL built-in functions -------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
11 #include <CL/sycl/builtins.hpp>
16 
17 #include <CL/__spirv/spirv_ops.hpp>
18 
19 // TODO Decide whether to mark functions with this attribute.
20 #define __NOEXC /*noexcept*/
21 
22 #ifdef __SYCL_DEVICE_ONLY__
23 #define __SYCL_CONSTANT_AS __attribute__((opencl_constant))
24 #else
25 #define __SYCL_CONSTANT_AS
26 #endif
27 
29 namespace sycl {
30 namespace ext {
31 namespace oneapi {
32 namespace experimental {
33 
34 // Provides functionality to print data from kernels in a C way:
35 // - On non-host devices this function is directly mapped to printf from
36 // OpenCL C
37 // - On host device, this function should be equivalent to standard printf
38 // function from C/C++.
39 //
40 // Please refer to corresponding section in OpenCL C specification to find
41 // information about format string and its differences from standard C rules.
42 //
43 // This function is placed under 'experimental' namespace on purpose, because it
44 // has too much caveats you need to be aware of before using it. Please find
45 // them below and read carefully before using it:
46 //
47 // - According to the OpenCL spec, the format string must be
48 // resolvable at compile time i.e. cannot be dynamically created by the
49 // executing program.
50 //
51 // - According to the OpenCL spec, the format string must reside in constant
52 // address space. The constant address space declarations might get "tricky",
53 // see test/built-ins/printf.cpp for examples.
54 // In simple cases (compile-time known string contents, direct declaration of
55 // the format literal inside the printf call, etc.), the compiler should handle
56 // the automatic address space conversion.
57 // FIXME: Once the extension to generic address space is fully supported, the
58 // constant AS version may need to be deprecated.
59 //
60 // - The format string is interpreted according to the OpenCL C spec, where all
61 // data types has fixed size, opposed to C++ types which doesn't guarantee
62 // the exact width of particular data types (except, may be, char). This might
63 // lead to unexpected result, for example: %ld in OpenCL C means that printed
64 // argument has 'long' type which is 64-bit wide by the OpenCL C spec. However,
65 // by C++ spec long is just at least 32-bit wide, so, you need to ensure (by
66 // performing a cast, for example) that if you use %ld specifier, you pass
67 // 64-bit argument to the cl::sycl::experimental::printf
68 //
69 // - OpenCL spec defines several additional features, like, for example, 'v'
70 // modifier which allows to print OpenCL vectors: note that these features are
71 // not available on host device and therefore their usage should be either
72 // guarded using __SYCL_DEVICE_ONLY__ preprocessor macro or avoided in favor
73 // of more portable solutions if needed
74 //
75 template <typename FormatT, typename... Args>
76 int printf(const FormatT *__format, Args... args) {
77 #if defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__)
78  return __spirv_ocl_printf(__format, args...);
79 #else
80  return ::printf(__format, args...);
81 #endif // defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__)
82 }
83 
84 namespace native {
85 
86 // genfloatfh tanh (genfloatfh x)
87 template <typename T>
91  T>
92  tanh(T x) __NOEXC {
93 #if defined(__NVPTX__)
95  _ocl_T arg1 = cl::sycl::detail::convertDataToType<T, _ocl_T>(x);
96  return cl::sycl::detail::convertDataToType<_ocl_T, T>(
97  __clc_native_tanh(arg1));
98 #else
99  return __sycl_std::__invoke_tanh<T>(x);
100 #endif
101 }
102 
103 // genfloath exp2 (genfloath x)
104 template <typename T>
107  exp2(T x) __NOEXC {
108 #if defined(__NVPTX__)
110  _ocl_T arg1 = cl::sycl::detail::convertDataToType<T, _ocl_T>(x);
111  return cl::sycl::detail::convertDataToType<_ocl_T, T>(
112  __clc_native_exp2(arg1));
113 #else
114  return __sycl_std::__invoke_exp2<T>(x);
115 #endif
116 }
117 
118 } // namespace native
119 
120 } // namespace experimental
121 } // namespace oneapi
122 } // namespace ext
123 
124 } // namespace sycl
125 } // __SYCL_INLINE_NAMESPACE(cl)
126 
127 #undef __SYCL_CONSTANT_AS
spirv_ops.hpp
type_traits.hpp
builtins.hpp
cl::sycl::detail::is_contained
Definition: type_list.hpp:54
sycl
Definition: invoke_simd.hpp:68
generic_type_lists.hpp
cl::sycl::detail::ConvertToOpenCLType_t
conditional_t< TryToGetVectorT< SelectMatchingOpenCLType_t< T > >::value, typename TryToGetVectorT< SelectMatchingOpenCLType_t< T > >::type, conditional_t< TryToGetPointerT< SelectMatchingOpenCLType_t< T > >::value, typename TryToGetPointerVecT< SelectMatchingOpenCLType_t< T > >::type, SelectMatchingOpenCLType_t< T > >> ConvertToOpenCLType_t
Definition: generic_type_traits.hpp:472
__SYCL_ALWAYS_INLINE
#define __SYCL_ALWAYS_INLINE
Definition: defines_elementary.hpp:29
generic_type_traits.hpp
__NOEXC
#define __NOEXC
Definition: builtins.hpp:20
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
builtins.hpp
__simd_abi
Definition: simd.hpp:718
cl::sycl::ext::oneapi::experimental::native::tanh
__SYCL_ALWAYS_INLINE sycl::detail::enable_if_t< sycl::detail::is_genfloatf< T >::value||sycl::detail::is_genfloath< T >::value, T > tanh(T x) __NOEXC
Definition: builtins.hpp:92
cl::sycl::ext::oneapi::experimental::printf
int printf(const FormatT *__format, Args... args)
Definition: builtins.hpp:76
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::sycl::ext::oneapi::experimental::native::exp2
__SYCL_ALWAYS_INLINE sycl::detail::enable_if_t< sycl::detail::is_genfloath< T >::value, T > exp2(T x) __NOEXC
Definition: builtins.hpp:107
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12