Allow initializing SYCL execution space from sycl::queue and SYCL::impl_static_fence #3767

masterleinad · 2021-02-01T17:17:21Z

Based on #3759.

This pull request allows initializing SYCL execution space from a sycl::queue corresponding to streams for CUDA and SYCL. It also provides a valid implementation for impl_static_fence that simply fences all queues.
I was running into some issues with the current TeamPolicy implementation but didn't bother to look too closely since the alternative implementation in #3759 doesn't show this problem and I believe that we should move forward with that one.
Also, we need to make sure that all memory we allocate using sycl::allocate_* references the correct sycl::queue. Memory access across queues doesn't work with this kind of allocation. Thirs turned out to be a problem in the parallel_reduce implementation.

nliber · 2021-02-01T23:20:50Z

core/src/SYCL/Kokkos_SYCL_Instance.cpp

@@ -55,6 +55,8 @@ namespace Impl {

 int SYCLInternal::was_finalized = 0;

+std::vector<std::optional<sycl::queue>*> SYCLInternal::all_queues;


Why not just a vector<queue>?

m_queue is an optional<queue> to get around vendor bugs when assigning one queue to another. As long as we aren't inserting/emplacing/erasing from the beginning or middle of the vector or assigning to elements in the vector, we won't run into that issue, and it won't be a concern at all once the vendor bugs are fixed.

Also, is m_queue going to eventually go away (I assume so but an checking here)? (Even now it could go away as it should always be equivalent to all_queues.back().

I was running into the same (or what looked like the same) vendor bug and storing pointers seemed to be the only way to make it work, see db9dc39 (#3767) for the fixing commit and https://cloud.cees.ornl.gov/jenkins-ci/blue/organizations/jenkins/Kokkos/detail/Kokkos/4186/pipeline for the error when just using std::vector<sycl::queue>.

Also, is m_queue going to eventually go away (I assume so but an checking here)? (Even now it could go away as it should always be equivalent to all_queues.back().

What do you mean? We still must have access to the queue corresponding to the current instance so I don't see a way to eliminate it. Or am I misunderstanding?

As I read lines 119 and 120, m_queue and all_queues.back() always refer to the same queue, don't they?

That's only true for the last instance that adds a queue to the static variable.

My mistake. I didn't realize all_queues was a static member.

Please confirm that std::vector<std::optional<sycl::queue>> does not work

Comment to capture this discussion

core/src/SYCL/Kokkos_SYCL.cpp

masterleinad · 2021-02-02T15:36:12Z

The last commit provides a remedy for the inaccessibility of memory across multiple queues. We just need to make sure that all queues use the same sycl::context. Hence, I added a member function that returns the sycl::context associated to the default queue that can be used when initializing additional queues.

core/src/Kokkos_SYCL.hpp

dalg24 · 2021-02-04T15:00:49Z

core/src/SYCL/Kokkos_SYCL_Instance.cpp

@@ -55,6 +55,8 @@ namespace Impl {

 int SYCLInternal::was_finalized = 0;

+std::vector<std::optional<sycl::queue>*> SYCLInternal::all_queues;


Please confirm that std::vector<std::optional<sycl::queue>> does not work

core/src/SYCL/Kokkos_SYCL.cpp

masterleinad · 2021-02-04T15:35:43Z

Please confirm that std::vector<std::optionalsycl::queue> does not work

PI CUDA ERROR:
	Value:           4
	Name:            terminate called after throwing an instance of 'cl::sycl::runtime_error'
  what():  OpenCL API failed. OpenCL API returns: -999 (Unknown OpenCL error code) -999 (Unknown OpenCL error code)
Aborted (core dumped)

masterleinad · 2021-02-05T13:15:59Z

Retest this please.

crtrott

looks ok but need to protect the pushback into all_queues. Also all_queues will strictly grow forever. would maybe a set or soemthing where we can delete again from be better? Or maybe a counter on how many are empty and if more than 20% are empty the vector gets condensed and then shrunk? I know that sucks a bit because it means the impl_global_fence then also needs to use the lock.

crtrott · 2021-02-05T16:51:25Z

core/src/SYCL/Kokkos_SYCL_Instance.cpp

-    };
-    m_queue.emplace(d, exception_handler);
+    m_queue = q;
+    all_queues.push_back(&m_queue);


so we made the counter thread safe, i.e. we try to support multi threaded environments, but here we use push_back on a global? Lets protect this with a global mutex or so. Maybe have a static member function to insert into all_queues.

masterleinad · 2021-02-05T18:53:14Z

I guarded access to all_queues with a std::mutex now. Assuming that we traverse all_queues (using impl_static_fence()) much more often than pushing into it or removing elements, I left it as a std::vector.
Calling SYCLInternal::finalize now removes the respective sycl::queue from all_queues.

nliber · 2021-02-02T17:14:20Z

core/src/SYCL/Kokkos_SYCL_Instance.cpp

@@ -55,6 +55,8 @@ namespace Impl {

 int SYCLInternal::was_finalized = 0;

+std::vector<std::optional<sycl::queue>*> SYCLInternal::all_queues;


My mistake. I didn't realize all_queues was a static member.

core/src/Kokkos_SYCL.hpp

core/src/SYCL/Kokkos_SYCL.cpp

core/src/Kokkos_SYCL.hpp

core/src/SYCL/Kokkos_SYCL_Instance.cpp

core/src/SYCL/Kokkos_SYCL.cpp

core/src/Kokkos_SYCL.hpp

core/src/SYCL/Kokkos_SYCL_Parallel_Reduce.hpp

masterleinad · 2021-02-09T17:24:49Z

@nliber I believe I addressed all your comments.

masterleinad · 2021-02-09T18:39:55Z

Retest this please.

core/src/SYCL/Kokkos_SYCL.cpp

core/src/SYCL/Kokkos_SYCL_Instance.hpp

core/src/SYCL/Kokkos_SYCL_Parallel_Reduce.hpp

core/unit_test/sycl/TestSYCL_InterOp_Streams.cpp

dalg24 · 2021-02-10T15:38:18Z

core/src/SYCL/Kokkos_SYCL_Instance.cpp

@@ -55,6 +55,8 @@ namespace Impl {

 int SYCLInternal::was_finalized = 0;

+std::vector<std::optional<sycl::queue>*> SYCLInternal::all_queues;


Comment to capture this discussion

core/src/SYCL/Kokkos_SYCL_Instance.cpp

masterleinad · 2021-02-10T18:12:03Z

@dalg24 I think I addressed all your comments.

dalg24

Looks good

dalg24 · 2021-02-10T21:01:33Z

core/src/SYCL/Kokkos_SYCL_Instance.hpp

@@ -75,6 +75,12 @@ class SYCLInternal {

  std::optional<sycl::queue> m_queue;

+  // Using std::vector<std::optional<sycl::queue>> reveals a compiler bug when
+  // compiling for the CUDA backend. Storing pointers instead works around this.


Where did we report this bug? @nliber

I don't know that we have (where do SYCL+CUDA bugs get reported)?

core/unit_test/sycl/TestSYCL_InterOp_Streams.cpp

core/unit_test/sycl/TestSYCL_InterOp_Init_Context.cpp

core/unit_test/sycl/TestSYCL_InterOp_Streams.cpp

masterleinad · 2021-02-10T22:46:56Z

Retest this please.

dalg24 · 2021-02-11T22:13:43Z

Plz cleanup history

masterleinad · 2021-02-11T22:51:38Z

Plz cleanup history

@dalg24 Done.

dalg24 · 2021-02-12T02:25:41Z

Retest this please

masterleinad requested a review from nliber February 1, 2021 17:17

nliber requested changes Feb 1, 2021

View reviewed changes

masterleinad force-pushed the sycl_static_fence branch from d9ce787 to 5b72390 Compare February 2, 2021 15:37

nliber approved these changes Feb 2, 2021

View reviewed changes

masterleinad added the Blocks Promotion Overview issue for release-blocking bugs label Feb 3, 2021

masterleinad force-pushed the sycl_static_fence branch from 35feba8 to ae70fdb Compare February 4, 2021 04:54

dalg24 reviewed Feb 4, 2021

View reviewed changes

masterleinad force-pushed the sycl_static_fence branch from a75fab3 to fd0faec Compare February 4, 2021 23:34

dalg24 requested a review from nliber February 5, 2021 14:57

crtrott requested changes Feb 5, 2021

View reviewed changes

crtrott approved these changes Feb 5, 2021

View reviewed changes

nliber requested changes Feb 8, 2021

View reviewed changes

masterleinad force-pushed the sycl_static_fence branch from 7fc8d01 to 4fd5d82 Compare February 9, 2021 17:13

masterleinad requested a review from nliber February 9, 2021 17:24

masterleinad force-pushed the sycl_static_fence branch 2 times, most recently from 8c7376f to e377894 Compare February 10, 2021 13:55

dalg24 reviewed Feb 10, 2021

View reviewed changes

masterleinad force-pushed the sycl_static_fence branch from e61205b to 0c8fb50 Compare February 10, 2021 16:42

dalg24 mentioned this pull request Feb 10, 2021

Smart pointer with shared object ownership compatible with host and device #3791

Closed

dalg24 reviewed Feb 10, 2021

View reviewed changes

Rombur reviewed Feb 10, 2021

View reviewed changes

core/unit_test/sycl/TestSYCL_InterOp_Init_Context.cpp Show resolved Hide resolved

core/unit_test/sycl/TestSYCL_InterOp_Streams.cpp Show resolved Hide resolved

nliber approved these changes Feb 11, 2021

View reviewed changes

masterleinad added 6 commits February 11, 2021 17:36

Fix typo

1ec1c38

Avoid copying policies to SYCL kernels

ce21714

Allow initializing SYCL execution space instances with external queues

6fe9077

Implement SYCL::impl_static_fence

c27469b

Fix memory leaks in HIP and CUDA stream tests

187b059

Add tests for using multiple SYCL queues

65d270e

masterleinad force-pushed the sycl_static_fence branch from 2d4f340 to 65d270e Compare February 11, 2021 22:51

dalg24 merged commit efd8560 into kokkos:develop Feb 12, 2021

masterleinad deleted the sycl_static_fence branch February 12, 2021 13:49

masterleinad mentioned this pull request Mar 1, 2021

Memory leak in CUDA and HIP stream execution space instances #3817

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Allow initializing SYCL execution space from sycl::queue and SYCL::impl_static_fence #3767

Allow initializing SYCL execution space from sycl::queue and SYCL::impl_static_fence #3767

masterleinad commented Feb 1, 2021

nliber Feb 1, 2021 •

edited

masterleinad Feb 1, 2021

masterleinad Feb 2, 2021

nliber Feb 2, 2021

dalg24 Feb 4, 2021

dalg24 Feb 10, 2021

masterleinad commented Feb 2, 2021

dalg24 Feb 4, 2021

masterleinad commented Feb 4, 2021

masterleinad commented Feb 5, 2021

crtrott left a comment

crtrott Feb 5, 2021

masterleinad commented Feb 5, 2021

nliber Feb 2, 2021

masterleinad commented Feb 9, 2021

masterleinad commented Feb 9, 2021

dalg24 Feb 10, 2021

masterleinad commented Feb 10, 2021

dalg24 left a comment

dalg24 Feb 10, 2021

nliber Feb 11, 2021

masterleinad commented Feb 10, 2021

dalg24 commented Feb 11, 2021

masterleinad commented Feb 11, 2021

dalg24 commented Feb 12, 2021

		@@ -55,6 +55,8 @@ namespace Impl {

		int SYCLInternal::was_finalized = 0;

		std::vector<std::optional<sycl::queue>*> SYCLInternal::all_queues;

Allow initializing SYCL execution space from sycl::queue and SYCL::impl_static_fence #3767

Allow initializing SYCL execution space from sycl::queue and SYCL::impl_static_fence #3767

Conversation

masterleinad commented Feb 1, 2021

nliber Feb 1, 2021 • edited

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

masterleinad commented Feb 2, 2021

Choose a reason for hiding this comment

masterleinad commented Feb 4, 2021

masterleinad commented Feb 5, 2021

crtrott left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

masterleinad commented Feb 5, 2021

Choose a reason for hiding this comment

masterleinad commented Feb 9, 2021

masterleinad commented Feb 9, 2021

Choose a reason for hiding this comment

masterleinad commented Feb 10, 2021

dalg24 left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

masterleinad commented Feb 10, 2021

dalg24 commented Feb 11, 2021

masterleinad commented Feb 11, 2021

dalg24 commented Feb 12, 2021

nliber Feb 1, 2021 •

edited