From e72b85cdf2c336b92cf6e9d8e16c54d8630916de Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Alberto=20Cabrera=20P=C3=A9rez?= Date: Tue, 27 Feb 2024 15:58:38 +0000 Subject: [PATCH] [SYCL][COMPAT] Add set_default_queue functionality (#12835) This PR adds mechanisms to change the default queue of the current device both via the device extension and a free function. --------- Signed-off-by: Alberto Cabrera --- sycl/doc/syclcompat/README.md | 20 +++ sycl/include/syclcompat/device.hpp | 20 +++ sycl/test-e2e/syclcompat/device/device.cpp | 11 ++ .../memory/memory_management_diff_queues.cpp | 128 ++++++++++++++++++ 4 files changed, 179 insertions(+) create mode 100644 sycl/test-e2e/syclcompat/memory/memory_management_diff_queues.cpp diff --git a/sycl/doc/syclcompat/README.md b/sycl/doc/syclcompat/README.md index bafbb8efa4d11..d5f36a689fde3 100644 --- a/sycl/doc/syclcompat/README.md +++ b/sycl/doc/syclcompat/README.md @@ -742,6 +742,14 @@ sycl::queue create_queue(bool print_on_async_exceptions = false, // device manager. sycl::queue get_default_queue(); +// Util function to set the default queue of the current device in the +// device manager. +// If the device extension saved queue is the default queue, +// the previous saved queue will be overwritten as well. +// This function will be blocking if there are submitted kernels in the +// previous default queue. +void set_default_queue(const sycl::queue &q); + // Util function to wait for the queued kernels. void wait(sycl::queue q = get_default_queue()); @@ -785,6 +793,17 @@ destructor waits on a set of `sycl::event` which can be added to via schedule release of memory after a kernel or `mempcy`. SYCL device properties can be queried through `device_ext` as well. +Users can manage queues through the `syclcompat::set_default_queue(sycl::queue q)` +free function, and the `device_ext` `set_saved_queue`, `set_default_queue`, +and `get_saved_queue` member functions. +`set_default_queue` is blocking, and +overwrites the previous default queue with a user defined one, waiting for any +submitted kernels to finish. +The `device_ext` automatically sets the saved queue to the default queue. +Therefore, it's important to note that if the previous default queue was the +device's saved queue, setting a new default queue will update the reference of +the saved queue to the new default one to keep the state of the class consistent. + The class is exposed as follows: ```c++ @@ -807,6 +826,7 @@ class device_ext : public sycl::device { void reset(); sycl::queue *default_queue(); + void set_default_queue(const sycl::queue &q); void queues_wait_and_throw(); sycl::queue *create_queue(bool print_on_async_exceptions = false, bool in_order = true); diff --git a/sycl/include/syclcompat/device.hpp b/sycl/include/syclcompat/device.hpp index be02a42aff4cf..0a30debf9032f 100644 --- a/sycl/include/syclcompat/device.hpp +++ b/sycl/include/syclcompat/device.hpp @@ -309,6 +309,15 @@ class device_ext : public sycl::device { _saved_queue = _default_queue = _queues.front().get(); } + void set_default_queue(const sycl::queue &q) { + std::lock_guard lock(m_mutex); + _queues.front().get()->wait_and_throw(); + _queues[0] = std::make_shared(q); + if (_saved_queue == _default_queue) + _saved_queue = _queues.front().get(); + _default_queue = _queues.front().get(); + } + sycl::queue *default_queue() { return _default_queue; } void queues_wait_and_throw() { @@ -517,6 +526,17 @@ static inline sycl::queue get_default_queue() { return *detail::dev_mgr::instance().current_device().default_queue(); } +/// Util function to change the default queue of the current device in the +/// device manager +/// If the device extension saved queue is the default queue, +/// the previous saved queue will be overwritten as well. +/// This function will be blocking if there are submitted kernels in the +/// previous default queue. +/// @param q New user-defined queue +static inline void set_default_queue(const sycl::queue &q) { + detail::dev_mgr::instance().current_device().set_default_queue(q); +} + inline void wait(sycl::queue q = get_default_queue()) { q.wait(); } /// Util function to get the id of current device in diff --git a/sycl/test-e2e/syclcompat/device/device.cpp b/sycl/test-e2e/syclcompat/device/device.cpp index c807340301586..c12052fd2045e 100644 --- a/sycl/test-e2e/syclcompat/device/device.cpp +++ b/sycl/test-e2e/syclcompat/device/device.cpp @@ -36,6 +36,17 @@ #include "device_fixt.hpp" +void test_set_default_queue() { + std::cout << __PRETTY_FUNCTION__ << std::endl; + + DeviceExtFixt dev_ext; + auto &dev_ = dev_ext.get_dev_ext(); + sycl::queue old_default_queue = syclcompat::get_default_queue(); + dev_.set_default_queue(syclcompat::create_queue()); + assert(*dev_.default_queue() == *dev_.get_saved_queue()); + assert(*dev_.default_queue() != old_default_queue); +} + int main() { /* Device Tests diff --git a/sycl/test-e2e/syclcompat/memory/memory_management_diff_queues.cpp b/sycl/test-e2e/syclcompat/memory/memory_management_diff_queues.cpp new file mode 100644 index 0000000000000..008cdb14ec36c --- /dev/null +++ b/sycl/test-e2e/syclcompat/memory/memory_management_diff_queues.cpp @@ -0,0 +1,128 @@ +/*************************************************************************** + * + * Copyright (C) Codeplay Software Ltd. + * + * Part of the LLVM Project, under the Apache License v2.0 with LLVM + * Exceptions. See https://llvm.org/LICENSE.txt for license information. + * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + * + * SYCLcompat API + * + * memory_management_diff_queue.cpp + * + * Description: + * memory operations tests for operations when changing the default queue + **************************************************************************/ + +// RUN: %clangxx -std=c++20 -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out +// RUN: %{run} %t.out + +#include + +#include + +#include "../common.hpp" +#include "memory_common.hpp" +#include "memory_fixt.hpp" + +void test_memcpy() { + std::cout << __PRETTY_FUNCTION__ << std::endl; + sycl::queue q{{sycl::property::queue::in_order()}}; + + constexpr int ELEMENTS = 5000; + constexpr int N1 = 1000; + float *h_A = (float *)malloc(ELEMENTS * sizeof(float)); + float *h_B = (float *)malloc(ELEMENTS * sizeof(float)); + float *h_C = (float *)malloc(ELEMENTS * sizeof(float)); + + for (int i = 0; i < ELEMENTS; i++) { + h_A[i] = 1.0f; + h_B[i] = 2.0f; + } + + float *d_A = nullptr; + // hostA[0..999] -> deviceA[0..999] + // hostB[0..3999] -> deviceA[1000..4999] + // deviceA[0..4999] -> hostC[0..4999] + d_A = (float *)syclcompat::malloc(ELEMENTS * sizeof(float), q); + syclcompat::memcpy((void *)d_A, (void *)h_A, N1 * sizeof(float), q); + + syclcompat::set_default_queue(q); + syclcompat::memcpy((void *)(d_A + N1), (void *)h_B, + (ELEMENTS - N1) * sizeof(float)); + + syclcompat::memcpy((void *)h_C, (void *)d_A, ELEMENTS * sizeof(float)); + + // verify + for (int i = 0; i < N1; i++) { + assert(h_A[i] == h_C[i]); + } + + for (int i = N1; i < ELEMENTS; i++) { + assert(h_B[i] == h_C[i]); + } + + free(h_A); + free(h_B); + free(h_C); + syclcompat::free((void *)d_A); +} + +void test_memset() { + std::cout << __PRETTY_FUNCTION__ << std::endl; + sycl::queue q{{sycl::property::queue::in_order()}}; + + constexpr int PORTION = 5; + constexpr int ELEMENTS = PORTION * 3; + + int *h_A = (int *)malloc(ELEMENTS * sizeof(int)); + for (int i = 0; i < ELEMENTS; i++) { + h_A[i] = 4; + } + + int *d_A = nullptr; + + d_A = (int *)syclcompat::malloc(ELEMENTS * sizeof(int)); + // hostA -> deviceA + syclcompat::memcpy((void *)d_A, (void *)h_A, ELEMENTS * sizeof(int), q); + + // set d_A[0,..., PORTION - 1] = 0 + syclcompat::memset((void *)d_A, 0, PORTION * sizeof(int), q); + + syclcompat::set_default_queue(q); + // set d_A[PORTION,..., 2 * PORTION - 1] = 0x01010101 + syclcompat::memset((void *)(d_A + PORTION), 1, PORTION * sizeof(int)); + // deviceA -> hostA + syclcompat::memcpy((void *)h_A, (void *)d_A, ELEMENTS * sizeof(int)); + + // check d_A[0,..., PORTION - 1] = 0 + for (int i = 0; i < PORTION; i++) { + assert(h_A[i] == 0); + } + + // check d_A[PORTION,..., 2 * PORTION - 1] = 0x01010101 + for (int i = PORTION; i < (2 * PORTION - 1); i++) { + assert(h_A[i] == 0x01010101); + } + + // check d_A[2 * PORTION,..., ELEMENTS] = 4 + for (int i = 2 * PORTION; i < ELEMENTS; i++) { + assert(h_A[i] == 4); + } + + free(h_A); + syclcompat::free((void *)d_A); +} + +int main() { + test_memcpy(); + test_memset(); + + return 0; +}