[SYCL] Added the support for std::numeric_limits<cl::sycl::half> #536

shiltian · 2019-08-22T03:26:25Z

Implemented the partial specilization of the template class
std::numeric_limits<HALF_TYPE> where HALF_TYPE is _Float16 on
device side and cl::sycl::detail::half_impl::half on host side.
Also defined some marcros corresponding to its FP32 counterpart
like HLF_MIN, HLF_MAX, etc.

Signed-off-by: Shilei Tian [email protected]

sycl/include/CL/sycl/half_type.hpp

keryell

Just a couple of questions.

sycl/include/CL/sycl/half_type.hpp

bader · 2019-09-06T21:40:38Z

@keryell, @ilyastepykin, could you revisit, please?

sycl/include/CL/sycl/half_type.hpp

sycl/test/basic_tests/half_type.cpp

bader · 2019-09-12T13:16:20Z

@tianshilei1992, we recently enforced new policy "no new warnings" by treating warnings as errors and this patch introduces new warnings on Windows.

D:\buildbot\Product_worker_intel\sycl-win-x64-pr\llvm.src\sycl\include\CL/sycl/half_type.hpp(208): error C2220: warning treated as error - no 'object' file generated
D:\buildbot\Product_worker_intel\sycl-win-x64-pr\llvm.src\sycl\include\CL/sycl/half_type.hpp(208): warning C4305: 'argument': truncation from 'double' to 'const float'
D:\buildbot\Product_worker_intel\sycl-win-x64-pr\llvm.src\sycl\include\CL/sycl/half_type.hpp(212): warning C4305: 'argument': truncation from 'double' to 'const float'

Could you fix these warnings, please?

shiltian · 2019-09-12T13:17:59Z

@bader Will do.

alexeyvoronov-intel · 2019-09-12T13:21:53Z

@tianshilei1992
Can you please not use cl::sycl::half? Just half from global namespace.
Now it is w/a because it is used by Khronos CTS.

Implemented the partial specilization of the template class `std::numeric_limits<HALF_TYPE>` where HALF_TYPE is `_Float16` on device side and `cl::sycl::detail::half_impl::half` on host side. Also defined some marcros corresponding to its FP32 counterpart like `HLF_MIN`, `HLF_MAX`, etc. Signed-off-by: Shilei Tian <[email protected]>

shiltian · 2019-09-12T13:23:55Z

@alexeyvoronov Sorry I didn't get your point. Do you mean in the test case?

bader · 2019-09-12T13:25:09Z

@tianshilei1992
Can you please not use cl::sycl::half? Just half from global namespace.

Why?

I've opened an issue on the SYCL specification document: KhronosGroup/SYCL-Docs#40.

According to the first comment, we shouldn't use ::half.

alexeyvoronov-intel · 2019-09-12T13:26:29Z

@alexeyvoronov Sorry I didn't get your point.

By SYCL specification half is fundamental builtin type and available inside global namespace. But in the specification cl::sycl::half does not exist.

alexeyvoronov-intel · 2019-09-12T13:27:22Z

@tianshilei1992
Can you please not use cl::sycl::half? Just half from global namespace.

Why?

I've opened an issue on the SYCL specification document: KhronosGroup/SYCL-Docs#40.

According to the first comment, we shouldn't use ::half.

I didn't say add a global space specifier, just half.

bader · 2019-09-12T13:29:04Z

@tianshilei1992
Can you please not use cl::sycl::half? Just half from global namespace.

Why?
I've opened an issue on the SYCL specification document: KhronosGroup/SYCL-Docs#40.
According to the first comment, we shouldn't use ::half.

I didn't say add a global space specifier, just half.

What's the difference?

alexeyvoronov-intel · 2019-09-12T13:32:01Z

@tianshilei1992
Can you please not use cl::sycl::half? Just half from global namespace.

Why?
I've opened an issue on the SYCL specification document: KhronosGroup/SYCL-Docs#40.
According to the first comment, we shouldn't use ::half.

I didn't say add a global space specifier, just half.

What's the difference?
well I didn't see the remark(the first comment), but still this is not the specification..
otherwise there is a mixture of using half in this patch . Both half and cl::sycl::half are used.

shiltian · 2019-09-12T13:33:36Z

@alexeyvoronov The half in the global namespace is not added by me. If it is clear, I'll remove the definition in global namespace.

bader · 2019-09-12T13:36:34Z

@alexeyvoronov The half in the global namespace is not added by me. If it is clear, I'll remove the definition in global namespace.

I suggest @alexeyvoronov to make the change he proposed in a separate patch, as it's not related to goal of this change.
It deserves separate review.

sycl/include/CL/sycl/half_type.hpp

alexeyvoronov-intel · 2019-09-12T13:41:13Z

sycl/include/CL/sycl/half_type.hpp

+  static CONSTEXPR_QUALIFIER half round_error() noexcept { return 0.5F; }
+
+  static CONSTEXPR_QUALIFIER half infinity() noexcept {
+    return __builtin_huge_valf();


Does MSVC compiler have __builtin_huge_valf?
Can we use std::numeric_limitscl::sycl::half::infinity() in SYCL lib? For Microsoft Windows we compile SYCL via MSVC compiler.

Isn't std::numeric_limits<cl::sycl::half>::infinity() what we're implementing here?

Hm, I have an idea for an experiment.
There is a code with W/A for half type in source\detail\builtins_math.cpp.
Can you try fix it?

s::cl_half nan(s::cl_ushort nancode) __NOEXC { return s::cl_half(d::quiet_NaN<s::cl_float>()); }

to

s::cl_half nan(s::cl_ushort nancode) __NOEXC { return d::quiet_NaN<s::cl_half>(); // or td::numeric_limits<s::cl_half>::quiet_NaN() if you have a problem with constexpr }

If you can pass the tests, then all is well.
Otherwise your implementation differs from for double and float.

It does. The pre-check-in passes. :) I'll also take a look at the W/A later.

It does. The pre-check-in passes. :) I'll also take a look at the W/A later.

