[SYCL] Renew and sync with SYCL2020 the implementation of queue::para… #4352

v-klochkov · 2021-08-16T22:12:10Z

…llel_for()

Signed-off-by: Vyacheslav N Klochkov vyacheslav.n.klochkov@intel.com

v-klochkov · 2021-08-18T22:26:53Z

This patch is a DRAFT now because it must NOT be merged as is. It is supposed to be combined with XPTI instrumentation changes handling the code-location that is missing (with this patch) in case of call of queue::parallel_for().

…llel_for() Signed-off-by: Vyacheslav N Klochkov <vyacheslav.n.klochkov@intel.com>

My previous commit in this PR removed it and caused 3 LIT test fails. I decided to restore those shortcuts accepting 'offset' to to not break backward compatibility. Signed-off-by: Vyacheslav N Klochkov <vyacheslav.n.klochkov@intel.com>

…reduction_queue_N_reds

Signed-off-by: Vyacheslav N Klochkov <vyacheslav.n.klochkov@intel.com>

…reduction_queue_N_reds

v-klochkov · 2022-04-27T02:18:08Z

Testing is passed except one build failed with unrelated linkage error:
"pi_level_zero.cpp:(.text.piextProgramGetNativeHandle+0x54): undefined reference to `pthread_rwlock_rdlock'"

@tovinkere - please review, thank you.

…reduction_queue_N_reds

v-klochkov · 2022-04-28T05:15:09Z

@tovinkere - from XPTI + performance point of view, would it worth having additional code-duplication and instead of having Variant1 (as in this PR now) have Variant2 with additional code-duplication:

Current/Previous code:
template <typename KernelName = detail::auto_name, typename KernelType> event parallel_for(range<1> NumWorkItems, _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc));

Variant1 (proposed in the current PR):
template <typename KernelName = detail::auto_name, typename... RestT> event parallel_for(range<1> Range, RestT &&...Rest);

Variant2 (with code-duplication. parallel_for without reduction still has CodeLoc parameter):

 template <typename KernelName = detail::auto_name, typename KernelType>
  event parallel_for(range<1> NumWorkItems,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc));

  template <typename KernelName = detail::auto_name, typename... RestT>
  std::enable_if_t<(sizeof...(RestT) > 1 && ext::oneapi::detail::AreAllButLastReductions<RestT...>::value), event>
  parallel_for(range<1> Range, RestT &&...Rest);

If having explicit CodeLoc is noticably better/faster, then Variant2 makes sense.
Please comment/answer this question.

steffenlarsen

LGTM!

tovinkere · 2022-04-28T13:27:25Z

@tovinkere - from XPTI + performance point of view, would it worth having additional code-duplication and instead of having Variant1 (as in this PR now) have Variant2 with additional code-duplication:

Current/Previous code: template <typename KernelName = detail::auto_name, typename KernelType> event parallel_for(range<1> NumWorkItems, _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc));

Variant1 (proposed in the current PR): template <typename KernelName = detail::auto_name, typename... RestT> event parallel_for(range<1> Range, RestT &&...Rest);

Variant2 (with code-duplication. parallel_for without reduction still has CodeLoc parameter):
 template <typename KernelName = detail::auto_name, typename KernelType>
  event parallel_for(range<1> NumWorkItems,
                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc));

  template <typename KernelName = detail::auto_name, typename... RestT>
  std::enable_if_t<(sizeof...(RestT) > 1 && ext::oneapi::detail::AreAllButLastReductions<RestT...>::value), event>
  parallel_for(range<1> Range, RestT &&...Rest);
If having explicit CodeLoc is noticably better/faster, then Variant2 makes sense. Please comment/answer this question.

Explicit CodeLoc is always faster as the capture happens at compile time. I haven't measured the cost of retrieving it from KernelInfo object, so I cannot comment on the overall perf difference.

tovinkere · 2022-04-28T15:22:51Z

