Skip to content

Commit f34ba2c

Browse files
authored
[ESIMD] Add infra to support half, bfloat, etc, support sycl::half. (#5123)
- implement infrastructure for non-standard element type support to support a new type, the following must be implemented: * element type traits for the new type * scalar and vector conversions to/from selected std C++ type * std operations (or default with promotion to std C++ type can be used) - esimd::simd<sycl::half, N> can now be used on host and device * most of the operations, except extended math are supported for half Signed-off-by: Konstantin S Bobrovsky <[email protected]>
1 parent 27f59d8 commit f34ba2c

File tree

17 files changed

+1554
-670
lines changed

17 files changed

+1554
-670
lines changed

sycl/include/CL/sycl/half_type.hpp

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -35,6 +35,19 @@
3535

3636
__SYCL_INLINE_NAMESPACE(cl) {
3737
namespace sycl {
38+
39+
namespace ext {
40+
namespace intel {
41+
namespace experimental {
42+
namespace esimd {
43+
namespace detail {
44+
class WrapperElementTypeProxy;
45+
} // namespace detail
46+
} // namespace esimd
47+
} // namespace experimental
48+
} // namespace intel
49+
} // namespace ext
50+
3851
namespace detail {
3952

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

271+
friend class sycl::ext::intel::experimental::esimd::detail::
272+
WrapperElementTypeProxy;
273+
258274
private:
259275
uint16_t Buf;
260276
};
@@ -391,6 +407,9 @@ class half {
391407

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

410+
friend class sycl::ext::intel::experimental::esimd::detail::
411+
WrapperElementTypeProxy;
412+
394413
private:
395414
StorageT Data;
396415
};

sycl/include/sycl/ext/intel/experimental/esimd/common.hpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -29,6 +29,9 @@
2929
#define ESIMD_REGISTER(n) __attribute__((register_num(n)))
3030

3131
#define __ESIMD_API ESIMD_NODEBUG ESIMD_INLINE
32+
33+
#define __ESIMD_UNSUPPORTED_ON_HOST
34+
3235
#else // __SYCL_DEVICE_ONLY__
3336
#define SYCL_ESIMD_KERNEL
3437
#define SYCL_ESIMD_FUNCTION
@@ -41,6 +44,9 @@
4144
#define ESIMD_REGISTER(n)
4245

4346
#define __ESIMD_API ESIMD_INLINE
47+
48+
#define __ESIMD_UNSUPPORTED_ON_HOST throw cl::sycl::feature_not_supported()
49+
4450
#endif // __SYCL_DEVICE_ONLY__
4551

4652
// Mark a function being noinline

0 commit comments

Comments
 (0)