[SYCL] Added support of rounding modes for floating and integer types#1576
Conversation
67329a5 to
1cd7097
Compare
|
+@Naghasan to review SPIRV built-ins tablegen changes. |
clang/lib/Sema/SPIRVBuiltins.td
Outdated
There was a problem hiding this comment.
| if !ne(OutType.ElementSize, InType.ElementSize) then { | |
| def : SPVBuiltin<"SConvert_R" # OutType.Name, [OutType, InType], Attr.Const>; | |
| if !ne(OutType.ElementSize, InType.ElementSize) then { | |
| def : SPVBuiltin<"SConvert_R" # OutType.Name, [OutType, InType], Attr.Const>; |
sycl/include/CL/sycl/types.hpp
Outdated
There was a problem hiding this comment.
I don't see why you do need this include.
AlexeySachkov
left a comment
There was a problem hiding this comment.
Part about conversions is getting bigger and bigger. We should probably consider outlining it into a separate header file, just to reduce the size of types.hpp and make it easier to read
sycl/include/CL/sycl/types.hpp
Outdated
There was a problem hiding this comment.
!istd::is_unsigned -> std::is_signed
There was a problem hiding this comment.
std::is_signed include floating point types also, is not?
There was a problem hiding this comment.
Yes, is_signed includes is_arithmetic, but you already have is_integral here - I guess this should be enough to reject floating-point types, right?
There was a problem hiding this comment.
I will use "is_sigeninteger" in new version any way.
sycl/include/CL/sycl/types.hpp
Outdated
There was a problem hiding this comment.
I think that you can use the following type trait that we have:
sycl/include/CL/sycl/types.hpp
Outdated
There was a problem hiding this comment.
You have very similar chunk of code for signed to signed. I suggest you to adjust enable_if a bit so it can be re-used for this unsigned to unsigned as well.
Something like
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<R, DestType>::value && \
std::is_same<cl::sycl::detail::ConvertToOpenCLType_t<T>, R>::value, \
R>
sycl/include/CL/sycl/types.hpp
Outdated
There was a problem hiding this comment.
This looks like a lot of work on templates for the compiler. Can we somehow simplify it?
Is it possible to convert SYCL types to OpenCL types somewhere else in the call stack (let's say in a function which calls convertImpl) and re-design convertImpl so it already operates on OpenCL types?
sycl/include/CL/sycl/types.hpp
Outdated
There was a problem hiding this comment.
This std::is_same<R, DestType::value doesn't seem like a right thing to do here. Is it possible to avoid adding this check into each enable_if?
What about creating partial specializations of convertImpl?
template <typename T, typename R, rounding_mode roundingMode>
R convertImp(T Value) {
// the most generic one
// static_cast as fallback?
}
template <typename T, rounding_mode roundingMode>
half convertImpl<T, half, roundingMode>(half Value) {
// actual implementation here
}
// Actual implementations can still be generated by preprocessor:
#define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType) \
detail::enable_if_t< is_signed_arithmetic<T>::value, T> \
template <typename T, rounding_mode roundingMode> \
convertImpl<T, DestType, roundingMode>(DestType Value) { \
// actual implementation here \
}
__SYCL_GENERATE_CONVERT_IMPL(SToF, half)
__SYCL_GENERATE_CONVERT_IMPL(SToF, float)
__SYCL_GENERATE_CONVERT_IMPL(SToF, double)
#undef __SYCL_GENERATE_CONVERT_IMPLThere was a problem hiding this comment.
How about just get name of a type as string and add it to a spirv func name?
There was a problem hiding this comment.
I don't understand how to implement your idea
There was a problem hiding this comment.
C++ doesn't support function template partial specialization.
There was a problem hiding this comment.
| test<double, float, 8, rounding_mode::automatic>( | |
| test<double, float, 8, rounding_mode::rte>( |
There was a problem hiding this comment.
Since you are adding separate tests for particular data types, I suggest you to just remove this test case from here - it is anyway duplicated in vec_convert_f_to_f.cpp
There was a problem hiding this comment.
Please outline duplicated code into a header file with helpers
There was a problem hiding this comment.
There should I place the header?
There was a problem hiding this comment.
Since you are adding separate files for conversions between different types, I suggest you to leave in this file only float to int conversions - just to avoid testing the same combinations in different files
There was a problem hiding this comment.
Maybe I should rename this file then?
There was a problem hiding this comment.
Maybe I should rename this file then?
Makes sense
clang/lib/Sema/SPIRVBuiltins.td
Outdated
There was a problem hiding this comment.
Why this restriction ? It is valid to convert a signed int to an unsigned one.
There was a problem hiding this comment.
Yes, it is valid, but SatConvertSToU exists for this purpose.
There was a problem hiding this comment.
UConvert/SConvert does not saturate the result unless decorated...
There was a problem hiding this comment.
I will check it, but I have doubt, that UConvert gets non unsigned arguments.
There was a problem hiding this comment.
When I run test that converts int to uint, it writes "call to '__spirv_UConvert_Ruint' is ambiguous", that means it make implicit conversion form signed to unsigned so that push int argument, but compiler have many options for that that's why it writes the error.
There was a problem hiding this comment.
I think the document https://www.khronos.org/registry/spir-v/specs/unified1/SPIRV.html#OpUConvert proves my gusses.
There was a problem hiding this comment.
I think the document https://www.khronos.org/registry/spir-v/specs/unified1/SPIRV.html#OpUConvert proves my gusses.
I hardly see what it proves, from the spec This is either a truncate or a zero extend. whether the destination type is bigger or smaller, this is not a saturating operation: Uconvert 0x0F00 to uchar yields 0x0, SatConvertStoU yields 0xFF.
that converts int to uint, it writes "call to '__spirv_UConvert_Ruint' is ambiguous"
This is expected as this operation does not exist in SPIR-V, so the overload does not exist. You are trying to do a 32 to 32 bits conversion (see https://www.khronos.org/registry/spir-v/specs/unified1/SPIRV.html#OpUConvert: The component width cannot equal the component width in Result Type.). The operation you are looking for in this cast is simply a reinterpret_cast in that case.
sycl/include/CL/sycl/types.hpp
Outdated
sycl/include/CL/sycl/types.hpp
Outdated
There was a problem hiding this comment.
Why do you need this as macro? The main idea was to auto-generate some code by using pre-processor features. Now you have no arguments to this macro and call it only once - in that case having a macro just doesn't make sense
|
@turinevgeny, @erichkeane, @Fznamznon, @Naghasan, please, make review. |
erichkeane
left a comment
There was a problem hiding this comment.
The CFE changes are fine, but RT reviewers are going to need to do approval for the rest.
turinevgeny
left a comment
There was a problem hiding this comment.
types.hpp looks fine.
|
@fadeeval, please, resolve merge conflict. |
There was a problem hiding this comment.
vec_convert.cpp -> vec_convert_f_to_f.cpp
There was a problem hiding this comment.
Please, fix the comment in a separate PR.
There was a problem hiding this comment.
vec_convert.cpp -> vec_convert_f_to_i.cpp
There was a problem hiding this comment.
Please, fix the comment in a separate PR.
There was a problem hiding this comment.
| test<float, int, 8, rounding_mode::automatic>( | |
| float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, | |
| int8{2, 2, 3, -2, -2, -3, 0, 0}); | |
| test<float, int, 8, rounding_mode::automatic>( | |
| float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f}, | |
| int8{2, 2, 2, -2, -2, -2, 0, 0}); |
RTZ is the default for conversions to integer
There was a problem hiding this comment.
I think so too, but in the sycl-1.2.1.pdf spec on the 226 page the automatic description is following: "Default rounding mode for the SYCL vec class element type. rtz (round toward zero) for integer types and rte (round to nearest even) for floating-point types", which sounds confusing. And the person who wrote conversions for HOST implemented so that RTE is automatic mode for converting from floating-point types. That is why I followed the idea that automatic mode is RTE. Should I change it to RTZ?
There was a problem hiding this comment.
"Default rounding mode for the SYCL vec class element type. rtz (round toward zero) for integer types and rte (round to nearest even) for floating-point types", which sounds confusing
I agree, RTZ for conversions to integer types could be a better wording for this.
Should I change it to RTZ?
Well that's what the spec mandates. But according to what you are saying, changing this will go well out of scope of your patch. The bug is already there anyway and the patch consistent with it, so I would personally lean toward merging as is and do a PR to fully fix the automatic mode in one go after. But that's more a CO decision now.
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
f834594 to
d2c7a22
Compare
Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>
|
@turinevgeny, @erichkeane, @Fznamznon, @Naghasan, @sergey-semenov, make review, please,. |
turinevgeny
left a comment
There was a problem hiding this comment.
LGTM in general, but I'd like someone else to approve who knows more details.
|
@erichkeane, @Fznamznon, @Naghasan, @sergey-semenov, approve if no objections, please. |
|
I don't see any FE changes, please don't wait for approve from me. |
Naghasan
left a comment
There was a problem hiding this comment.
Looks ok to me. I think RTZ issue for integers should be handled separately.
sergey-semenov
left a comment
There was a problem hiding this comment.
LGTM aside from the unresolved file header comments by @bader, but those can be addressed separately.
Implementing rounding models for cl::sycl::vec type for non host devices.
Signed-off-by: Aleksander Fadeev aleksander.fadeev@intel.com