Skip to content

[SYCL] Added support of rounding modes for floating and integer types #1576

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 38 commits into from
May 14, 2020
Merged

[SYCL] Added support of rounding modes for floating and integer types #1576

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
38 commits
Select commit Hold shift + click to select a range
8b4058d
Unsupported on cuda buffer_dev_to_dev.cpp
fadeeval Apr 14, 2020
b70cd8b
Comments fix
fadeeval Apr 15, 2020
8d0b71e
Fix
fadeeval Apr 15, 2020
f8ed32b
Origin
fadeeval Apr 7, 2020
dd50990
Test
fadeeval Apr 7, 2020
132d5d0
Commit
fadeeval Apr 7, 2020
5a0c215
New test
fadeeval Apr 8, 2020
6d25b89
Commit
fadeeval Apr 8, 2020
482e295
New Tests and converts
fadeeval Apr 8, 2020
53db721
Fix types.hpp
fadeeval Apr 9, 2020
eca6c6b
Convert setting
fadeeval Apr 9, 2020
174acc0
Commit
fadeeval Apr 10, 2020
9b842d4
Half test
fadeeval Apr 10, 2020
825874a
Tests fix
fadeeval Apr 10, 2020
e4791bd
Tests fix
fadeeval Apr 20, 2020
97c0f69
half test implementation
fadeeval Apr 20, 2020
df7e870
half test fix
fadeeval Apr 20, 2020
dda9f46
types.hpp fix
fadeeval Apr 21, 2020
756cc9c
long long types converions
fadeeval Apr 22, 2020
59f6745
Tests fix
fadeeval Apr 23, 2020
5ee4be6
convert test fix
fadeeval Apr 23, 2020
8d5bdde
Formatting
fadeeval Apr 23, 2020
582ddbf
Formatting 2
fadeeval Apr 23, 2020
77ccdd1
Formatting 3
fadeeval Apr 23, 2020
b8f770b
redoing templates
fadeeval Apr 24, 2020
aa76104
Tests reconstruction
fadeeval Apr 27, 2020
6d204e0
Fix convert to long long type
fadeeval Apr 27, 2020
62db4e5
Convert to long long correction
fadeeval Apr 27, 2020
9038cc9
Builtins fix
fadeeval Apr 27, 2020
39c7c19
generic_type_traits modifying
fadeeval Apr 27, 2020
58c128b
tepes.hpp improving
fadeeval Apr 28, 2020
1ee5c1a
types.hpp improvment 2
fadeeval Apr 28, 2020
27b79ef
Formatting
fadeeval Apr 29, 2020
dbad799
Formatting 2
fadeeval Apr 29, 2020
3a3fcbe
FIX
fadeeval Apr 30, 2020
c8484e0
vec_convert_half.cpp fix
fadeeval Apr 30, 2020
d2c7a22
tepes.cpp fix
fadeeval Apr 30, 2020
7bb9767
FIX rebase
fadeeval May 12, 2020
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
197 changes: 188 additions & 9 deletions sycl/include/CL/sycl/types.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -199,6 +199,32 @@ using is_int_to_int =
std::integral_constant<bool, std::is_integral<T>::value &&
std::is_integral<R>::value>;

template <typename T, typename R>
using is_sint_to_sint =
std::integral_constant<bool, is_sigeninteger<T>::value &&
is_sigeninteger<R>::value>;

template <typename T, typename R>
using is_uint_to_uint =
std::integral_constant<bool, is_sugeninteger<T>::value &&
is_sugeninteger<R>::value>;

template <typename T, typename R>
using is_sint_to_from_uint = std::integral_constant<
bool, is_sugeninteger<T>::value && is_sigeninteger<R>::value ||
is_sigeninteger<T>::value && is_sugeninteger<R>::value>;

template <typename T, typename R>
using is_sint_to_float =
std::integral_constant<bool, std::is_integral<T>::value &&
!(std::is_unsigned<T>::value) &&
detail::is_floating_point<R>::value>;

template <typename T, typename R>
using is_uint_to_float =
std::integral_constant<bool, std::is_unsigned<T>::value &&
detail::is_floating_point<R>::value>;

template <typename T, typename R>
using is_int_to_float =
std::integral_constant<bool, std::is_integral<T>::value &&
Expand All @@ -213,15 +239,23 @@ template <typename T, typename R>
using is_float_to_float =
std::integral_constant<bool, detail::is_floating_point<T>::value &&
detail::is_floating_point<R>::value>;
template <typename T>
using is_standard_type = std::integral_constant<
bool, detail::is_sgentype<T>::value && !std::is_same<T, long long>::value &&
!std::is_same<T, unsigned long long>::value>;

