@@ -804,13 +804,9 @@ NativeToT ConvertToBF16(NativeFromT val) {
804
804
// / conversion kind. It is expected to be \c DataT template argument of a vector
805
805
// / we are trying to convert \a to
806
806
// / \tparam NativeFromT \b scalar or \b vector internal type corresponding to
807
- // / \c FromT, which is used to hold vector data. It is expected to be
808
- // / vec<FromT, VecSize>::vector_t of a vector we are trying to convert \a from
809
- // / if VecSize > 1, or result of detail::ConvertToOpenCLType_t<FromT>
807
+ // / \c FromT, which is used to hold vector data.
810
808
// / \tparam NativeToT \b scalar or \b vector internal type corresponding to
811
- // / \c ToT, which is used to hold vector data. It is expected to be
812
- // / vec<ToT, VecSize>::vector_t of a vector we are trying to convert \a from
813
- // / if VecSize > 1, or result of detail::ConvertToOpenCLType_t<ToT>
809
+ // / \c ToT
814
810
// /
815
811
// / \note Each pair of types FromT, ToT and NativeFromT, NativeToT can't contain
816
812
// / the same type, because there are no no-op convert instructions in SPIR-V.
@@ -911,9 +907,6 @@ vec<convertT, NumElements> vec<DataT, NumElements>::convert() const {
911
907
using OpenCLVecT = OpenCLT __attribute__ ((ext_vector_type (NumElements)));
912
908
using OpenCLVecR = OpenCLR __attribute__ ((ext_vector_type (NumElements)));
913
909
914
- auto NativeVector = sycl::bit_cast<vector_t >(*this );
915
- using ConvertTVecType = typename vec<convertT, NumElements>::vector_t ;
916
-
917
910
// Whole vector conversion can only be done, if:
918
911
constexpr bool canUseNativeVectorConvert =
919
912
#ifdef __NVPTX__
@@ -923,11 +916,6 @@ vec<convertT, NumElements> vec<DataT, NumElements>::convert() const {
923
916
false &&
924
917
#endif
925
918
NumElements > 1 &&
926
- // - vec storage has an equivalent OpenCL native vector it is
927
- // implicitly convertible to. There are some corner cases where it
928
- // is not the case with char, long and long long types.
929
- std::is_convertible_v<vector_t , OpenCLVecT> &&
930
- std::is_convertible_v<ConvertTVecType, OpenCLVecR> &&
931
919
// - it is not a signed to unsigned (or vice versa) conversion
932
920
// see comments within 'convertImpl' for more details;
933
921
!detail::is_sint_to_from_uint<T, R>::value &&
@@ -939,8 +927,9 @@ vec<convertT, NumElements> vec<DataT, NumElements>::convert() const {
939
927
!std::is_same_v<convertT, bool >;
940
928
941
929
if constexpr (canUseNativeVectorConvert) {
942
- auto val = detail::convertImpl<T, R, roundingMode, NumElements,
943
- OpenCLVecT, OpenCLVecR>(NativeVector);
930
+ auto val =
931
+ detail::convertImpl<T, R, roundingMode, NumElements, OpenCLVecT,
932
+ OpenCLVecR>(bit_cast<OpenCLVecT>(*this ));
944
933
Result.m_Data = sycl::bit_cast<decltype (Result.m_Data )>(val);
945
934
} else
946
935
#endif // __SYCL_DEVICE_ONLY__
0 commit comments