diff --git a/sycl/doc/syclcompat/README.md b/sycl/doc/syclcompat/README.md index 08196908e1277..06b5001e7c0f1 100644 --- a/sycl/doc/syclcompat/README.md +++ b/sycl/doc/syclcompat/README.md @@ -399,7 +399,7 @@ static void destroy_event(event_ptr event); } // syclcompat ``` -### Memory Allocation +### Memory Operations This library provides interfaces to allocate memory to be accessed within kernel functions and on the host. The `syclcompat::malloc` function allocates device @@ -510,6 +510,64 @@ public: } // syclcompat ``` +The `syclcompat::experimental` namespace contains currently unsupported `memcpy` overloads which take a `syclcompat::experimental::memcpy_parameter` argument. These are included for forwards compatibility and currently throw a `std::runtime_error`. + +```cpp +namespace syclcompat { +namespace experimental { +// Forward declarations for types relating to unsupported memcpy_parameter API: + +enum memcpy_direction { + host_to_host, + host_to_device, + device_to_host, + device_to_device, + automatic +}; + +#ifdef SYCL_EXT_ONEAPI_BINDLESS_IMAGES +class image_mem_wrapper; +#endif +class image_matrix; + +/// Memory copy parameters for 2D/3D memory data. +struct memcpy_parameter { + struct data_wrapper { + pitched_data pitched{}; + sycl::id<3> pos{}; +#ifdef SYCL_EXT_ONEAPI_BINDLESS_IMAGES + experimental::image_mem_wrapper *image_bindless{nullptr}; +#endif + image_matrix *image{nullptr}; + }; + data_wrapper from{}; + data_wrapper to{}; + sycl::range<3> size{}; + syclcompat::detail::memcpy_direction direction{syclcompat::detail::memcpy_direction::automatic}; +}; + +/// [UNSUPPORTED] Synchronously copies 2D/3D memory data specified by \p param . +/// The function will return after the copy is completed. +/// +/// \param param Memory copy parameters. +/// \param q Queue to execute the copy task. +/// \returns no return value. +static inline void memcpy(const memcpy_parameter ¶m, + sycl::queue q = get_default_queue()); + +/// [UNSUPPORTED] Asynchronously copies 2D/3D memory data specified by \p param +/// . The return of the function does NOT guarantee the copy is completed. +/// +/// \param param Memory copy parameters. +/// \param q Queue to execute the copy task. +/// \returns no return value. +static inline void memcpy_async(const memcpy_parameter ¶m, + sycl::queue q = get_default_queue()); + +} // namespace experimental +} // namespace syclcompat +``` + Finally, the class `pitched_data`, which manages memory allocation for 3D spaces, padded to avoid uncoalesced memory accesses. diff --git a/sycl/include/syclcompat/memory.hpp b/sycl/include/syclcompat/memory.hpp index a8f7e89c52ab6..ad33ce9a9bdf8 100644 --- a/sycl/include/syclcompat/memory.hpp +++ b/sycl/include/syclcompat/memory.hpp @@ -77,7 +77,7 @@ template auto *local_mem() { return As; } -namespace detail { +namespace experimental { enum memcpy_direction { host_to_host, host_to_device, @@ -85,7 +85,7 @@ enum memcpy_direction { device_to_device, automatic }; -} // namespace detail +} enum class memory_region { global = 0, // device global memory @@ -122,6 +122,42 @@ class pitched_data { size_t _pitch, _x, _y; }; +namespace experimental { +#ifdef SYCL_EXT_ONEAPI_BINDLESS_IMAGES +class image_mem_wrapper; +namespace detail { +static sycl::event memcpy(const image_mem_wrapper *src, + const sycl::id<3> &src_id, pitched_data &dest, + const sycl::id<3> &dest_id, + const sycl::range<3> ©_extend, sycl::queue q); +static sycl::event memcpy(const pitched_data src, const sycl::id<3> &src_id, + image_mem_wrapper *dest, const sycl::id<3> &dest_id, + const sycl::range<3> ©_extend, sycl::queue q); +} // namespace detail +#endif +class image_matrix; +namespace detail { +static pitched_data to_pitched_data(image_matrix *image); +} + +/// Memory copy parameters for 2D/3D memory data. +struct memcpy_parameter { + struct data_wrapper { + pitched_data pitched{}; + sycl::id<3> pos{}; +#ifdef SYCL_EXT_ONEAPI_BINDLESS_IMAGES + experimental::image_mem_wrapper *image_bindless{nullptr}; +#endif + image_matrix *image{nullptr}; + }; + data_wrapper from{}; + data_wrapper to{}; + sycl::range<3> size{}; + syclcompat::experimental::memcpy_direction direction{ + syclcompat::experimental::memcpy_direction::automatic}; +}; +} // namespace experimental + namespace detail { template class accessor; @@ -263,21 +299,16 @@ static pointer_access_attribute get_pointer_attribute(sycl::queue q, } } -static memcpy_direction deduce_memcpy_direction(sycl::queue q, void *to_ptr, - const void *from_ptr) { +static experimental::memcpy_direction +deduce_memcpy_direction(sycl::queue q, void *to_ptr, const void *from_ptr) { // table[to_attribute][from_attribute] + using namespace experimental; // for memcpy_direction static const memcpy_direction direction_table[static_cast(pointer_access_attribute::end)] [static_cast(pointer_access_attribute::end)] = { - {memcpy_direction::host_to_host, - memcpy_direction::device_to_host, - memcpy_direction::host_to_host}, - {memcpy_direction::host_to_device, - memcpy_direction::device_to_device, - memcpy_direction::device_to_device}, - {memcpy_direction::host_to_host, - memcpy_direction::device_to_device, - memcpy_direction::device_to_device}}; + {host_to_host, device_to_host, host_to_host}, + {host_to_device, device_to_device, device_to_device}, + {host_to_host, device_to_device, device_to_device}}; return direction_table[static_cast(get_pointer_attribute( q, to_ptr))][static_cast(get_pointer_attribute(q, from_ptr))]; } @@ -300,6 +331,28 @@ static inline size_t get_offset(sycl::id<3> id, size_t slice, size_t pitch) { return slice * id.get(2) + pitch * id.get(1) + id.get(0); } +// RAII for host pointer +class host_buffer { + void *_buf; + size_t _size; + sycl::queue _q; + const std::vector &_deps; // free operation depends + +public: + host_buffer(size_t size, sycl::queue q, const std::vector &deps) + : _buf(std::malloc(size)), _size(size), _q(q), _deps(deps) {} + void *get_ptr() const { return _buf; } + size_t get_size() const { return _size; } + ~host_buffer() { + if (_buf) { + _q.submit([&](sycl::handler &cgh) { + cgh.depends_on(_deps); + cgh.host_task([buf = _buf] { std::free(buf); }); + }); + } + } +}; + /// copy 3D matrix specified by \p size from 3D matrix specified by \p from_ptr /// and \p from_range to another specified by \p to_ptr and \p to_range. static inline std::vector @@ -307,28 +360,7 @@ memcpy(sycl::queue q, void *to_ptr, const void *from_ptr, sycl::range<3> to_range, sycl::range<3> from_range, sycl::id<3> to_id, sycl::id<3> from_id, sycl::range<3> size, const std::vector &dep_events = {}) { - // RAII for host pointer - class host_buffer { - void *_buf; - size_t _size; - sycl::queue _q; - const std::vector &_deps; // free operation depends - - public: - host_buffer(size_t size, sycl::queue q, - const std::vector &deps) - : _buf(std::malloc(size)), _size(size), _q(q), _deps(deps) {} - void *get_ptr() const { return _buf; } - size_t get_size() const { return _size; } - ~host_buffer() { - if (_buf) { - _q.submit([&](sycl::handler &cgh) { - cgh.depends_on(_deps); - cgh.host_task([buf = _buf] { std::free(buf); }); - }); - } - } - }; + std::vector event_list; size_t to_slice = to_range.get(1) * to_range.get(0); @@ -343,6 +375,7 @@ memcpy(sycl::queue q, void *to_ptr, const void *from_ptr, return {memcpy(q, to_surface, from_surface, to_slice * size.get(2), dep_events)}; } + using namespace experimental; // for memcpy_direction memcpy_direction direction = deduce_memcpy_direction(q, to_ptr, from_ptr); size_t size_slice = size.get(1) * size.get(0); switch (direction) { @@ -448,6 +481,56 @@ static sycl::event combine_events(std::vector &events, } // namespace detail +namespace experimental { +namespace detail { +static inline std::vector +memcpy(sycl::queue q, const experimental::memcpy_parameter ¶m) { + auto to = param.to.pitched; + auto from = param.from.pitched; +#ifdef SYCL_EXT_ONEAPI_BINDLESS_IMAGES + if (param.to.image_bindless != nullptr && + param.from.image_bindless != nullptr) { + throw std::runtime_error( + "[SYCLcompat] memcpy: Unsupported bindless_image API."); + // TODO: Need change logic when sycl support image_mem to image_mem copy. + std::vector event_list; + syclcompat::detail::host_buffer buf(param.size.size(), q, event_list); + to.set_data_ptr(buf.get_ptr()); + experimental::detail::memcpy(param.from.image_bindless, param.from.pos, to, + sycl::id<3>(0, 0, 0), param.size, q); + from.set_data_ptr(buf.get_ptr()); + event_list.push_back(experimental::detail::memcpy( + from, sycl::id<3>(0, 0, 0), param.to.image_bindless, param.to.pos, + param.size, q)); + return event_list; + } else if (param.to.image_bindless != nullptr) { + throw std::runtime_error( + "[SYCLcompat] memcpy: Unsupported bindless_image API."); + return {experimental::detail::memcpy(from, param.from.pos, + param.to.image_bindless, param.to.pos, + param.size, q)}; + } else if (param.from.image_bindless != nullptr) { + throw std::runtime_error( + "[SYCLcompat] memcpy: Unsupported bindless_image API."); + return {experimental::detail::memcpy(param.from.image_bindless, + param.from.pos, to, param.to.pos, + param.size, q)}; + } +#endif + if (param.to.image != nullptr) { + throw std::runtime_error("[SYCLcompat] memcpy: Unsupported image API."); + to = experimental::detail::to_pitched_data(param.to.image); + } + if (param.from.image != nullptr) { + throw std::runtime_error("[SYCLcompat] memcpy: Unsupported image API."); + from = experimental::detail::to_pitched_data(param.from.image); + } + return syclcompat::detail::memcpy(q, to, param.to.pos, from, param.from.pos, + param.size); +} +} // namespace detail +} // namespace experimental + /// Allocate memory block on the device. /// \param num_bytes Number of bytes to allocate. /// \param q Queue to execute the allocate task. @@ -757,6 +840,31 @@ static sycl::event inline fill_async(void *dev_ptr, const T &pattern, return detail::fill(q, dev_ptr, pattern, count); } +namespace experimental { + +/// [UNSUPPORTED] Synchronously copies 2D/3D memory data specified by \p param . +/// The function will return after the copy is completed. +/// +/// \param param Memory copy parameters. +/// \param q Queue to execute the copy task. +/// \returns no return value. +static inline void memcpy(const memcpy_parameter ¶m, + sycl::queue q = get_default_queue()) { + sycl::event::wait(syclcompat::experimental::detail::memcpy(q, param)); +} + +/// [UNSUPPORTED] Asynchronously copies 2D/3D memory data specified by \p param +/// . The return of the function does NOT guarantee the copy is completed. +/// +/// \param param Memory copy parameters. +/// \param q Queue to execute the copy task. +/// \returns no return value. +static inline void memcpy_async(const memcpy_parameter ¶m, + sycl::queue q = get_default_queue()) { + syclcompat::experimental::detail::memcpy(q, param); +} +} // namespace experimental + /// Synchronously sets \p value to the first \p size bytes starting from \p /// dev_ptr. The function will return after the memset operation is completed. /// diff --git a/sycl/test-e2e/syclcompat/memory/memory_image.cpp b/sycl/test-e2e/syclcompat/memory/memory_image.cpp new file mode 100644 index 0000000000000..5dc2ac8d5ed8a --- /dev/null +++ b/sycl/test-e2e/syclcompat/memory/memory_image.cpp @@ -0,0 +1,257 @@ +/*************************************************************************** + * + * 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_image.cpp + * + * Description: + * 3D memory copy tests for new image/memcpy_parameter API + **************************************************************************/ + +// RUN: %clangxx -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out +// RUN: %{run} %t.out + +// FIXME(@intel/syclcompat-lib-reviewers): These are some limited tests for the +// new syclcompat::experimental::memcpy API. These aren't officially supported +// at present, but we can test the pitched_data variants easily. Once this +// moves out of experimental, let's test these APIs thoroughly + +#include +#include +#include + +#include + +#include "memory_common.hpp" + +void test_memcpy3D_parameter_offset() { + std::cout << __PRETTY_FUNCTION__ << std::endl; + size_t width = 4; + size_t height = 4; + size_t depth = 5; + float *h_data; + + syclcompat::pitched_data d_data; + sycl::range<3> extent = sycl::range<3>(sizeof(float) * 1, 1, 1); + syclcompat::pitched_data cpyParm_from_data_ct1, cpyParm_to_data_ct1; + sycl::id<3> cpyParm_from_pos_ct1(0, 0, 0), cpyParm_to_pos_ct1(0, 0, 0); + sycl::range<3> cpyParm_size_ct1(0, 0, 0); + + h_data = + (float *)syclcompat::malloc_host(sizeof(float) * width * height * depth); + /* + 0.000000 1.000000 2.000000 3.000000 + 4.000000 5.000000 6.000000 7.000000 + 8.000000 9.000000 10.000000 11.000000 + 12.000000 13.000000 14.000000 15.000000 + + 16.000000 17.000000 18.000000 19.000000 + 20.000000 21.000000 22.000000 23.000000 + 24.000000 25.000000 26.000000 27.000000 + 28.000000 29.000000 30.000000 31.000000 + + 32.000000 33.000000 34.000000 35.000000 + 36.000000 37.000000 38.000000 39.000000 + 40.000000 41.000000 42.000000 43.000000 + 44.000000 45.000000 46.000000 47.000000 + + 48.000000 49.000000 50.000000 51.000000 + 52.000000 53.000000 54.000000 55.000000 + 56.000000 57.000000 58.000000 59.000000 + 60.000000 61.000000 62.000000 63.000000 + + 64.000000 65.000000 66.000000 67.000000 + 68.000000 69.000000 70.000000 71.000000 + 72.000000 73.000000 74.000000 75.000000 + 76.000000 77.000000 78.000000 79.000000 + */ + for (int i = 0; i < width * height * depth; i++) + h_data[i] = (float)i; + + /* + 5.000000 6.000000 + 9.000000 10.000000 + + 21.000000 22.000000 + 25.000000 26.000000 + + 37.000000 38.000000 + 41.000000 42.000000 + */ + float Ref[12] = {5, 6, 9, 10, 21, 22, 25, 26, 37, 38, 41, 42}; + + size_t out_width = 2; + size_t out_height = 2; + size_t out_depth = 3; + + // alloc memory. + extent = sycl::range<3>(sizeof(float) * width, height, depth); + d_data = (syclcompat::pitched_data)syclcompat::malloc(extent); + + // copy to Device. + cpyParm_from_data_ct1 = syclcompat::pitched_data( + (void *)h_data, sizeof(float) * width, width, height); + cpyParm_to_data_ct1 = d_data; + cpyParm_size_ct1 = extent; + + { + syclcompat::experimental::memcpy_parameter params{}; + params.to.pitched = cpyParm_to_data_ct1; + params.to.pos = cpyParm_to_pos_ct1; + params.from.pitched = cpyParm_from_data_ct1; + params.from.pos = cpyParm_from_pos_ct1; + params.size = cpyParm_size_ct1; + syclcompat::experimental::memcpy(params); + } + + cpyParm_from_pos_ct1 = {1 * sizeof(float), 1, 0}; // set offset on x/y/z. + cpyParm_size_ct1 = {out_width * sizeof(float), out_height, out_depth}; + + for (int i = 0; i < out_width * out_height * out_depth; i++) + h_data[i] = -1; + // copy back to host. + cpyParm_from_data_ct1 = d_data; + cpyParm_to_data_ct1 = syclcompat::pitched_data( + (void *)h_data, sizeof(float) * out_width, out_width, out_height); + + { + syclcompat::experimental::memcpy_parameter params{}; + params.to.pitched = cpyParm_to_data_ct1; + params.to.pos = cpyParm_to_pos_ct1; + params.from.pitched = cpyParm_from_data_ct1; + params.from.pos = cpyParm_from_pos_ct1; + params.size = cpyParm_size_ct1; + syclcompat::experimental::memcpy(params); + } + + // Copy back to host data. + check(h_data, Ref, out_width * out_height * out_depth); + syclcompat::free(h_data); + sycl::free(d_data.get_data_ptr(), syclcompat::get_default_context()); +} + +void test_memcpy3D_async_parameter_offset() { + std::cout << __PRETTY_FUNCTION__ << std::endl; + + size_t width = 4; + size_t height = 4; + size_t depth = 5; + float *h_data; + + syclcompat::pitched_data d_data; + sycl::range<3> extent = sycl::range<3>(sizeof(float) * 1, 1, 1); + syclcompat::pitched_data cpyParm_from_data_ct1, cpyParm_to_data_ct1; + sycl::id<3> cpyParm_from_pos_ct1(0, 0, 0), cpyParm_to_pos_ct1(0, 0, 0); + sycl::range<3> cpyParm_size_ct1(0, 0, 0); + + h_data = + (float *)syclcompat::malloc_host(sizeof(float) * width * height * depth); + /* + 0.000000 1.000000 2.000000 3.000000 + 4.000000 5.000000 6.000000 7.000000 + 8.000000 9.000000 10.000000 11.000000 + 12.000000 13.000000 14.000000 15.000000 + + 16.000000 17.000000 18.000000 19.000000 + 20.000000 21.000000 22.000000 23.000000 + 24.000000 25.000000 26.000000 27.000000 + 28.000000 29.000000 30.000000 31.000000 + + 32.000000 33.000000 34.000000 35.000000 + 36.000000 37.000000 38.000000 39.000000 + 40.000000 41.000000 42.000000 43.000000 + 44.000000 45.000000 46.000000 47.000000 + + 48.000000 49.000000 50.000000 51.000000 + 52.000000 53.000000 54.000000 55.000000 + 56.000000 57.000000 58.000000 59.000000 + 60.000000 61.000000 62.000000 63.000000 + + 64.000000 65.000000 66.000000 67.000000 + 68.000000 69.000000 70.000000 71.000000 + 72.000000 73.000000 74.000000 75.000000 + 76.000000 77.000000 78.000000 79.000000 + */ + for (int i = 0; i < width * height * depth; i++) + h_data[i] = (float)i; + + /* + 5.000000 6.000000 + 9.000000 10.000000 + + 21.000000 22.000000 + 25.000000 26.000000 + + 37.000000 38.000000 + 41.000000 42.000000 + */ + float Ref[12] = {5, 6, 9, 10, 21, 22, 25, 26, 37, 38, 41, 42}; + + size_t out_width = 2; + size_t out_height = 2; + size_t out_depth = 3; + + // alloc memory. + extent = sycl::range<3>(sizeof(float) * width, height, depth); + // test_feature:malloc + d_data = (syclcompat::pitched_data)syclcompat::malloc(extent); + + // copy to Device. + cpyParm_from_data_ct1 = syclcompat::pitched_data( + (void *)h_data, sizeof(float) * width, width, height); + cpyParm_to_data_ct1 = d_data; + cpyParm_size_ct1 = extent; + + { + syclcompat::experimental::memcpy_parameter params{}; + params.to.pitched = cpyParm_to_data_ct1; + params.to.pos = cpyParm_to_pos_ct1; + params.from.pitched = cpyParm_from_data_ct1; + params.from.pos = cpyParm_from_pos_ct1; + params.size = cpyParm_size_ct1; + syclcompat::experimental::memcpy_async(params); + } + syclcompat::get_default_queue().wait_and_throw(); + cpyParm_from_pos_ct1 = {1 * sizeof(float), 1, 0}; // set offset on x/y/z. + cpyParm_size_ct1 = {out_width * sizeof(float), out_height, out_depth}; + + for (int i = 0; i < out_width * out_height * out_depth; i++) + h_data[i] = -1; + // copy back to host. + cpyParm_from_data_ct1 = d_data; + cpyParm_to_data_ct1 = syclcompat::pitched_data( + (void *)h_data, sizeof(float) * out_width, out_width, out_height); + { + syclcompat::experimental::memcpy_parameter params{}; + params.to.pitched = cpyParm_to_data_ct1; + params.to.pos = cpyParm_to_pos_ct1; + params.from.pitched = cpyParm_from_data_ct1; + params.from.pos = cpyParm_from_pos_ct1; + params.size = cpyParm_size_ct1; + syclcompat::experimental::memcpy_async(params); + } + syclcompat::get_default_queue().wait_and_throw(); + // Copy back to host data. + check(h_data, Ref, out_width * out_height * out_depth); + syclcompat::free(h_data); + sycl::free(d_data.get_data_ptr(), syclcompat::get_default_context()); +} + +int main() { + // Copied and modified from memcpy_3d.cpp test_memcpy3D_offset() + test_memcpy3D_parameter_offset(); + test_memcpy3D_async_parameter_offset(); + return 0; +} diff --git a/sycl/test-e2e/syclcompat/memory/memory_image_xfails.cpp b/sycl/test-e2e/syclcompat/memory/memory_image_xfails.cpp new file mode 100644 index 0000000000000..a011e1da9e407 --- /dev/null +++ b/sycl/test-e2e/syclcompat/memory/memory_image_xfails.cpp @@ -0,0 +1,142 @@ +/*************************************************************************** + * + * 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_async.cpp + * + * Description: + * Asynchronous memory operations event dependency tests + **************************************************************************/ + +// The original source was under the license below: +// ====------ memory_async.cpp------------------- -*- C++ -* ----===// +// +// 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 +// +// +// ===---------------------------------------------------------------------===// + +// RUN: %clangxx -std=c++20 -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out +// RUN: %{run} %t.out + +// Tests for the sycl::events returned from syclcompat::*Async API calls + +#include "sycl/exception.hpp" +#include +#include + +#include + +#include + +void test_memcpy_parameter_async( + syclcompat::experimental::memcpy_parameter param, bool xpass) { + std::cout << __PRETTY_FUNCTION__ << std::endl; + + try { + syclcompat::experimental::memcpy_async(param); + assert(xpass); + } catch (std::runtime_error &) { + assert(!xpass); + } +} + +void test_memcpy_parameter(syclcompat::experimental::memcpy_parameter param, + bool xpass) { + std::cout << __PRETTY_FUNCTION__ << std::endl; + try { + syclcompat::experimental::memcpy(param); + assert(xpass); + } catch (std::runtime_error &) { + assert(!xpass); + } +} + +// Check (most) memcpy_parameter APIs raise std::runtime_error. +void test_memcpy_parameter_xfails() { + + { + // Empty `memcpy_params` passes in no bindless_image + // or image pointers. This is the code path that ought to pass. + syclcompat::experimental::memcpy_parameter params; + test_memcpy_parameter(params, true); + test_memcpy_parameter_async(params, true); + } + + { + // Mimick passing a bindless image for source + syclcompat::experimental::memcpy_parameter params; + params.from.image_bindless = + reinterpret_cast(1); + test_memcpy_parameter(params, false); + test_memcpy_parameter_async(params, false); + } + + { + // Mimick passing a bindless image for dest + syclcompat::experimental::memcpy_parameter params; + params.to.image_bindless = + reinterpret_cast(1); + test_memcpy_parameter(params, false); + test_memcpy_parameter_async(params, false); + } + + { + // Mimick passing a bindless image for source & dest + syclcompat::experimental::memcpy_parameter params; + params.from.image_bindless = + reinterpret_cast(1); + params.to.image_bindless = + reinterpret_cast(1); + test_memcpy_parameter(params, false); + test_memcpy_parameter_async(params, false); + } + + { + // Mimick passing an image for source + syclcompat::experimental::memcpy_parameter params; + params.from.image = + reinterpret_cast(1); + test_memcpy_parameter(params, false); + test_memcpy_parameter_async(params, false); + } + + { + // Mimick passing an image for dest + syclcompat::experimental::memcpy_parameter params; + params.to.image = + reinterpret_cast(1); + test_memcpy_parameter(params, false); + test_memcpy_parameter_async(params, false); + } + + { + // Mimick passing an image for source & dest + syclcompat::experimental::memcpy_parameter params; + params.from.image = + reinterpret_cast(1); + params.to.image = + reinterpret_cast(1); + test_memcpy_parameter(params, false); + test_memcpy_parameter_async(params, false); + } +} + +int main() { + test_memcpy_parameter_xfails(); + return 0; +} diff --git a/sycl/test-e2e/syclcompat/memory/usm_allocations.cpp b/sycl/test-e2e/syclcompat/memory/usm_allocations.cpp index 532163dda263a..2de67aeebb251 100644 --- a/sycl/test-e2e/syclcompat/memory/usm_allocations.cpp +++ b/sycl/test-e2e/syclcompat/memory/usm_allocations.cpp @@ -88,7 +88,7 @@ void test_non_templated_host() { void test_deduce() { std::cout << __PRETTY_FUNCTION__ << std::endl; - using memcpy_direction = syclcompat::detail::memcpy_direction; + using namespace syclcompat::experimental; // for memcpy_direction auto default_queue = syclcompat::get_default_queue(); if (!default_queue.get_device().has(sycl::aspect::usm_host_allocations)) return; // Skip unsupported diff --git a/sycl/test-e2e/syclcompat/memory/usm_shared_allocations.cpp b/sycl/test-e2e/syclcompat/memory/usm_shared_allocations.cpp index 6737614549863..22c45a1e874a4 100644 --- a/sycl/test-e2e/syclcompat/memory/usm_shared_allocations.cpp +++ b/sycl/test-e2e/syclcompat/memory/usm_shared_allocations.cpp @@ -63,7 +63,7 @@ void test_non_templated_shared() { void test_deduce_shared() { std::cout << __PRETTY_FUNCTION__ << std::endl; - using memcpy_direction = syclcompat::detail::memcpy_direction; + using namespace syclcompat::experimental; auto default_queue = syclcompat::get_default_queue(); int *h_ptr = (int *)syclcompat::malloc_host(sizeof(int));