Good. On Microsoft Windows too(I'm worried about)?

Yes. The pre-check-in includes testing on both Linux and Windows.

it's wonderful

alexeyvoronov-intel · 2019-09-12T13:58:39Z

@alexeyvoronov The half in the global namespace is not added by me. If it is clear, I'll remove the definition in global namespace.

I suggest @alexeyvoronov to make the change he proposed in a separate patch, as it's not related to goal of this change.
It deserves separate review.

I can do appropriate edits later if this patch is urgent.
You can ignore my comments.

bader · 2019-09-12T14:08:20Z

@alexeyvoronov The half in the global namespace is not added by me. If it is clear, I'll remove the definition in global namespace.

I suggest @alexeyvoronov to make the change he proposed in a separate patch, as it's not related to goal of this change.
It deserves separate review.

@alexeyvoronov The half in the global namespace is not added by me. If it is clear, I'll remove the definition in global namespace.

I suggest @alexeyvoronov to make the change he proposed in a separate patch, as it's not related to goal of this change.
It deserves separate review.

I can do appropriate edits later if this patch is urgent.
You can ignore my comments.

It's not urgent, but they seems to be unrelated to std::numeric_limits support, so I suggest doing it in separate patch.

alexeyvoronov-intel · 2019-09-12T14:43:11Z

sycl/test/basic_tests/half_type.cpp

@@ -103,17 +103,97 @@ void verify_vec(queue &q) {
  assert(e.get_access<access::mode::read>()[0] == 0);
 }

+void verify_numeric_limits(queue &q) {
+  // Verify on host side


I believe that you can delete this part, because it will tested in kernel side code via // RUN: env SYCL_DEVICE_TYPE=HOST %t.out

No. They have different types during the two compilation stages. One is for _Float16 and another is for the half I implemented.

bader · 2019-09-13T10:01:46Z

Thanks a lot for working on this!

shiltian · 2019-09-13T13:55:33Z

@bader You're welcome! My another PR will be fixed once I get GPU machine. Will be soon. :)

bader requested a review from ilyastepykin August 22, 2019 07:59

ilyastepykin reviewed Aug 23, 2019

View reviewed changes

sycl/include/CL/sycl/half_type.hpp Outdated Show resolved Hide resolved

sycl/include/CL/sycl/half_type.hpp Outdated Show resolved Hide resolved

sycl/include/CL/sycl/half_type.hpp Outdated Show resolved Hide resolved

bader requested a review from ilyastepykin August 29, 2019 17:35

shiltian commented Aug 31, 2019

View reviewed changes

sycl/include/CL/sycl/half_type.hpp Outdated Show resolved Hide resolved

keryell reviewed Aug 31, 2019

View reviewed changes

sycl/include/CL/sycl/half_type.hpp Outdated Show resolved Hide resolved

sycl/include/CL/sycl/half_type.hpp Outdated Show resolved Hide resolved

keryell reviewed Sep 6, 2019

View reviewed changes

sycl/include/CL/sycl/half_type.hpp Outdated Show resolved Hide resolved

sycl/include/CL/sycl/half_type.hpp Outdated Show resolved Hide resolved

sycl/include/CL/sycl/half_type.hpp Outdated Show resolved Hide resolved

keryell previously approved these changes Sep 11, 2019

View reviewed changes

ilyastepykin reviewed Sep 11, 2019

View reviewed changes

sycl/test/basic_tests/half_type.cpp Outdated Show resolved Hide resolved

shiltian dismissed keryell’s stale review via b421718 September 11, 2019 22:49

bader requested a review from ilyastepykin September 12, 2019 06:31

ilyastepykin previously approved these changes Sep 12, 2019

View reviewed changes

shiltian dismissed ilyastepykin’s stale review via 456486b September 12, 2019 13:22

bader approved these changes Sep 12, 2019

View reviewed changes

alexeyvoronov-intel reviewed Sep 12, 2019

View reviewed changes

sycl/include/CL/sycl/half_type.hpp Show resolved Hide resolved

alexeyvoronov-intel reviewed Sep 12, 2019

View reviewed changes

bader merged commit 6edca52 into intel:sycl Sep 13, 2019

shiltian deleted the numeric branch September 13, 2019 12:37

[SYCL] Added the support for std::numeric_limits<cl::sycl::half> #536

[SYCL] Added the support for std::numeric_limits<cl::sycl::half> #536

Uh oh!

Conversation

shiltian commented Aug 22, 2019

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

keryell left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

bader commented Sep 6, 2019

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

bader commented Sep 12, 2019

Uh oh!

shiltian commented Sep 12, 2019

Uh oh!

alexeyvoronov-intel commented Sep 12, 2019

Uh oh!

shiltian commented Sep 12, 2019 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

bader commented Sep 12, 2019

Uh oh!

alexeyvoronov-intel commented Sep 12, 2019

Uh oh!

alexeyvoronov-intel commented Sep 12, 2019

Uh oh!

bader commented Sep 12, 2019

Uh oh!

alexeyvoronov-intel commented Sep 12, 2019

Uh oh!

shiltian commented Sep 12, 2019

Uh oh!

bader commented Sep 12, 2019

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

alexeyvoronov-intel commented Sep 12, 2019

Uh oh!

bader commented Sep 12, 2019

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

bader commented Sep 13, 2019

Uh oh!

shiltian commented Sep 13, 2019

Uh oh!

Uh oh!

shiltian commented Sep 12, 2019 •

edited

Loading