Skip to content

[ESIMD] Add infra to support half, bfloat, etc, support sycl::half. #5123

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 9 commits into from
Dec 28, 2021
Merged

[ESIMD] Add infra to support half, bfloat, etc, support sycl::half. #5123

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
5f6043d
[ESIMD] Add infra for half, bfloat, etc. support, support sycl::half.
kbobrovs Dec 7, 2021
108963c
Improve non-std type support
kbobrovs Dec 26, 2021
c2e1c75
Address review comments.
kbobrovs Dec 27, 2021
b71e0f7
Address review comments.
kbobrovs Dec 27, 2021
c07e6ac
clang format
kbobrovs Dec 27, 2021
bd7ea4d
mark some APIs as ESIMD_INLINE to avoid symbol duplication linker error.
kbobrovs Dec 28, 2021
d189cdf
Fixed bug in 1-elem simd_view element type definition.
kbobrovs Dec 28, 2021
7f45e13
clang format
kbobrovs Dec 28, 2021
0c9fe3e
Fix test failures
kbobrovs Dec 28, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 19 additions & 0 deletions sycl/include/CL/sycl/half_type.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -35,6 +35,19 @@

__SYCL_INLINE_NAMESPACE(cl) {
namespace sycl {

namespace ext {
namespace intel {
namespace experimental {
namespace esimd {
namespace detail {
class WrapperElementTypeProxy;
} // namespace detail
} // namespace esimd
} // namespace experimental
} // namespace intel
} // namespace ext

namespace detail {

inline __SYCL_CONSTEXPR_HALF uint16_t float2Half(const float &Val) {
Expand Down Expand Up @@ -255,6 +268,9 @@ class __SYCL_EXPORT half_v2 {
// Initialize underlying data
constexpr explicit half_v2(uint16_t x) : Buf(x) {}

friend class sycl::ext::intel::experimental::esimd::detail::
WrapperElementTypeProxy;

private:
uint16_t Buf;
};
Expand Down Expand Up @@ -391,6 +407,9 @@ class half {

template <typename Key> friend struct std::hash;

friend class sycl::ext::intel::experimental::esimd::detail::
WrapperElementTypeProxy;

private:
StorageT Data;
};
Expand Down
6 changes: 6 additions & 0 deletions sycl/include/sycl/ext/intel/experimental/esimd/common.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,9 @@
#define ESIMD_REGISTER(n) __attribute__((register_num(n)))

#define __ESIMD_API ESIMD_NODEBUG ESIMD_INLINE

#define __ESIMD_UNSUPPORTED_ON_HOST

#else // __SYCL_DEVICE_ONLY__
#define SYCL_ESIMD_KERNEL
#define SYCL_ESIMD_FUNCTION
Expand All @@ -41,6 +44,9 @@
#define ESIMD_REGISTER(n)

#define __ESIMD_API ESIMD_INLINE

#define __ESIMD_UNSUPPORTED_ON_HOST throw cl::sycl::feature_not_supported()

#endif // __SYCL_DEVICE_ONLY__

// Mark a function being noinline
Expand Down
Loading