Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
6 changes: 4 additions & 2 deletions sycl/include/sycl/detail/builtins/builtins.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -63,11 +63,13 @@

#pragma once

#include <sycl/detail/helpers.hpp>
#include <sycl/detail/type_traits.hpp>
#include <sycl/detail/type_traits/vec_marray_traits.hpp>
#include <sycl/detail/vector_convert.hpp>
#include <sycl/marray.hpp> // for marray
#include <sycl/vector.hpp> // for vec
#include <sycl/marray.hpp>
#include <sycl/multi_ptr.hpp>
#include <sycl/vector.hpp>

namespace sycl {
inline namespace _V1 {
Expand Down
51 changes: 51 additions & 0 deletions sycl/include/sycl/detail/fwd/half.hpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,51 @@
//===----------------------------------------------------------------------===//
//
// 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

#include <cstdint>

namespace sycl {
inline namespace _V1 {
namespace detail::half_impl {
class half;

// Several aliases are defined below:
// - StorageT: actual representation of half data type. It is used by scalar
// half values. On device side, it points to some native half data type, while
// on host it is represented by a 16-bit integer that the implementation
// manipulates to emulate half-precision floating-point behavior.
//
// - BIsRepresentationT: data type which is used by built-in functions. It is
// distinguished from StorageT, because on host, we can still operate on the
// wrapper itself and there is no sense in direct usage of underlying data
// type (too many changes required for BIs implementation without any
// foreseeable profits)
//
// - VecElemT: representation of each element in the vector. On device it is
// the same as StorageT to carry a native vector representation, while on
// host it stores the sycl::half implementation directly.
//
// - VecNStorageT: representation of N-element vector of halfs. Follows the
// same logic as VecElemT.
#ifdef __SYCL_DEVICE_ONLY__
using StorageT = _Float16;
using BIsRepresentationT = _Float16;
using VecElemT = _Float16;
#else // SYCL_DEVICE_ONLY
using StorageT = uint16_t;
// No need to extract underlying data type for built-in functions operating on
// host
using BIsRepresentationT = half;
using VecElemT = half;
#endif // SYCL_DEVICE_ONLY
} // namespace detail::half_impl
using half = detail::half_impl::half;

} // namespace _V1
} // namespace sycl
23 changes: 13 additions & 10 deletions sycl/include/sycl/detail/generic_type_traits.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -8,14 +8,11 @@

#pragma once

#include <sycl/access/access.hpp> // for decorated, address_space
#include <sycl/aliases.hpp> // for half, cl_char, cl_double
#include <sycl/detail/helpers.hpp> // for marray
#include <sycl/detail/type_traits.hpp> // for is_gen_based_on_type_s...
#include <sycl/half_type.hpp> // for BIsRepresentationT
#include <sycl/multi_ptr.hpp> // for multi_ptr, address_spa...

#include <sycl/ext/oneapi/bfloat16.hpp> // for bfloat16 storage type.
#include <sycl/access/access.hpp>
#include <sycl/aliases.hpp>
#include <sycl/bit_cast.hpp>
#include <sycl/detail/fwd/half.hpp>
#include <sycl/detail/type_traits.hpp>

#include <cstddef> // for byte
#include <cstdint> // for uint8_t
Expand All @@ -24,6 +21,9 @@

namespace sycl {
inline namespace _V1 {
namespace ext::oneapi {
class bfloat16;
}
namespace detail {
template <typename T>
using is_byte = typename
Expand Down Expand Up @@ -166,13 +166,16 @@ template <typename T> auto convertToOpenCLType(T &&x) {
static_assert(sizeof(OpenCLType) == sizeof(T));
return static_cast<OpenCLType>(x);
} else if constexpr (std::is_same_v<no_ref, half>) {
using OpenCLType = sycl::detail::half_impl::BIsRepresentationT;
// Make it template-param-dependent to compile with incomplete `half`:
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is fine, because for this to be instantiated with sycl::half, the caller has to pass it, meaning that caller has to have a definition for it.

using OpenCLType =
std::enable_if_t<std::is_same_v<no_ref, half>,
sycl::detail::half_impl::BIsRepresentationT>;
static_assert(sizeof(OpenCLType) == sizeof(T));
return static_cast<OpenCLType>(x);
} else if constexpr (std::is_same_v<no_ref, ext::oneapi::bfloat16>) {
// On host, don't interpret BF16 as uint16.
#ifdef __SYCL_DEVICE_ONLY__
using OpenCLType = sycl::ext::oneapi::bfloat16::Bfloat16StorageT;
using OpenCLType = typename no_ref::Bfloat16StorageT;
return sycl::bit_cast<OpenCLType>(x);
#else
return std::forward<T>(x);
Expand Down
1 change: 1 addition & 0 deletions sycl/include/sycl/detail/spirv.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -20,6 +20,7 @@
#include <sycl/access/access.hpp>
#include <sycl/detail/generic_type_traits.hpp>
#include <sycl/id.hpp>
#include <sycl/memory_enums.hpp>
#include <sycl/multi_ptr.hpp>

#if defined(__NVPTX__)
Expand Down
1 change: 1 addition & 0 deletions sycl/include/sycl/detail/usm_impl.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -14,6 +14,7 @@
namespace sycl {
inline namespace _V1 {
class device;
class context;

namespace detail::usm {

Expand Down
3 changes: 2 additions & 1 deletion sycl/include/sycl/detail/vector_convert.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -58,7 +58,8 @@
#include <sycl/exception.hpp> // for errc

#include <sycl/detail/memcpy.hpp>
#include <sycl/ext/oneapi/bfloat16.hpp> // bfloat16
#include <sycl/ext/oneapi/bfloat16.hpp>
#include <sycl/half_type.hpp>
#include <sycl/vector.hpp>

#ifndef __SYCL_DEVICE_ONLY__
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -11,6 +11,7 @@
#pragma once

#include <sycl/detail/address_space_cast.hpp>
#include <sycl/detail/helpers.hpp>
#include <sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp>
#include <sycl/ext/oneapi/properties/properties.hpp>
#include <sycl/group_barrier.hpp>
Expand Down
2 changes: 2 additions & 0 deletions sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,9 @@
#pragma once

#include <sycl/__spirv/spirv_ops.hpp>
#include <sycl/detail/address_space_cast.hpp>
#include <sycl/ext/oneapi/properties/properties.hpp>
#include <sycl/id.hpp>
#include <sycl/vector.hpp>

namespace sycl {
Expand Down
31 changes: 1 addition & 30 deletions sycl/include/sycl/half_type.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,7 @@

#include <sycl/bit_cast.hpp> // for bit_cast
#include <sycl/detail/export.hpp> // for __SYCL_EXPORT
#include <sycl/detail/fwd/half.hpp>

#ifdef __SYCL_DEVICE_ONLY__
#include <sycl/aspects.hpp>
Expand Down Expand Up @@ -154,36 +155,6 @@ inline __SYCL_CONSTEXPR_HALF float half2Float(const uint16_t &Val) {
namespace half_impl {
class half;

// Several aliases are defined below:
// - StorageT: actual representation of half data type. It is used by scalar
// half values. On device side, it points to some native half data type, while
// on host it is represented by a 16-bit integer that the implementation
// manipulates to emulate half-precision floating-point behavior.
//
// - BIsRepresentationT: data type which is used by built-in functions. It is
// distinguished from StorageT, because on host, we can still operate on the
// wrapper itself and there is no sense in direct usage of underlying data
// type (too many changes required for BIs implementation without any
// foreseeable profits)
//
// - VecElemT: representation of each element in the vector. On device it is
// the same as StorageT to carry a native vector representation, while on
// host it stores the sycl::half implementation directly.
//
// - VecNStorageT: representation of N-element vector of halfs. Follows the
// same logic as VecElemT.
#ifdef __SYCL_DEVICE_ONLY__
using StorageT = _Float16;
using BIsRepresentationT = _Float16;
using VecElemT = _Float16;
#else // SYCL_DEVICE_ONLY
using StorageT = uint16_t;
// No need to extract underlying data type for built-in functions operating on
// host
using BIsRepresentationT = half;
using VecElemT = half;
#endif // SYCL_DEVICE_ONLY

// Creation token to disambiguate constructors.
struct RawHostHalfToken {
constexpr explicit RawHostHalfToken(uint16_t Val) : Value{Val} {}
Expand Down
9 changes: 5 additions & 4 deletions sycl/include/sycl/vector.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -31,10 +31,11 @@
#error "SYCL device compiler is built without ext_vector_type support"
#endif

#include <sycl/access/access.hpp> // for decorated, address_space
#include <sycl/aliases.hpp> // for half, cl_char, cl_int
#include <sycl/detail/common.hpp> // for ArrayCreator
#include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEPRECATED
#include <sycl/access/access.hpp> // for decorated, address_space
#include <sycl/aliases.hpp> // for half, cl_char, cl_int
#include <sycl/detail/common.hpp> // for ArrayCreator
#include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEPRECATED
#include <sycl/detail/fwd/accessor.hpp>
#include <sycl/detail/generic_type_traits.hpp> // for is_sigeninteger, is_s...
#include <sycl/detail/memcpy.hpp> // for memcpy
#include <sycl/detail/named_swizzles_mixin.hpp>
Expand Down
11 changes: 6 additions & 5 deletions sycl/test/include_deps/sycl_accessor.hpp.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -43,18 +43,18 @@
// CHECK-NEXT: detail/accessor_iterator.hpp
// CHECK-NEXT: detail/generic_type_traits.hpp
// CHECK-NEXT: aliases.hpp
// CHECK-NEXT: bit_cast.hpp
// CHECK-NEXT: detail/fwd/half.hpp
// CHECK-NEXT: detail/type_traits.hpp
// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
// CHECK-NEXT: detail/fwd/multi_ptr.hpp
// CHECK-NEXT: detail/handler_proxy.hpp
// CHECK-NEXT: multi_ptr.hpp
// CHECK-NEXT: detail/address_space_cast.hpp
// CHECK-NEXT: half_type.hpp
// CHECK-NEXT: bit_cast.hpp
// CHECK-NEXT: aspects.hpp
// CHECK-NEXT: info/aspects.def
// CHECK-NEXT: info/aspects_deprecated.def
// CHECK-NEXT: multi_ptr.hpp
// CHECK-NEXT: detail/address_space_cast.hpp
// CHECK-NEXT: ext/oneapi/bfloat16.hpp
// CHECK-NEXT: detail/handler_proxy.hpp
// CHECK-NEXT: pointers.hpp
// CHECK-NEXT: properties/accessor_properties.hpp
// CHECK-NEXT: properties/runtime_accessor_properties.def
Expand All @@ -67,6 +67,7 @@
// CHECK-NEXT: ext/oneapi/experimental/device_architecture.def
// CHECK-NEXT: ext/oneapi/experimental/forward_progress.hpp
// CHECK-NEXT: ext/oneapi/matrix/query-types.hpp
// CHECK-NEXT: ext/oneapi/bfloat16.hpp
// CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp
// CHECK-NEXT: info/platform_traits.def
// CHECK-NEXT: info/context_traits.def
Expand Down
11 changes: 6 additions & 5 deletions sycl/test/include_deps/sycl_detail_core.hpp.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -44,18 +44,18 @@
// CHECK-NEXT: detail/accessor_iterator.hpp
// CHECK-NEXT: detail/generic_type_traits.hpp
// CHECK-NEXT: aliases.hpp
// CHECK-NEXT: bit_cast.hpp
// CHECK-NEXT: detail/fwd/half.hpp
// CHECK-NEXT: detail/type_traits.hpp
// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
// CHECK-NEXT: detail/fwd/multi_ptr.hpp
// CHECK-NEXT: detail/handler_proxy.hpp
// CHECK-NEXT: multi_ptr.hpp
// CHECK-NEXT: detail/address_space_cast.hpp
// CHECK-NEXT: half_type.hpp
// CHECK-NEXT: bit_cast.hpp
// CHECK-NEXT: aspects.hpp
// CHECK-NEXT: info/aspects.def
// CHECK-NEXT: info/aspects_deprecated.def
// CHECK-NEXT: multi_ptr.hpp
// CHECK-NEXT: detail/address_space_cast.hpp
// CHECK-NEXT: ext/oneapi/bfloat16.hpp
// CHECK-NEXT: detail/handler_proxy.hpp
// CHECK-NEXT: pointers.hpp
// CHECK-NEXT: properties/accessor_properties.hpp
// CHECK-NEXT: properties/runtime_accessor_properties.def
Expand All @@ -68,6 +68,7 @@
// CHECK-NEXT: ext/oneapi/experimental/device_architecture.def
// CHECK-NEXT: ext/oneapi/experimental/forward_progress.hpp
// CHECK-NEXT: ext/oneapi/matrix/query-types.hpp
// CHECK-NEXT: ext/oneapi/bfloat16.hpp
// CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp
// CHECK-NEXT: info/platform_traits.def
// CHECK-NEXT: info/context_traits.def
Expand Down