diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index ad8e1cf4f2571..125944e6c10eb 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -1223,42 +1223,69 @@ class accessor : #if __cplusplus > 201402L -template -accessor(buffer, Ts...) +template +accessor(buffer) ->accessor; -template -accessor(buffer, handler, Ts...) +template +accessor(buffer, Type1) + ->accessor(), + detail::deduceAccessTarget(target::global_buffer), + access::placeholder::true_t>; + +template +accessor(buffer, Type1, Type2) + ->accessor(), + detail::deduceAccessTarget(target::global_buffer), + access::placeholder::true_t>; + +template +accessor(buffer, Type1, Type2, Type3) + ->accessor(), + detail::deduceAccessTarget(target::global_buffer), + access::placeholder::true_t>; + +template +accessor(buffer, Type1, Type2, Type3, Type4) + ->accessor(), + detail::deduceAccessTarget(target::global_buffer), + access::placeholder::true_t>; + +template +accessor(buffer, handler) ->accessor; -template -accessor(buffer, Ts..., mode_tag_t, - property_list = {}) - ->accessor; +template +accessor(buffer, handler, Type1) + ->accessor(), + detail::deduceAccessTarget(target::global_buffer), + access::placeholder::false_t>; -template -accessor(buffer, handler, Ts..., - mode_tag_t, property_list = {}) - ->accessor +accessor(buffer, handler, Type1, Type2) + ->accessor(), + detail::deduceAccessTarget(target::global_buffer), access::placeholder::false_t>; -template -accessor(buffer, Ts..., - mode_target_tag_t, property_list = {}) - ->accessor; +template +accessor(buffer, handler, Type1, Type2, Type3) + ->accessor(), + detail::deduceAccessTarget(target::global_buffer), + access::placeholder::false_t>; -template -accessor(buffer, handler, Ts..., - mode_target_tag_t, property_list = {}) - ->accessor +accessor(buffer, handler, Type1, Type2, Type3, + Type4) + ->accessor(), + detail::deduceAccessTarget(target::global_buffer), access::placeholder::false_t>; #endif @@ -1658,15 +1685,32 @@ class host_accessor #if __cplusplus > 201402L -template -host_accessor(buffer, Ts...) +template +host_accessor(buffer) ->host_accessor; -template -host_accessor(buffer, Ts..., - mode_tag_t, property_list = {}) - ->host_accessor; +template +host_accessor(buffer, Type1) + ->host_accessor()>; + +template +host_accessor(buffer, Type1, Type2) + ->host_accessor()>; + +template +host_accessor(buffer, Type1, Type2, Type3) + ->host_accessor()>; + +template +host_accessor(buffer, Type1, Type2, Type3, Type4) + ->host_accessor()>; #endif diff --git a/sycl/include/CL/sycl/buffer.hpp b/sycl/include/CL/sycl/buffer.hpp index 55eec92e637a5..98b4c217bac39 100644 --- a/sycl/include/CL/sycl/buffer.hpp +++ b/sycl/include/CL/sycl/buffer.hpp @@ -277,6 +277,23 @@ class buffer { accessOffset); } +#if __cplusplus > 201402L + + template auto get_access(Ts... args) { + return accessor{*this, args...}; + } + + template + auto get_access(handler &commandGroupHandler, Ts... args) { + return accessor{*this, commandGroupHandler, args...}; + } + + template auto get_host_access(Ts... args) { + return host_accessor{*this, args...}; + } + +#endif + template void set_final_data(Destination finalData = nullptr) { impl->set_final_data(finalData); diff --git a/sycl/include/CL/sycl/detail/accessor_impl.hpp b/sycl/include/CL/sycl/detail/accessor_impl.hpp index 83f842b622a19..5ef72ce5f57b9 100644 --- a/sycl/include/CL/sycl/detail/accessor_impl.hpp +++ b/sycl/include/CL/sycl/detail/accessor_impl.hpp @@ -183,6 +183,57 @@ using Requirement = AccessorImplHost; void __SYCL_EXPORT addHostAccessorAndWait(Requirement *Req); +#if __cplusplus > 201402L + +template +constexpr access::mode deduceAccessMode() { + // property_list = {} is not properly detected by deduction guide, + // when parameter is passed without curly braces: access(buffer, noinit) + // thus simplest approach is to check 2 last arguments for being a tag + if constexpr (std::is_same>::value || + std::is_same>::value) { + return access::mode::read; + } + + if constexpr (std::is_same>::value || + std::is_same>::value) { + return access::mode::write; + } + + if constexpr ( + std::is_same>::value || + std::is_same>::value) { + return access::mode::read; + } + + return access::mode::read_write; +} + +template +constexpr access::target deduceAccessTarget(access::target defaultTarget) { + if constexpr ( + std::is_same>::value || + std::is_same>::value) { + return access::target::constant_buffer; + } + + return defaultTarget; +} + +#endif + } // namespace detail } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/CL/sycl/detail/buffer_impl.hpp b/sycl/include/CL/sycl/detail/buffer_impl.hpp index 09783b9d3713f..d66692a5ea688 100644 --- a/sycl/include/CL/sycl/detail/buffer_impl.hpp +++ b/sycl/include/CL/sycl/detail/buffer_impl.hpp @@ -31,6 +31,8 @@ template class buffer; +template +class host_accessor; using buffer_allocator = detail::sycl_memory_object_allocator; diff --git a/sycl/test/basic_tests/accessor/Inputs/device_accessor.cpp b/sycl/test/basic_tests/accessor/Inputs/device_accessor.cpp new file mode 100755 index 0000000000000..32bb9b7a67b1f --- /dev/null +++ b/sycl/test/basic_tests/accessor/Inputs/device_accessor.cpp @@ -0,0 +1,315 @@ +//==---------- device_accessor.cpp - SYCL accessor basic test --------------==// +// +// 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 +// +//===----------------------------------------------------------------------===// +#include +#include + +int main() { + // Non-placeholder accessors. + { + int data[9] = {1, 2, 3, 4, 5, 6, 7, 8, 9}; + + sycl::buffer buf_data(data, sycl::range<1>(9), + {cl::sycl::property::buffer::use_host_ptr()}); + + sycl::queue Queue; + + Queue.submit([&](sycl::handler &cgh) { + +#if defined(accessor_new_api_test) + sycl::accessor acc_1(buf_data, cgh); + sycl::accessor acc_2(buf_data, cgh, sycl::range<1>(8)); + sycl::accessor acc_3(buf_data, cgh, sycl::range<1>(8), sycl::id<1>(1)); + sycl::accessor acc_4(buf_data, cgh, sycl::read_only); + sycl::accessor acc_5(buf_data, cgh, sycl::range<1>(8), sycl::read_only); + sycl::accessor acc_6(buf_data, cgh, sycl::range<1>(8), sycl::id<1>(1), + sycl::read_only); + sycl::accessor acc_7(buf_data, cgh, sycl::write_only); + sycl::accessor acc_8(buf_data, cgh, sycl::range<1>(8), sycl::write_only); + sycl::accessor acc_9(buf_data, cgh, sycl::range<1>(8), sycl::id<1>(1), + sycl::write_only); +#elif defined(buffer_new_api_test) + auto acc_1 = buf_data.get_access(cgh); + auto acc_2 = buf_data.get_access(cgh, sycl::range<1>(8)); + auto acc_3 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::id<1>(1)); + auto acc_4 = buf_data.get_access(cgh, sycl::read_only); + auto acc_5 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::read_only); + auto acc_6 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::id<1>(1), + sycl::read_only); + auto acc_7 = buf_data.get_access(cgh, sycl::write_only); + auto acc_8 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::write_only); + auto acc_9 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::id<1>(1), + sycl::write_only); +#endif + + assert(!acc_1.is_placeholder()); + assert(!acc_2.is_placeholder()); + assert(!acc_3.is_placeholder()); + assert(!acc_4.is_placeholder()); + assert(!acc_5.is_placeholder()); + assert(!acc_6.is_placeholder()); + assert(!acc_7.is_placeholder()); + assert(!acc_8.is_placeholder()); + assert(!acc_9.is_placeholder()); + + cgh.single_task( + [=]() { + acc_7[6] = acc_1[0]; + acc_8[7] = acc_2[1]; + acc_9[7] = acc_3[1]; + acc_1[0] = acc_4[3]; + acc_2[1] = acc_5[4]; + acc_3[1] = acc_6[4]; + }); + }); + Queue.wait(); + +#if defined(accessor_new_api_test) + sycl::host_accessor host_acc(buf_data, sycl::read_only); +#elif defined(buffer_new_api_test) + auto host_acc = buf_data.get_host_access(sycl::read_only); +#endif + assert(host_acc[0] == 4 && host_acc[1] == 5 && host_acc[2] == 6); + assert(host_acc[3] == 4 && host_acc[4] == 5 && host_acc[5] == 6); + assert(host_acc[6] == 1 && host_acc[7] == 2 && host_acc[8] == 3); + } + + // Placeholder accessors. + { + int data[9] = {1, 2, 3, 4, 5, 6, 7, 8, 9}; + + sycl::buffer buf_data(data, sycl::range<1>(9), + {cl::sycl::property::buffer::use_host_ptr()}); + +#if defined(accessor_new_api_test) + sycl::accessor acc_1(buf_data); + sycl::accessor acc_2(buf_data, sycl::range<1>(8)); + sycl::accessor acc_3(buf_data, sycl::range<1>(8), sycl::id<1>(1)); + sycl::accessor acc_4(buf_data, sycl::read_only); + sycl::accessor acc_5(buf_data, sycl::range<1>(8), sycl::read_only); + sycl::accessor acc_6(buf_data, sycl::range<1>(8), sycl::id<1>(1), + sycl::read_only); + sycl::accessor acc_7(buf_data, sycl::write_only); + sycl::accessor acc_8(buf_data, sycl::range<1>(8), sycl::write_only); + sycl::accessor acc_9(buf_data, sycl::range<1>(8), sycl::id<1>(1), + sycl::write_only); +#elif defined(buffer_new_api_test) + auto acc_1 = buf_data.get_access(); + auto acc_2 = buf_data.get_access(sycl::range<1>(8)); + auto acc_3 = buf_data.get_access(sycl::range<1>(8), sycl::id<1>(1)); + auto acc_4 = buf_data.get_access(sycl::read_only); + auto acc_5 = buf_data.get_access(sycl::range<1>(8), sycl::read_only); + auto acc_6 = buf_data.get_access(sycl::range<1>(8), sycl::id<1>(1), + sycl::read_only); + auto acc_7 = buf_data.get_access(sycl::write_only); + auto acc_8 = buf_data.get_access(sycl::range<1>(8), sycl::write_only); + auto acc_9 = buf_data.get_access(sycl::range<1>(8), sycl::id<1>(1), + sycl::write_only); +#endif + + assert(acc_1.is_placeholder()); + assert(acc_2.is_placeholder()); + assert(acc_3.is_placeholder()); + assert(acc_4.is_placeholder()); + assert(acc_5.is_placeholder()); + assert(acc_6.is_placeholder()); + assert(acc_7.is_placeholder()); + assert(acc_8.is_placeholder()); + assert(acc_9.is_placeholder()); + + sycl::queue Queue; + + Queue.submit([&](sycl::handler &cgh) { + cgh.require(acc_1); + cgh.require(acc_2); + cgh.require(acc_3); + cgh.require(acc_4); + cgh.require(acc_5); + cgh.require(acc_6); + cgh.require(acc_7); + cgh.require(acc_8); + cgh.require(acc_9); + + cgh.single_task( + [=]() { + acc_7[6] = acc_1[0]; + acc_8[7] = acc_2[1]; + acc_9[7] = acc_3[1]; + acc_1[0] = acc_4[3]; + acc_2[1] = acc_5[4]; + acc_3[1] = acc_6[4]; + }); + }); + Queue.wait(); + +#if defined(accessor_new_api_test) + sycl::host_accessor host_acc(buf_data, sycl::read_only); +#elif defined(buffer_new_api_test) + auto host_acc = buf_data.get_host_access(sycl::read_only); +#endif + assert(host_acc[0] == 4 && host_acc[1] == 5 && host_acc[2] == 6); + assert(host_acc[3] == 4 && host_acc[4] == 5 && host_acc[5] == 6); + assert(host_acc[6] == 1 && host_acc[7] == 2 && host_acc[8] == 3); + } + + // Non-placeholder noinit and constant_buffer accessors. + { + int data[9] = {1, 2, 3, 4, 5, 6, 7, 8, 9}; + + sycl::buffer buf_data(data, sycl::range<1>(9), + {cl::sycl::property::buffer::use_host_ptr()}); + + sycl::queue Queue; + + Queue.submit([&](sycl::handler &cgh) { + +#if defined(accessor_new_api_test) + sycl::accessor acc_1(buf_data, cgh, sycl::noinit); + sycl::accessor acc_2(buf_data, cgh, sycl::range<1>(8), sycl::noinit); + sycl::accessor acc_3(buf_data, cgh, sycl::range<1>(8), sycl::id<1>(1), + sycl::noinit); + sycl::accessor acc_4(buf_data, cgh, sycl::read_constant); + sycl::accessor acc_5(buf_data, cgh, sycl::range<1>(8), + sycl::read_constant); + sycl::accessor acc_6(buf_data, cgh, sycl::range<1>(8), sycl::id<1>(1), + sycl::read_constant); + sycl::accessor acc_7(buf_data, cgh, sycl::write_only, sycl::noinit); + sycl::accessor acc_8(buf_data, cgh, sycl::range<1>(8), sycl::write_only, + sycl::noinit); + sycl::accessor acc_9(buf_data, cgh, sycl::range<1>(8), sycl::id<1>(1), + sycl::write_only, sycl::noinit); +#elif defined(buffer_new_api_test) + auto acc_1 = buf_data.get_access(cgh, sycl::noinit); + auto acc_2 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::noinit); + auto acc_3 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::id<1>(1), + sycl::noinit); + auto acc_4 = buf_data.get_access(cgh, sycl::read_constant); + auto acc_5 = buf_data.get_access(cgh, sycl::range<1>(8), + sycl::read_constant); + auto acc_6 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::id<1>(1), + sycl::read_constant); + auto acc_7 = buf_data.get_access(cgh, sycl::write_only, sycl::noinit); + auto acc_8 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::write_only, + sycl::noinit); + auto acc_9 = buf_data.get_access(cgh, sycl::range<1>(8), sycl::id<1>(1), + sycl::write_only, sycl::noinit); +#endif + + assert(!acc_1.is_placeholder()); + assert(!acc_2.is_placeholder()); + assert(!acc_3.is_placeholder()); + assert(!acc_4.is_placeholder()); + assert(!acc_5.is_placeholder()); + assert(!acc_6.is_placeholder()); + assert(!acc_7.is_placeholder()); + assert(!acc_8.is_placeholder()); + assert(!acc_9.is_placeholder()); + + cgh.single_task( + [=]() { + acc_7[6] = 1; + acc_8[7] = 2; + acc_9[7] = 3; + acc_1[0] = acc_4[3]; + acc_2[1] = acc_5[4]; + acc_3[1] = acc_6[4]; + }); + }); + Queue.wait(); + +#if defined(accessor_new_api_test) + sycl::host_accessor host_acc(buf_data, sycl::read_only); +#elif defined(buffer_new_api_test) + auto host_acc = buf_data.get_host_access(sycl::read_only); +#endif + assert(host_acc[0] == 4 && host_acc[1] == 5 && host_acc[2] == 6); + assert(host_acc[3] == 4 && host_acc[4] == 5 && host_acc[5] == 6); + assert(host_acc[6] == 1 && host_acc[7] == 2 && host_acc[8] == 3); + } + + // Placeholder noinit and constant_buffer accessors. + { + int data[9] = {1, 2, 3, 4, 5, 6, 7, 8, 9}; + + sycl::buffer buf_data(data, sycl::range<1>(9), + {cl::sycl::property::buffer::use_host_ptr()}); + +#if defined(accessor_new_api_test) + sycl::accessor acc_1(buf_data, sycl::noinit); + sycl::accessor acc_2(buf_data, sycl::range<1>(8), sycl::noinit); + sycl::accessor acc_3(buf_data, sycl::range<1>(8), sycl::id<1>(1), + sycl::noinit); + sycl::accessor acc_4(buf_data, sycl::read_constant); + sycl::accessor acc_5(buf_data, sycl::range<1>(8), sycl::read_constant); + sycl::accessor acc_6(buf_data, sycl::range<1>(8), sycl::id<1>(1), + sycl::read_constant); + sycl::accessor acc_7(buf_data, sycl::write_only, sycl::noinit); + sycl::accessor acc_8(buf_data, sycl::range<1>(8), sycl::write_only, + sycl::noinit); + sycl::accessor acc_9(buf_data, sycl::range<1>(8), sycl::id<1>(1), + sycl::write_only, sycl::noinit); +#elif defined(buffer_new_api_test) + auto acc_1 = buf_data.get_access(sycl::noinit); + auto acc_2 = buf_data.get_access(sycl::range<1>(8), sycl::noinit); + auto acc_3 = buf_data.get_access(sycl::range<1>(8), sycl::id<1>(1), + sycl::noinit); + auto acc_4 = buf_data.get_access(sycl::read_constant); + auto acc_5 = buf_data.get_access(sycl::range<1>(8), sycl::read_constant); + auto acc_6 = buf_data.get_access(sycl::range<1>(8), sycl::id<1>(1), + sycl::read_constant); + auto acc_7 = buf_data.get_access(sycl::write_only, sycl::noinit); + auto acc_8 = buf_data.get_access(sycl::range<1>(8), sycl::write_only, + sycl::noinit); + auto acc_9 = buf_data.get_access(sycl::range<1>(8), sycl::id<1>(1), + sycl::write_only, sycl::noinit); +#endif + + assert(acc_1.is_placeholder()); + assert(acc_2.is_placeholder()); + assert(acc_3.is_placeholder()); + assert(acc_4.is_placeholder()); + assert(acc_5.is_placeholder()); + assert(acc_6.is_placeholder()); + assert(acc_7.is_placeholder()); + assert(acc_8.is_placeholder()); + assert(acc_9.is_placeholder()); + + sycl::queue Queue; + + Queue.submit([&](sycl::handler &cgh) { + cgh.require(acc_1); + cgh.require(acc_2); + cgh.require(acc_3); + cgh.require(acc_4); + cgh.require(acc_5); + cgh.require(acc_6); + cgh.require(acc_7); + cgh.require(acc_8); + cgh.require(acc_9); + + cgh.single_task( + [=]() { + acc_7[6] = 1; + acc_8[7] = 2; + acc_9[7] = 3; + acc_1[0] = acc_4[3]; + acc_2[1] = acc_5[4]; + acc_3[1] = acc_6[4]; + }); + }); + Queue.wait(); + +#if defined(accessor_new_api_test) + sycl::host_accessor host_acc(buf_data, sycl::read_only); +#elif defined(buffer_new_api_test) + auto host_acc = buf_data.get_host_access(sycl::read_only); +#endif + assert(host_acc[0] == 4 && host_acc[1] == 5 && host_acc[2] == 6); + assert(host_acc[3] == 4 && host_acc[4] == 5 && host_acc[5] == 6); + assert(host_acc[6] == 1 && host_acc[7] == 2 && host_acc[8] == 3); + } +} diff --git a/sycl/test/basic_tests/accessor/Inputs/host_accessor.cpp b/sycl/test/basic_tests/accessor/Inputs/host_accessor.cpp new file mode 100755 index 0000000000000..1f87f2d7f848a --- /dev/null +++ b/sycl/test/basic_tests/accessor/Inputs/host_accessor.cpp @@ -0,0 +1,137 @@ +//==----------------host_accessor.cpp - SYCL accessor basic test -----------==// +// +// 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 +// +//===----------------------------------------------------------------------===// +#include +#include + +int main() { + { + int data[3] = {3, 7, 9}; + + sycl::buffer buf_data(data, sycl::range<1>(3), + {cl::sycl::property::buffer::use_host_ptr()}); + + { +#if defined(accessor_new_api_test) + sycl::host_accessor acc(buf_data); +#elif defined(buffer_new_api_test) + auto acc = buf_data.get_host_access(); +#endif + + assert(acc.get_count() == 3); + assert(acc.get_range() == sycl::range<1>(3)); + assert(acc[0] == 3); + + acc[0] = 2; + + assert(data[0] == 2 && data[1] == 7 && data[2] == 9); + } + { +#if defined(accessor_new_api_test) + sycl::host_accessor acc(buf_data, sycl::read_only); +#elif defined(buffer_new_api_test) + auto acc = buf_data.get_host_access(sycl::read_only); +#endif + + assert(acc.get_count() == 3); + assert(acc.get_range() == sycl::range<1>(3)); + assert(acc[0] == 2); + } + { +#if defined(accessor_new_api_test) + sycl::host_accessor acc(buf_data, sycl::write_only); +#elif defined(buffer_new_api_test) + auto acc = buf_data.get_host_access(sycl::write_only); +#endif + + assert(acc.get_count() == 3); + assert(acc.get_range() == sycl::range<1>(3)); + acc[0] = 1; + assert(data[0] == 1 && data[1] == 7 && data[2] == 9); + } + { +#if defined(accessor_new_api_test) + sycl::host_accessor acc(buf_data, sycl::range<1>(2)); +#elif defined(buffer_new_api_test) + auto acc = buf_data.get_host_access(sycl::range<1>(2)); +#endif + + assert(acc.get_count() == 2); + assert(acc.get_range() == sycl::range<1>(2)); + assert(acc[0] == 1); + + acc[0] = 2; + + assert(data[0] == 2 && data[1] == 7 && data[2] == 9); + } + { +#if defined(accessor_new_api_test) + sycl::host_accessor acc(buf_data, sycl::range<1>(2), sycl::read_only); +#elif defined(buffer_new_api_test) + auto acc = buf_data.get_host_access(sycl::range<1>(2), sycl::read_only); +#endif + + assert(acc.get_count() == 2); + assert(acc.get_range() == sycl::range<1>(2)); + assert(acc[0] == 2); + } + { +#if defined(accessor_new_api_test) + sycl::host_accessor acc(buf_data, sycl::range<1>(2), sycl::write_only); +#elif defined(buffer_new_api_test) + auto acc = buf_data.get_host_access(sycl::range<1>(2), sycl::write_only); +#endif + + assert(acc.get_count() == 2); + assert(acc.get_range() == sycl::range<1>(2)); + acc[0] = 1; + assert(data[0] == 1 && data[1] == 7 && data[2] == 9); + } + { +#if defined(accessor_new_api_test) + sycl::host_accessor acc(buf_data, sycl::range<1>(2), sycl::id<1>(1)); +#elif defined(buffer_new_api_test) + auto acc = buf_data.get_host_access(sycl::range<1>(2), sycl::id<1>(1)); +#endif + + assert(acc.get_count() == 2); + assert(acc.get_range() == sycl::range<1>(2)); + assert(acc[0] == 7); + + acc[0] = 6; + + assert(data[0] == 1 && data[1] == 6 && data[2] == 9); + } + { +#if defined(accessor_new_api_test) + sycl::host_accessor acc(buf_data, sycl::range<1>(2), sycl::id<1>(1), + sycl::read_only); +#elif defined(buffer_new_api_test) + auto acc = buf_data.get_host_access(sycl::range<1>(2), sycl::id<1>(1), + sycl::read_only); +#endif + + assert(acc.get_count() == 2); + assert(acc.get_range() == sycl::range<1>(2)); + assert(acc[0] == 6); + } + { +#if defined(accessor_new_api_test) + sycl::host_accessor acc(buf_data, sycl::range<1>(2), sycl::id<1>(1), + sycl::write_only); +#elif defined(buffer_new_api_test) + auto acc = buf_data.get_host_access(sycl::range<1>(2), sycl::id<1>(1), + sycl::write_only); +#endif + + assert(acc.get_count() == 2); + assert(acc.get_range() == sycl::range<1>(2)); + acc[0] = 5; + assert(data[0] == 1 && data[1] == 5 && data[2] == 9); + } + } +} diff --git a/sycl/test/basic_tests/accessor/accessor.cpp b/sycl/test/basic_tests/accessor/accessor.cpp index b5ac254e91966..84286c6117eff 100644 --- a/sycl/test/basic_tests/accessor/accessor.cpp +++ b/sycl/test/basic_tests/accessor/accessor.cpp @@ -1,13 +1,8 @@ -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Dsimplification_test -std=c++17 %s -o %t.s.out +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: env SYCL_DEVICE_TYPE=HOST %t.s.out // RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.s.out // RUN: %GPU_RUN_PLACEHOLDER %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.s.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -// RUN: %ACC_RUN_PLACEHOLDER %t.s.out //==----------------accessor.cpp - SYCL accessor basic test ----------------==// // @@ -19,10 +14,6 @@ #include #include -namespace sycl { -using namespace cl::sycl; -} - struct IdxID1 { int x; @@ -72,13 +63,9 @@ int main() { {cl::sycl::property::buffer::use_host_ptr()}); sycl::id<1> id1(1); -#ifndef simplification_test auto acc_src = buf_src.get_access(); auto acc_dst = buf_dst.get_access(); -#else - sycl::host_accessor acc_src(buf_src, sycl::read_only); - sycl::host_accessor acc_dst(buf_dst); -#endif + assert(!acc_src.is_placeholder()); assert(acc_src.get_size() == sizeof(src)); assert(acc_src.get_count() == 2); @@ -104,11 +91,8 @@ int main() { data[i] = i; { sycl::buffer buf(data, sycl::range<3>(2, 3, 4)); -#ifndef simplification_test + auto acc = buf.get_access(); -#else - sycl::host_accessor acc(buf); -#endif assert(!acc.is_placeholder()); assert(acc.get_size() == sizeof(data)); @@ -133,11 +117,7 @@ int main() { {cl::sycl::property::buffer::use_host_ptr()}); Queue.submit([&](sycl::handler &cgh) { -#ifndef simplification_test auto acc = buf.get_access(cgh); -#else - sycl::accessor acc(buf, cgh); -#endif assert(!acc.is_placeholder()); assert(acc.get_size() == sizeof(int)); assert(acc.get_count() == 1); @@ -160,11 +140,7 @@ int main() { {cl::sycl::property::buffer::use_host_ptr()}); Queue.submit([&](sycl::handler &cgh) { -#ifndef simplification_test auto acc = buf.get_access(cgh); -#else - sycl::accessor acc(buf, cgh); -#endif cgh.parallel_for(Range, [=](sycl::item<2> itemID) { acc[itemID.get_id(0)][itemID.get_id(1)] += itemID.get_linear_id(); }); @@ -192,11 +168,7 @@ int main() { {cl::sycl::property::buffer::use_host_ptr()}); Queue.submit([&](sycl::handler &cgh) { -#ifndef simplification_test auto acc = buf.get_access(cgh); -#else - sycl::accessor acc(buf, cgh); -#endif cgh.parallel_for(Range, [=](sycl::item<3> itemID) { acc[itemID.get_id(0)][itemID.get_id(1)][itemID.get_id(2)] += itemID.get_linear_id(); @@ -223,22 +195,14 @@ int main() { sycl::buffer buf(sycl::range<1>(3)); Queue.submit([&](sycl::handler& cgh) { -#ifndef simplification_test auto dev_acc = buf.get_access(cgh); -#else - sycl::accessor dev_acc(buf, cgh, sycl::noinit); -#endif cgh.parallel_for( sycl::range<1>{3}, [=](sycl::id<1> index) { dev_acc[index] = 42; }); }); -#ifndef simplification_test auto host_acc = buf.get_access(); -#else - sycl::host_accessor host_acc(buf, sycl::read_only); -#endif for (int i = 0; i != 3; ++i) assert(host_acc[i] == 42); @@ -255,23 +219,15 @@ int main() { sycl::buffer buf(sycl::range<1>(3)); Queue.submit([&](sycl::handler& cgh) { -#ifndef simplification_test auto dev_acc = buf.get_access(cgh); -#else - sycl::accessor dev_acc(buf, cgh, sycl::write_only); -#endif cgh.parallel_for( sycl::range<1>{3}, [=](sycl::id<1> index) { dev_acc[index] = 42; }); }); -#ifndef simplification_test auto host_acc = buf.get_access(); -#else - sycl::host_accessor host_acc(buf, sycl::noinit); -#endif } catch (cl::sycl::exception e) { std::cout << "SYCL exception caught: " << e.what(); return 1; @@ -287,11 +243,7 @@ int main() { sycl::buffer buf((int *)array, sycl::range<1>(10), {cl::sycl::property::buffer::use_host_ptr()}); queue.submit([&](sycl::handler &cgh) { -#ifndef simplification_test auto acc = buf.get_access(cgh); -#else - sycl::accessor acc(buf, cgh); -#endif auto acc_wrapped = AccWrapper{acc}; cgh.parallel_for( sycl::range<1>(buf.get_count()), [=](sycl::item<1> it) { @@ -321,13 +273,8 @@ int main() { sycl::buffer buf2((int *)array2, sycl::range<1>(10), {cl::sycl::property::buffer::use_host_ptr()}); queue.submit([&](sycl::handler &cgh) { -#ifndef simplification_test auto acc1 = buf1.get_access(cgh); auto acc2 = buf2.get_access(cgh); -#else - sycl::accessor acc1(buf1, cgh); - sycl::accessor acc2(buf2, cgh); -#endif auto acc_wrapped = AccsWrapper{10, acc1, 5, acc2}; cgh.parallel_for( @@ -357,11 +304,7 @@ int main() { sycl::buffer buf((int *)array, sycl::range<1>(10), {cl::sycl::property::buffer::use_host_ptr()}); queue.submit([&](sycl::handler &cgh) { -#ifndef simplification_test auto acc = buf.get_access(cgh); -#else - sycl::accessor acc(buf, cgh); -#endif auto acc_wrapped = AccWrapper{acc}; Wrapper1 wr1; auto wr2 = Wrapper2{wr1, acc_wrapped}; @@ -389,24 +332,15 @@ int main() { sycl::buffer buf(array, sycl::range<1>(3)); queue.submit([&](sycl::handler& cgh) { -#ifndef simplification_test auto acc1 = buf.get_access(cgh); auto acc2 = buf.get_access(cgh); -#else - sycl::accessor acc1(buf, cgh, sycl::read_only); - sycl::accessor acc2(buf, cgh); -#endif cgh.parallel_for( sycl::range<1>{3}, [=](sycl::id<1> index) { acc2[index] = 41 + acc1[index]; }); }); -#ifndef simplification_test auto host_acc = buf.get_access(); -#else - sycl::host_accessor host_acc(buf, sycl::read_only); -#endif for (int i = 0; i != 3; ++i) assert(host_acc[i] == 42); @@ -459,17 +393,12 @@ int main() { sycl::accessor acc1(buf1, cgh); -#ifndef simplification_test sycl::accessor acc2(buf2, cgh); sycl::accessor acc3(buf3, cgh, sycl::range<1>(1)); -#else - sycl::accessor acc2(buf2, cgh); - sycl::accessor acc3(buf3, cgh, sycl::range<1>(1)); -#endif cgh.single_task([=]() { acc1 *= 2; acc2[0] *= 2; @@ -480,17 +409,12 @@ int main() { sycl::accessor acc4(buf1); -#ifndef simplification_test sycl::accessor acc5(buf2); sycl::accessor acc6(buf3, sycl::range<1>(1)); -#else - sycl::host_accessor acc5(buf2, sycl::read_only); - sycl::host_accessor acc6(buf3, sycl::range<1>(1), sycl::read_only); -#endif assert(acc4 == 2); assert(acc5[0] == 4); @@ -509,28 +433,19 @@ int main() { sycl::queue queue; queue.submit([&](sycl::handler &cgh) { -#ifndef simplification_test sycl::accessor D(d, cgh); sycl::accessor C(c, cgh); -#else - sycl::accessor D(d, cgh, sycl::write_only); - sycl::accessor C(c, cgh, sycl::read_constant); -#endif cgh.single_task([=]() { D[0] = C[0]; }); }); -#ifndef simplification_test auto host_acc = d.get_access(); -#else - sycl::host_accessor host_acc(d, sycl::read_only); -#endif assert(host_acc[0] == 399); } @@ -550,7 +465,6 @@ int main() { sycl::buffer d(&data, sycl::range<1>(1)); sycl::buffer c(&cnst, sycl::range<1>(1)); -#ifndef simplification_test sycl::accessor @@ -559,10 +473,6 @@ int main() { sycl::access::target::constant_buffer, sycl::access::placeholder::true_t> C(c); -#else - sycl::accessor D(d, sycl::write_only); - sycl::accessor C(c, sycl::read_constant); -#endif sycl::queue queue; queue.submit([&](sycl::handler &cgh) { @@ -574,11 +484,7 @@ int main() { }); }); -#ifndef simplification_test auto host_acc = d.get_access(); -#else - sycl::host_accessor host_acc(d, sycl::read_only); -#endif assert(host_acc[0] == 399); } diff --git a/sycl/test/basic_tests/accessor/device_accessor_deduction.cpp b/sycl/test/basic_tests/accessor/device_accessor_deduction.cpp new file mode 100755 index 0000000000000..d5081b0fbd344 --- /dev/null +++ b/sycl/test/basic_tests/accessor/device_accessor_deduction.cpp @@ -0,0 +1,5 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Daccessor_new_api_test -std=c++17 %S/Inputs/device_accessor.cpp -o %t.out +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out diff --git a/sycl/test/basic_tests/accessor/get_device_access_deduction.cpp b/sycl/test/basic_tests/accessor/get_device_access_deduction.cpp new file mode 100755 index 0000000000000..b0bb0a7668f1e --- /dev/null +++ b/sycl/test/basic_tests/accessor/get_device_access_deduction.cpp @@ -0,0 +1,5 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Dbuffer_new_api_test -std=c++17 %S/Inputs/device_accessor.cpp -o %t.out +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out diff --git a/sycl/test/basic_tests/accessor/get_host_access_deduction.cpp b/sycl/test/basic_tests/accessor/get_host_access_deduction.cpp new file mode 100755 index 0000000000000..3cb155791fd97 --- /dev/null +++ b/sycl/test/basic_tests/accessor/get_host_access_deduction.cpp @@ -0,0 +1,5 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Dbuffer_new_api_test -std=c++17 %S/Inputs/host_accessor.cpp -o %t.out +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out diff --git a/sycl/test/basic_tests/accessor/host_accessor_deduction.cpp b/sycl/test/basic_tests/accessor/host_accessor_deduction.cpp new file mode 100755 index 0000000000000..2d60da1edf23d --- /dev/null +++ b/sycl/test/basic_tests/accessor/host_accessor_deduction.cpp @@ -0,0 +1,5 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Daccessor_new_api_test -std=c++17 %S/Inputs/host_accessor.cpp -o %t.out +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out