sycl/include/CL/sycl/queue.hpp

+  std::enable_if_t<
+      ext::oneapi::detail::AreAllButLastReductions<RestT...>::value, event>
+  parallel_for(nd_range<Dims> Range, RestT &&...Rest) {
+    const detail::code_location CodeLoc = {};


We are sending a dummy code location for now, so can you please add a comment here that says "Actual code location needs to be captured from KernelInfo object" ?

I added the comment in this additional commit: fd064e2
In the process of adding the comment, I noticed some redundancy in parallel_for_impl accepting 'offset' (deprecated shortcuts)

…reduction_queue_N_reds

The 3 deprecated shortcuts accepting 'offset' actually never get code-location argument, thus removed code-loc from arguments there. Signed-off-by: Vyacheslav N Klochkov <vyacheslav.n.klochkov@intel.com>

againull · 2022-04-28T17:08:31Z

sycl/include/CL/sycl/queue.hpp

-                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
-    _CODELOCARG(&CodeLoc);
+  __SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
+  event parallel_for_impl(range<Dims> Range, id<Dims> WorkItemOffset,


Slava, did you intentionally change parallel_for to parallel_for_impl here?
It looks like you wanted to deprecate this API but still leave it as parallel_for because someone might still be using it.

Yes, I did that intentionally. The main entry point for this args-configuration is the func at the line 761.
It would be ambiguous call/functions error without renaming this code here to parallel_for_impl.

Because the entry point is at 761 that does not accept code-loc in the argument, I had to remove CODELOCPARAM(&CodeLoc) at L896

someone might still be using it

May be someone is using it, maybe not, but for sure we have 3 LIT test using parallel_for with 'offset'.
Fails on those tests made me rename this function to parallel_for_impl.

Ok, thank you!

tovinkere

LGTM!

v-klochkov force-pushed the public_vklochkov_reduction_queue_N_reds branch 4 times, most recently from 59d3253 to 654a003 Compare August 18, 2021 19:55

v-klochkov force-pushed the public_vklochkov_reduction_queue_N_reds branch from 654a003 to 9587d2b Compare September 8, 2021 14:08

github-actions bot added the Stale label Mar 8, 2022

v-klochkov force-pushed the public_vklochkov_reduction_queue_N_reds branch from 9587d2b to a83c328 Compare March 22, 2022 19:35

github-actions bot closed this Apr 22, 2022

v-klochkov reopened this Apr 22, 2022

v-klochkov force-pushed the public_vklochkov_reduction_queue_N_reds branch from a83c328 to 9d764bc Compare April 22, 2022 03:24

[SYCL] Renew and sync with SYCL2020 the implementation of queue::para…

5106950

…llel_for() Signed-off-by: Vyacheslav N Klochkov <vyacheslav.n.klochkov@intel.com>

v-klochkov force-pushed the public_vklochkov_reduction_queue_N_reds branch from 9d764bc to 5106950 Compare April 22, 2022 04:13

v-klochkov mentioned this pull request Apr 22, 2022

[SYCL] Add tests for SYCL2020 queue::parallel_for() shortcuts intel/llvm-test-suite#405

Merged

v-klochkov marked this pull request as ready for review April 22, 2022 04:19

v-klochkov requested a review from a team as a code owner April 22, 2022 04:19

v-klochkov requested review from againull and tovinkere April 22, 2022 04:19

v-klochkov added 5 commits April 25, 2022 12:10

Merge remote-tracking branch 'intel_llvm/sycl' into public_vklochkov_…

10acd05

…reduction_queue_N_reds

Fix clang-format issues

05f137d

Signed-off-by: Vyacheslav N Klochkov <vyacheslav.n.klochkov@intel.com>

Merge remote-tracking branch 'intel_llvm/sycl' into public_vklochkov_…

69360c0

…reduction_queue_N_reds

Merge remote-tracking branch 'intel_llvm/sycl' into public_vklochkov_…

08961c4

…reduction_queue_N_reds

v-klochkov mentioned this pull request Apr 27, 2022

[SYCL] Add missing reduction parallel_for shortcuts to queue #6070

Closed

github-actions bot removed the Stale label Apr 28, 2022

Merge remote-tracking branch 'intel_llvm/sycl' into public_vklochkov_…

88ab707

…reduction_queue_N_reds

v-klochkov requested a review from steffenlarsen April 28, 2022 05:02

steffenlarsen previously approved these changes Apr 28, 2022

View reviewed changes

tovinkere requested changes Apr 28, 2022

View reviewed changes

v-klochkov added 2 commits April 28, 2022 09:15

Merge remote-tracking branch 'intel_llvm/sycl' into public_vklochkov_…

b6f6a30

…reduction_queue_N_reds

Added a comment about dummy code-loc. Updated 3 deprecated shortcuts.

fd064e2

The 3 deprecated shortcuts accepting 'offset' actually never get code-location argument, thus removed code-loc from arguments there. Signed-off-by: Vyacheslav N Klochkov <vyacheslav.n.klochkov@intel.com>

v-klochkov dismissed steffenlarsen’s stale review via fd064e2 April 28, 2022 16:50

againull reviewed Apr 28, 2022

View reviewed changes

tovinkere approved these changes Apr 29, 2022

View reviewed changes

steffenlarsen approved these changes Apr 29, 2022

View reviewed changes

againull merged commit e59fb89 into intel:sycl Apr 29, 2022

v-klochkov mentioned this pull request May 3, 2022

[SYCL] Add tests for span reductions intel/llvm-test-suite#1009

Merged

v-klochkov deleted the public_vklochkov_reduction_queue_N_reds branch February 24, 2023 19:41

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[SYCL] Renew and sync with SYCL2020 the implementation of queue::para… #4352

[SYCL] Renew and sync with SYCL2020 the implementation of queue::para… #4352

Uh oh!

v-klochkov commented Aug 16, 2021

Uh oh!

v-klochkov commented Aug 18, 2021

Uh oh!

v-klochkov commented Apr 27, 2022

Uh oh!

v-klochkov commented Apr 28, 2022

Uh oh!

steffenlarsen left a comment

Uh oh!

tovinkere commented Apr 28, 2022

Uh oh!

tovinkere Apr 28, 2022

Uh oh!

v-klochkov Apr 28, 2022

Uh oh!

againull Apr 28, 2022 •

edited

Loading

Uh oh!

v-klochkov Apr 28, 2022 •

edited

Loading

Uh oh!

v-klochkov Apr 28, 2022

Uh oh!

againull Apr 29, 2022

Uh oh!

tovinkere left a comment

Uh oh!

Uh oh!

[SYCL] Renew and sync with SYCL2020 the implementation of queue::para… #4352

[SYCL] Renew and sync with SYCL2020 the implementation of queue::para… #4352

Uh oh!

Conversation

v-klochkov commented Aug 16, 2021

Uh oh!

v-klochkov commented Aug 18, 2021

Uh oh!

v-klochkov commented Apr 27, 2022

Uh oh!

v-klochkov commented Apr 28, 2022

Uh oh!

steffenlarsen left a comment

Choose a reason for hiding this comment

Uh oh!

tovinkere commented Apr 28, 2022

Uh oh!

tovinkere Apr 28, 2022

Choose a reason for hiding this comment

Uh oh!

v-klochkov Apr 28, 2022

Choose a reason for hiding this comment

Uh oh!

againull Apr 28, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

v-klochkov Apr 28, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

v-klochkov Apr 28, 2022

Choose a reason for hiding this comment

Uh oh!

againull Apr 29, 2022

Choose a reason for hiding this comment

Uh oh!

tovinkere left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

againull Apr 28, 2022 •

edited

Loading

v-klochkov Apr 28, 2022 •

edited

Loading