template <typename T, typename R, rounding_mode roundingMode>
template <typename T, typename R, rounding_mode roundingMode, typename OpenCLT,
typename OpenCLR>
detail::enable_if_t<std::is_same<T, R>::value, R> convertImpl(T Value) {
return Value;
}

#ifndef __SYCL_DEVICE_ONLY__

// Note for float to half conversions, static_cast calls the conversion operator
// implemented for host that takes care of the precision requirements.
template <typename T, typename R, rounding_mode roundingMode>
template <typename T, typename R, rounding_mode roundingMode, typename OpenCLT,
typename OpenCLR>
detail::enable_if_t<!std::is_same<T, R>::value &&
(is_int_to_int<T, R>::value ||
is_int_to_float<T, R>::value ||
Expand All @@ -231,9 +265,9 @@ convertImpl(T Value) {
return static_cast<R>(Value);
}

#ifndef __SYCL_DEVICE_ONLY__
// float to int
template <typename T, typename R, rounding_mode roundingMode>
template <typename T, typename R, rounding_mode roundingMode, typename OpenCLT,
typename OpenCLR>
detail::enable_if_t<is_float_to_int<T, R>::value, R> convertImpl(T Value) {
switch (roundingMode) {
// Round to nearest even is default rounding mode for floating-point types
Expand Down Expand Up @@ -280,16 +314,145 @@ using Rtp = detail::bool_constant<Mode == rounding_mode::rtp>;
template <rounding_mode Mode>
using Rtn = detail::bool_constant<Mode == rounding_mode::rtn>;

// Convert floating-point type to integer type
// convert signed and unsigned types with an equal size and diff names
template <typename T, typename R, rounding_mode roundingMode, typename OpenCLT,
typename OpenCLR>
detail::enable_if_t<!std::is_same<T, R>::value &&
(is_sint_to_sint<T, R>::value ||
is_uint_to_uint<T, R>::value) &&
std::is_same<OpenCLT, OpenCLR>::value,
R>
convertImpl(T Value) {
return static_cast<R>(Value);
}

// signed to signed
#define __SYCL_GENERATE_CONVERT_IMPL(DestType) \
template <typename T, typename R, rounding_mode roundingMode, \
typename OpenCLT, typename OpenCLR> \
detail::enable_if_t<!std::is_same<T, R>::value && \
is_sint_to_sint<T, R>::value && \
(std::is_same<OpenCLR, DestType>::value || \
std::is_same<OpenCLR, signed char>::value && \
std::is_same<DestType, char>::value) && \
!std::is_same<OpenCLT, OpenCLR>::value, \
R> \
convertImpl(T Value) { \
OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
return __spirv_SConvert##_R##DestType(OpValue); \
}

__SYCL_GENERATE_CONVERT_IMPL(char)
__SYCL_GENERATE_CONVERT_IMPL(short)
__SYCL_GENERATE_CONVERT_IMPL(int)
__SYCL_GENERATE_CONVERT_IMPL(long)
__SYCL_GENERATE_CONVERT_IMPL(longlong)

#undef __SYCL_GENERATE_CONVERT_IMPL

// unsigned to unsigned
#define __SYCL_GENERATE_CONVERT_IMPL(DestType) \
template <typename T, typename R, rounding_mode roundingMode, \
typename OpenCLT, typename OpenCLR> \
detail::enable_if_t<!std::is_same<T, R>::value && \
is_uint_to_uint<T, R>::value && \
std::is_same<OpenCLR, DestType>::value && \
!std::is_same<OpenCLT, OpenCLR>::value, \
R> \
convertImpl(T Value) { \
OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
return __spirv_UConvert##_R##DestType(OpValue); \
}

__SYCL_GENERATE_CONVERT_IMPL(uchar)
__SYCL_GENERATE_CONVERT_IMPL(ushort)
__SYCL_GENERATE_CONVERT_IMPL(uint)
__SYCL_GENERATE_CONVERT_IMPL(ulong)

#undef __SYCL_GENERATE_CONVERT_IMPL

// unsigned to (from) signed
template <typename T, typename R, rounding_mode roundingMode, typename OpenCLT,
typename OpenCLR>
detail::enable_if_t<is_sint_to_from_uint<T, R>::value, R> convertImpl(T Value) {
return static_cast<R>(Value);
}

// sint to float
#define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType) \
template <typename T, typename R, rounding_mode roundingMode, \
typename OpenCLT, typename OpenCLR> \
detail::enable_if_t< \
is_sint_to_float<T, R>::value && std::is_same<R, DestType>::value, R> \
convertImpl(T Value) { \
OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
return __spirv_Convert##SPIRVOp##_R##DestType(OpValue); \
}

__SYCL_GENERATE_CONVERT_IMPL(SToF, half)
__SYCL_GENERATE_CONVERT_IMPL(SToF, float)
__SYCL_GENERATE_CONVERT_IMPL(SToF, double)

#undef __SYCL_GENERATE_CONVERT_IMPL

// uint to float
#define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType) \
template <typename T, typename R, rounding_mode roundingMode, \
typename OpenCLT, typename OpenCLR> \
detail::enable_if_t< \
is_uint_to_float<T, R>::value && std::is_same<R, DestType>::value, R> \
convertImpl(T Value) { \
OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
return __spirv_Convert##SPIRVOp##_R##DestType(OpValue); \
}

__SYCL_GENERATE_CONVERT_IMPL(UToF, half)
__SYCL_GENERATE_CONVERT_IMPL(UToF, float)
__SYCL_GENERATE_CONVERT_IMPL(UToF, double)

#undef __SYCL_GENERATE_CONVERT_IMPL

// float to float
#define __SYCL_GENERATE_CONVERT_IMPL(DestType, RoundingMode, \
RoundingModeCondition) \
template <typename T, typename R, rounding_mode roundingMode, \
typename OpenCLT, typename OpenCLR> \
detail::enable_if_t<!std::is_same<T, R>::value && \
is_float_to_float<T, R>::value && \
std::is_same<R, DestType>::value && \
RoundingModeCondition<roundingMode>::value, \
R> \
convertImpl(T Value) { \
OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
return __spirv_FConvert##_R##DestType##_##RoundingMode(OpValue); \
}

#define __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(RoundingMode, \
RoundingModeCondition) \
__SYCL_GENERATE_CONVERT_IMPL(double, RoundingMode, RoundingModeCondition) \
__SYCL_GENERATE_CONVERT_IMPL(float, RoundingMode, RoundingModeCondition) \
__SYCL_GENERATE_CONVERT_IMPL(half, RoundingMode, RoundingModeCondition)

__SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rte, RteOrAutomatic)
__SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtz, Rtz)
__SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtp, Rtp)
__SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtn, Rtn)

#undef __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE
#undef __SYCL_GENERATE_CONVERT_IMPL

// float to int
#define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType, RoundingMode, \
RoundingModeCondition) \
template <typename T, typename R, rounding_mode roundingMode> \
template <typename T, typename R, rounding_mode roundingMode, \
typename OpenCLT, typename OpenCLR> \
detail::enable_if_t<is_float_to_int<T, R>::value && \
std::is_same<R, DestType>::value && \
(std::is_same<OpenCLR, DestType>::value || \
std::is_same<OpenCLR, signed char>::value && \
std::is_same<DestType, char>::value) && \
RoundingModeCondition<roundingMode>::value, \
R> \
convertImpl(T Value) { \
using OpenCLT = cl::sycl::detail::ConvertToOpenCLType_t<T>; \
OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
return __spirv_Convert##SPIRVOp##_R##DestType##_##RoundingMode(OpValue); \
}
Expand Down Expand Up @@ -319,6 +482,18 @@ __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtn, Rtn)
#undef __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE
#undef __SYCL_GENERATE_CONVERT_IMPL

// Back up
template <typename T, typename R, rounding_mode roundingMode, typename OpenCLT,
typename OpenCLR>
detail::enable_if_t<
(!is_standard_type<T>::value && !is_standard_type<OpenCLT>::value ||
!is_standard_type<R>::value && !is_standard_type<OpenCLR>::value) &&
!std::is_same<OpenCLT, OpenCLR>::value,
R>
convertImpl(T Value) {
return static_cast<R>(Value);
}

#endif // __SYCL_DEVICE_ONLY__

} // namespace detail
Expand Down Expand Up @@ -627,9 +802,13 @@ template <typename Type, int NumElements> class vec {
detail::is_floating_point<convertT>::value,
"Unsupported convertT");
vec<convertT, NumElements> Result;
using OpenCLT = detail::ConvertToOpenCLType_t<DataT>;
using OpenCLR = detail::ConvertToOpenCLType_t<convertT>;
for (size_t I = 0; I < NumElements; ++I) {
Result.setValue(
I, detail::convertImpl<DataT, convertT, roundingMode>(getValue(I)));
I,
detail::convertImpl<DataT, convertT, roundingMode, OpenCLT, OpenCLR>(
getValue(I)));
}
return Result;
}
Expand Down
Loading