From 7ba4b7fe993e6348a52dcfe0e3cf59ea01327885 Mon Sep 17 00:00:00 2001 From: iburylov Date: Mon, 8 Jun 2020 18:08:45 +0300 Subject: [PATCH 01/16] [SYCL] Accessor tags, CTAD and host_accessor Signed-off-by: iburylov --- sycl/include/CL/sycl/access/access.hpp | 36 +- sycl/include/CL/sycl/accessor.hpp | 511 +++++++++++++++-- sycl/include/CL/sycl/handler.hpp | 5 +- sycl/include/CL/sycl/property_list.hpp | 16 + .../accessor/accessor_simplification.cpp | 520 ++++++++++++++++++ 5 files changed, 1045 insertions(+), 43 deletions(-) create mode 100644 sycl/test/basic_tests/accessor/accessor_simplification.cpp diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index 240032c47ead8..2c28490c28708 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -11,7 +11,6 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { -namespace access { enum class target { global_buffer = 2014, @@ -23,6 +22,12 @@ enum class target { image_array }; +// Backward compatibility namespace nesting +namespace access { +using sycl::target; +} + +namespace access { enum class mode { read = 1024, write, @@ -31,6 +36,33 @@ enum class mode { discard_read_write, atomic }; +} + +using access_mode = access::mode; + +namespace access { +enum class placeholder { false_t, true_t }; +} + +#if __cplusplus > 201402L + +template struct mode_tag_t { + explicit mode_tag_t() = default; +}; + +template struct mode_target_tag_t { + explicit mode_target_tag_t() = default; +}; + +inline constexpr mode_tag_t read_only{}; +inline constexpr mode_tag_t read_write{}; +inline constexpr mode_tag_t write_only{}; +inline constexpr mode_target_tag_t + read_constant{}; + +#endif + +namespace access { enum class fence_space { local_space, @@ -38,8 +70,6 @@ enum class fence_space { global_and_local }; -enum class placeholder { false_t, true_t }; - enum class address_space : int { private_space = 0, global_space, diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 3fac2cb0c985e..9ab0b3fd1fdac 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -134,6 +134,7 @@ /// "image_accessor" -> a2; /// "image_accessor" -> a4; /// "image_accessor" -> a5; +/// a1 -> "host_accessor"; /// } /// \enddot /// @@ -156,6 +157,13 @@ // | | | | global_buffer | +-------------+ // | | | | constant_buffer | // | | | +-----------------+ +// | | | | +// | | | v +// | | | +-----------------+ +// | | | | | +// | | | | host_accessor | +// | | | | | +// | | | +-----------------+ // | | | // | | +------------------------------------+ // | | | @@ -190,7 +198,8 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { -template class accessor; @@ -709,7 +718,7 @@ class accessor : #endif public detail::accessor_common { - +protected: static_assert((AccessTarget == access::target::global_buffer || AccessTarget == access::target::constant_buffer || AccessTarget == access::target::host_buffer), @@ -756,6 +765,34 @@ class accessor : return Result; } + template static constexpr bool IsSameAsBuffer() { + return std::is_same::value && (Dims > 0) && (Dims == Dimensions); + } + + static access::mode getAdjustedMode(const property_list &PropertyList) { + access::mode AdjustedMode = AccessMode; + + if (PropertyList.has_property()) { + if (AdjustedMode == access::mode::write) { + AdjustedMode = access::mode::discard_write; + } else if (AdjustedMode == access::mode::read_write) { + AdjustedMode = access::mode::discard_read_write; + } + } + + return AdjustedMode; + } + +#if __cplusplus > 201402L + + template static constexpr bool IsValidTag() { + return std::is_same>::value || + std::is_same>::value; + } + +#endif + #ifdef __SYCL_DEVICE_ONLY__ id &getOffset() { return impl.Offset; } @@ -811,17 +848,45 @@ class accessor : using reference = DataT &; using const_reference = const DataT &; - template = 1 + // -------+---------+-------+----+-----+-------------- + // buffer | | | | | property_list + // buffer | | | | tag | property_list + // buffer | handler | | | | property_list + // buffer | handler | | | tag | property_list + // buffer | | range | | | property_list + // buffer | | range | | tag | property_list + // buffer | handler | range | | | property_list + // buffer | handler | range | | tag | property_list + // buffer | | range | id | | property_list + // buffer | | range | id | tag | property_list + // buffer | handler | range | id | | property_list + // buffer | handler | range | id | tag | property_list + // -------+---------+-------+----+-----+-------------- + +public: + template * = nullptr> - accessor(buffer &BufferRef) + std::is_same::value && Dims == 0 && + ((!IsPlaceH && IsHostBuf) || + (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))> * = nullptr> + accessor(buffer &BufferRef, + const property_list &PropertyList = {}) #ifdef __SYCL_DEVICE_ONLY__ : impl(id(), range<1>{1}, BufferRef.get_range()) { + (void)PropertyList; #else : AccessorBaseHost( /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}), - detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), AccessMode, + detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), + getAdjustedMode(PropertyList), detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT), BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) { if (!IsPlaceH) @@ -829,41 +894,45 @@ class accessor : #endif } - template - > - accessor(buffer &BufferRef, - handler &CommandGroupHandler) + template ::value && (Dims == 0) && + (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>> + accessor(buffer &BufferRef, handler &CommandGroupHandler, + const property_list &PropertyList = {}) #ifdef __SYCL_DEVICE_ONLY__ : impl(id(), range<1>{1}, BufferRef.get_range()) { (void)CommandGroupHandler; + (void)PropertyList; } #else : AccessorBaseHost( /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}), - detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), AccessMode, + detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), + getAdjustedMode(PropertyList), detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT), BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) { detail::associateWithHandler(CommandGroupHandler, this, AccessTarget); } #endif - template 0) && (Dims == Dimensions) && + template () && ((!IsPlaceH && IsHostBuf) || (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>> - accessor(buffer &BufferRef) + accessor(buffer &BufferRef, + const property_list &PropertyList = {}) #ifdef __SYCL_DEVICE_ONLY__ : impl(id(), BufferRef.get_range(), BufferRef.get_range()) { + (void)PropertyList; } #else : AccessorBaseHost( /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), - detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), AccessMode, + detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), + getAdjustedMode(PropertyList), detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT), BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) { if (!IsPlaceH) @@ -871,71 +940,182 @@ class accessor : } #endif - template 201402L + + template () && + IsValidTag() && IsPlaceH && + (IsGlobalBuf || IsConstantBuf)>> + accessor(buffer &BufferRef, TagT, + const property_list &PropertyList = {}) + : accessor(BufferRef, PropertyList) {} + +#endif + + template 0) && (Dims == Dimensions) && + IsSameAsBuffer() && (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>> - accessor(buffer &BufferRef, - handler &CommandGroupHandler) + accessor(buffer &BufferRef, handler &CommandGroupHandler, + const property_list &PropertyList = {}) #ifdef __SYCL_DEVICE_ONLY__ : impl(id(), BufferRef.get_range(), BufferRef.get_range()) { (void)CommandGroupHandler; + (void)PropertyList; } #else : AccessorBaseHost( /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), - detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), AccessMode, + detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), + getAdjustedMode(PropertyList), detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT), BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) { detail::associateWithHandler(CommandGroupHandler, this, AccessTarget); } #endif - template 0) && (Dims == Dimensions) && +#if __cplusplus > 201402L + + template () && + IsValidTag() && !IsPlaceH && + (IsGlobalBuf || IsConstantBuf)>> + accessor(buffer &BufferRef, handler &CommandGroupHandler, + TagT, const property_list &PropertyList = {}) + : accessor(BufferRef, CommandGroupHandler, PropertyList) {} + +#endif + + template () && + ((!IsPlaceH && IsHostBuf) || + (IsPlaceH && + (IsGlobalBuf || IsConstantBuf)))>> + accessor(buffer &BufferRef, + range AccessRange, + const property_list &PropertyList = {}) + : accessor(BufferRef, AccessRange, {}, PropertyList) {} + +#if __cplusplus > 201402L + + template () && + IsValidTag() && IsPlaceH && + (IsGlobalBuf || IsConstantBuf)>> + accessor(buffer &BufferRef, + range AccessRange, TagT, + const property_list &PropertyList = {}) + : accessor(BufferRef, AccessRange, {}, PropertyList) {} + +#endif + + template () && + (!IsPlaceH && + (IsGlobalBuf || IsConstantBuf))>> + accessor(buffer &BufferRef, handler &CommandGroupHandler, + range AccessRange, + const property_list &PropertyList = {}) + : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, + PropertyList) {} + +#if __cplusplus > 201402L + + template () && + IsValidTag() && !IsPlaceH && + (IsGlobalBuf || IsConstantBuf)>> + accessor(buffer &BufferRef, handler &CommandGroupHandler, + range AccessRange, TagT, + const property_list &PropertyList = {}) + : accessor(BufferRef, CommandGroupHandler, AccessRange, {}, + PropertyList) {} + +#endif + + template () && ((!IsPlaceH && IsHostBuf) || (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>> - accessor(buffer &BufferRef, - range AccessRange, id AccessOffset = {}) + accessor(buffer &BufferRef, + range AccessRange, id AccessOffset, + const property_list &PropertyList = {}) #ifdef __SYCL_DEVICE_ONLY__ : impl(AccessOffset, AccessRange, BufferRef.get_range()) { + (void)PropertyList; } #else : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset), detail::convertToArrayOfN<3, 1>(AccessRange), detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), - AccessMode, detail::getSyclObjImpl(BufferRef).get(), - Dimensions, sizeof(DataT), BufferRef.OffsetInBytes, + getAdjustedMode(PropertyList), + detail::getSyclObjImpl(BufferRef).get(), Dimensions, + sizeof(DataT), BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) { if (!IsPlaceH) addHostAccessorAndWait(AccessorBaseHost::impl.get()); } #endif - template 0) && (Dims == Dimensions) && +#if __cplusplus > 201402L + + template () && + IsValidTag() && IsPlaceH && + (IsGlobalBuf || IsConstantBuf)>> + accessor(buffer &BufferRef, + range AccessRange, id AccessOffset, TagT, + const property_list &PropertyList = {}) + : accessor(BufferRef, AccessRange, AccessOffset, PropertyList) {} + +#endif + + template () && (!IsPlaceH && (IsGlobalBuf || IsConstantBuf))>> - accessor(buffer &BufferRef, - handler &CommandGroupHandler, range AccessRange, - id AccessOffset = {}) + accessor(buffer &BufferRef, handler &CommandGroupHandler, + range AccessRange, id AccessOffset, + const property_list &PropertyList = {}) #ifdef __SYCL_DEVICE_ONLY__ : impl(AccessOffset, AccessRange, BufferRef.get_range()) { (void)CommandGroupHandler; + (void)PropertyList; } #else : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset), detail::convertToArrayOfN<3, 1>(AccessRange), detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), - AccessMode, detail::getSyclObjImpl(BufferRef).get(), - Dimensions, sizeof(DataT), BufferRef.OffsetInBytes, + getAdjustedMode(PropertyList), + detail::getSyclObjImpl(BufferRef).get(), Dimensions, + sizeof(DataT), BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) { detail::associateWithHandler(CommandGroupHandler, this, AccessTarget); } #endif +#if __cplusplus > 201402L + + template () && + IsValidTag() && !IsPlaceH && + (IsGlobalBuf || IsConstantBuf)>> + accessor(buffer &BufferRef, handler &CommandGroupHandler, + range AccessRange, id AccessOffset, TagT, + const property_list &PropertyList = {}) + : accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset, + PropertyList) {} + +#endif + constexpr bool is_placeholder() const { return IsPlaceH; } size_t get_size() const { return getAccessRange().size() * sizeof(DataT); } @@ -1055,6 +1235,48 @@ class accessor : bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); } }; +#if __cplusplus > 201402L + +template +accessor(buffer, Ts...) + -> accessor; + +template +accessor(buffer, handler, Ts...) + -> accessor; + +template +accessor(buffer, Ts..., mode_tag_t, + property_list = {}) + -> accessor; + +template +accessor(buffer, handler, Ts..., + mode_tag_t, property_list = {}) + -> accessor; + +template +accessor(buffer, Ts..., + mode_target_tag_t, property_list = {}) + -> accessor; + +template +accessor(buffer, handler, Ts..., + mode_target_tag_t, property_list = {}) + -> accessor; + +#endif + /// Local accessor /// /// \ingroup sycl_api_acc @@ -1329,6 +1551,221 @@ class accessor +class host_accessor + : public accessor { +protected: + using AccessorT = accessor; + + constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions; + + template static constexpr bool IsSameAsBuffer() { + return std::is_same::value && (Dims > 0) && (Dims == Dimensions); + } + +#if __cplusplus > 201402L + + template static constexpr bool IsValidTag() { + return std::is_same>::value; + } + +#endif + + void + __init(typename accessor::ConcreteASPtrType Ptr, + range AccessRange, range MemRange, + id Offset) { + AccessorT::__init(Ptr, AccessRange, MemRange, Offset); + } + +public: + host_accessor() : AccessorT() {} + + // The list of host_accessor constructors with their arguments + // -------+---------+-------+----+----------+-------------- + // Dimensions = 0 + // -------+---------+-------+----+----------+-------------- + // buffer | | | | | property_list + // buffer | handler | | | | property_list + // -------+---------+-------+----+----------+-------------- + // Dimensions >= 1 + // -------+---------+-------+----+----------+-------------- + // buffer | | | | | property_list + // buffer | | | | mode_tag | property_list + // buffer | handler | | | | property_list + // buffer | handler | | | mode_tag | property_list + // buffer | | range | | | property_list + // buffer | | range | | mode_tag | property_list + // buffer | handler | range | | | property_list + // buffer | handler | range | | mode_tag | property_list + // buffer | | range | id | | property_list + // buffer | | range | id | mode_tag | property_list + // buffer | handler | range | id | | property_list + // buffer | handler | range | id | mode_tag | property_list + // -------+---------+-------+----+----------+-------------- + + template ::value && Dims == 0>> + host_accessor(buffer &BufferRef, + const property_list &PropertyList = {}) + : AccessorT(BufferRef, PropertyList) {} + + template ::value && Dims == 0>> + host_accessor(buffer &BufferRef, + handler &CommandGroupHandler, + const property_list &PropertyList = {}) + : AccessorT(BufferRef, PropertyList) { + (void)CommandGroupHandler; + } + + template ()>> + host_accessor(buffer &BufferRef, + const property_list &PropertyList = {}) + : AccessorT(BufferRef, PropertyList) {} + +#if __cplusplus > 201402L + + template ()>> + host_accessor(buffer &BufferRef, + mode_tag_t, const property_list &PropertyList = {}) + : host_accessor(BufferRef, PropertyList) {} + +#endif + + template ()>> + host_accessor(buffer &BufferRef, + handler &CommandGroupHandler, + const property_list &PropertyList = {}) + : AccessorT(BufferRef, PropertyList) { + (void)CommandGroupHandler; + } + +#if __cplusplus > 201402L + + template ()>> + host_accessor(buffer &BufferRef, + handler &CommandGroupHandler, mode_tag_t, + const property_list &PropertyList = {}) + : host_accessor(BufferRef, CommandGroupHandler, PropertyList) {} + +#endif + + template ()>> + host_accessor(buffer &BufferRef, + range AccessRange, + const property_list &PropertyList = {}) + : AccessorT(BufferRef, AccessRange, {}, PropertyList) {} + +#if __cplusplus > 201402L + + template ()>> + host_accessor(buffer &BufferRef, + range AccessRange, mode_tag_t, + const property_list &PropertyList = {}) + : host_accessor(BufferRef, AccessRange, {}, PropertyList) {} + +#endif + + template ()>> + host_accessor(buffer &BufferRef, + handler &CommandGroupHandler, range AccessRange, + const property_list &PropertyList = {}) + : AccessorT(BufferRef, AccessRange, {}, PropertyList) { + (void)CommandGroupHandler; + } + +#if __cplusplus > 201402L + + template ()>> + host_accessor(buffer &BufferRef, + handler &CommandGroupHandler, range AccessRange, + mode_tag_t, const property_list &PropertyList = {}) + : host_accessor(BufferRef, CommandGroupHandler, AccessRange, {}, + PropertyList) {} + +#endif + + template ()>> + host_accessor(buffer &BufferRef, + range AccessRange, id AccessOffset, + const property_list &PropertyList = {}) + : AccessorT(BufferRef, AccessRange, AccessOffset, PropertyList) {} + +#if __cplusplus > 201402L + + template ()>> + host_accessor(buffer &BufferRef, + range AccessRange, id AccessOffset, + mode_tag_t, const property_list &PropertyList = {}) + : host_accessor(BufferRef, AccessRange, AccessOffset, PropertyList) {} + +#endif + + template ()>> + host_accessor(buffer &BufferRef, + handler &CommandGroupHandler, range AccessRange, + id AccessOffset, + const property_list &PropertyList = {}) + : AccessorT(BufferRef, AccessRange, AccessOffset, PropertyList) { + (void)CommandGroupHandler; + } + +#if __cplusplus > 201402L + + template ()>> + host_accessor(buffer &BufferRef, + handler &CommandGroupHandler, range AccessRange, + id AccessOffset, mode_tag_t, + const property_list &PropertyList = {}) + : host_accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset, + PropertyList) {} + +#endif +}; + +#if __cplusplus > 201402L + +template +host_accessor(buffer, Ts...) + -> host_accessor; + +template +host_accessor(buffer, handler, Ts...) + -> host_accessor; + +template +host_accessor(buffer, Ts..., + mode_tag_t, property_list = {}) + -> host_accessor; + +template +host_accessor(buffer, handler, Ts..., + mode_tag_t, property_list = {}) + -> host_accessor; + +#endif + } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index 0614ef3105c82..fee04f1588d18 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -749,9 +749,8 @@ class __SYCL_EXPORT handler { /// \param Acc is a SYCL accessor describing required memory region. template - void - require(accessor - Acc) { + void require(accessor &Acc) { #ifndef __SYCL_DEVICE_ONLY__ associateWithHandler(&Acc, AccTarget); #else diff --git a/sycl/include/CL/sycl/property_list.hpp b/sycl/include/CL/sycl/property_list.hpp index fbc3ef7bf5402..a7210f9199032 100644 --- a/sycl/include/CL/sycl/property_list.hpp +++ b/sycl/include/CL/sycl/property_list.hpp @@ -42,6 +42,8 @@ class enable_profiling; class in_order; } // namespace queue +class noinit; + namespace detail { // List of all properties' IDs. @@ -60,6 +62,9 @@ enum PropKind { QueueEnableProfiling, InOrder, + // Accessor + NoInit, + PropKindSize }; @@ -147,6 +152,9 @@ RegisterProp(PropKind::BufferContextBound, buffer::context_bound); RegisterProp(PropKind::QueueEnableProfiling, queue::enable_profiling); RegisterProp(PropKind::InOrder, queue::in_order); +// Buffer +RegisterProp(PropKind::NoInit, noinit); + // Sentinel, needed for automatic build of tuple in property_list. RegisterProp(PropKind::PropKindSize, PropBase); @@ -212,8 +220,16 @@ class enable_profiling class in_order : public detail::Prop {}; } // namespace queue +class noinit : public detail::Prop {}; + } // namespace property +#if __cplusplus > 201402L + +inline constexpr property::noinit noinit; + +#endif + class property_list { // The structs validate that all objects passed are base of PropBase class. diff --git a/sycl/test/basic_tests/accessor/accessor_simplification.cpp b/sycl/test/basic_tests/accessor/accessor_simplification.cpp new file mode 100644 index 0000000000000..fa978acf4be14 --- /dev/null +++ b/sycl/test/basic_tests/accessor/accessor_simplification.cpp @@ -0,0 +1,520 @@ +// RUN: %clangxx -fsycl -std=c++17 %s -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 +//==---------accessor_simplification.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 + +namespace sycl { +using namespace cl::sycl; +} + +struct IdxID1 { + int x; + + IdxID1(int x) : x(x) {} + operator sycl::id<1>() { return x; } +}; + +struct IdxID3 { + int x; + int y; + int z; + + IdxID3(int x, int y, int z) : x(x), y(y), z(z) {} + operator sycl::id<3>() { return sycl::id<3>(x, y, z); } +}; + +struct IdxSzT { + int x; + + IdxSzT(int x) : x(x) {} + operator size_t() { return x; } +}; + +template +struct AccWrapper { Acc accessor; }; + +template +struct AccsWrapper { + int a; + Acc1 accessor1; + int b; + Acc2 accessor2; +}; + +struct Wrapper1 { + int a; + int b; +}; + +template +struct Wrapper2 { + Wrapper1 w1; + AccWrapper wrapped; +}; + +template +struct Wrapper3 { Wrapper2 w2; }; + +int main() { + // Host accessor. + { + int src[2] = {3, 7}; + int dst[2]; + + sycl::buffer buf_src(src, sycl::range<1>(2), + {cl::sycl::property::buffer::use_host_ptr()}); + sycl::buffer buf_dst(dst, sycl::range<1>(2), + {cl::sycl::property::buffer::use_host_ptr()}); + + sycl::id<1> id1(1); + //auto acc_src = buf_src.get_access(); + sycl::host_accessor acc_src(buf_src, sycl::read_only); + //auto acc_dst = buf_dst.get_access(); + sycl::host_accessor acc_dst(buf_dst); + + assert(!acc_src.is_placeholder()); + assert(acc_src.get_size() == sizeof(src)); + assert(acc_src.get_count() == 2); + assert(acc_src.get_range() == sycl::range<1>(2)); + + // Make sure that operator[] is defined for both size_t and id<1>. + // Implicit conversion from IdxSzT to size_t guarantees that no + // implicit conversion from size_t to id<1> will happen. + assert(acc_src[IdxSzT(0)] + acc_src[IdxID1(1)] == 10); + + acc_dst[0] = acc_src[0] + acc_src[IdxID1(0)]; + acc_dst[id1] = acc_src[1] + acc_src[IdxSzT(1)]; + assert(dst[0] == 6 && dst[1] == 14); + } + + // Three-dimensional host accessor. + { + int data[24]; + for (int i = 0; i < 24; ++i) + data[i] = i; + { + sycl::buffer buf(data, sycl::range<3>(2, 3, 4)); + //auto acc = buf.get_access(); + sycl::host_accessor acc(buf); + + assert(!acc.is_placeholder()); + assert(acc.get_size() == sizeof(data)); + assert(acc.get_count() == 24); + assert(acc.get_range() == sycl::range<3>(2, 3, 4)); + + for (int i = 0; i < 2; ++i) + for (int j = 0; j < 3; ++j) + for (int k = 0; k < 4; ++k) + acc[IdxID3(i, j, k)] += acc[sycl::id<3>(i, j, k)]; + } + for (int i = 0; i < 24; ++i) { + assert(data[i] == 2 * i); + } + } + int data = 5; + // Device accessor. + { + sycl::queue Queue; + + sycl::buffer buf(&data, sycl::range<1>(1), + {cl::sycl::property::buffer::use_host_ptr()}); + + Queue.submit([&](sycl::handler &cgh) { + //auto acc = buf.get_access(cgh); + sycl::accessor acc(buf, cgh); + assert(!acc.is_placeholder()); + assert(acc.get_size() == sizeof(int)); + assert(acc.get_count() == 1); + assert(acc.get_range() == sycl::range<1>(1)); + cgh.single_task( + [=]() { acc[IdxSzT(0)] += acc[IdxID1(0)]; }); + }); + Queue.wait(); + } + assert(data == 10); + + // Device accessor with 2-dimensional subscript operators. + { + sycl::queue Queue; + if (!Queue.is_host()) { + int array[2][3] = {0}; + { + sycl::range<2> Range(2, 3); + sycl::buffer buf((int *)array, Range, + {cl::sycl::property::buffer::use_host_ptr()}); + + Queue.submit([&](sycl::handler &cgh) { + //auto acc = buf.get_access(cgh); + sycl::accessor acc(buf, cgh); + cgh.parallel_for(Range, [=](sycl::item<2> itemID) { + acc[itemID.get_id(0)][itemID.get_id(1)] += itemID.get_linear_id(); + }); + }); + Queue.wait(); + } + for (int i = 0; i < 2; i++) { + for (int j = 0; j < 3; j++) { + std::cout << "array[" << i << "][" << j << "]=" << array[i][j] + << std::endl; + assert(array[i][j] == i * 3 + j); + } + } + } + } + + // Device accessor with 3-dimensional subscript operators. + { + sycl::queue Queue; + if (!Queue.is_host()) { + int array[2][3][4] = {0}; + { + sycl::range<3> Range(2, 3, 4); + sycl::buffer buf((int *)array, Range, + {cl::sycl::property::buffer::use_host_ptr()}); + + Queue.submit([&](sycl::handler &cgh) { + //auto acc = buf.get_access(cgh); + sycl::accessor acc(buf, cgh); + 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(); + }); + }); + Queue.wait(); + } + for (int i = 0; i < 2; i++) { + for (int j = 0; j < 3; j++) { + for (int k = 0; k < 4; k++) { + std::cout << "array[" << i << "][" << j << "][" << k + << "]=" << array[i][j][k] << std::endl; + assert(array[i][j][k] == k + 4 * (j + 3 * i)); + } + } + } + } + } + + // Discard write accessor. + { + try { + sycl::queue Queue; + sycl::buffer buf(sycl::range<1>(3)); + + Queue.submit([&](sycl::handler &cgh) { + //auto dev_acc = buf.get_access(cgh); + sycl::accessor dev_acc(buf, cgh, sycl::noinit); + + cgh.parallel_for( + sycl::range<1>{3}, + [=](sycl::id<1> index) { dev_acc[index] = 42; }); + }); + + //auto host_acc = buf.get_access(); + sycl::host_accessor host_acc(buf, sycl::read_only); // TODO: it is read_write now - is there a reason to have read only accessor? + + for (int i = 0; i != 3; ++i) + assert(host_acc[i] == 42); + + } catch (cl::sycl::exception e) { + std::cout << "SYCL exception caught: " << e.what(); + return 1; + } + } + + // Discard read-write accessor. + { + try { + sycl::queue Queue; + sycl::buffer buf(sycl::range<1>(3)); + + Queue.submit([&](sycl::handler &cgh) { + //auto dev_acc = buf.get_access(cgh); + sycl::accessor dev_acc(buf, cgh, sycl::write_only); + + cgh.parallel_for( + sycl::range<1>{3}, + [=](sycl::id<1> index) { dev_acc[index] = 42; }); + }); + + //auto host_acc = buf.get_access(); + sycl::host_accessor host_acc(buf, sycl::noinit); + } catch (cl::sycl::exception e) { + std::cout << "SYCL exception caught: " << e.what(); + return 1; + } + } + + // Check that accessor is initialized when accessor is wrapped to some class. + { + sycl::queue queue; + if (!queue.is_host()) { + int array[10] = {0}; + { + sycl::buffer buf((int *)array, sycl::range<1>(10), + {cl::sycl::property::buffer::use_host_ptr()}); + queue.submit([&](sycl::handler &cgh) { + //auto acc = buf.get_access(cgh); + sycl::accessor acc(buf, cgh); + auto acc_wrapped = AccWrapper{acc}; + cgh.parallel_for( + sycl::range<1>(buf.get_count()), [=](sycl::item<1> it) { + auto idx = it.get_linear_id(); + acc_wrapped.accessor[idx] = 333; + }); + }); + queue.wait(); + } + for (int i = 0; i < 10; i++) { + std::cout << "array[" << i << "]=" << array[i] << std::endl; + assert(array[i] == 333); + } + } + } + + // Case when several accessors are wrapped to some class. Check that they are + // initialized in proper way and value is assigned. + { + sycl::queue queue; + if (!queue.is_host()) { + int array1[10] = {0}; + int array2[10] = {0}; + { + sycl::buffer buf1((int *)array1, sycl::range<1>(10), + {cl::sycl::property::buffer::use_host_ptr()}); + sycl::buffer buf2((int *)array2, sycl::range<1>(10), + {cl::sycl::property::buffer::use_host_ptr()}); + queue.submit([&](sycl::handler &cgh) { + //auto acc1 = buf1.get_access(cgh); + sycl::accessor acc1(buf1, cgh); + //auto acc2 = buf2.get_access(cgh); + sycl::accessor acc2(buf2, cgh); + auto acc_wrapped = + AccsWrapper{10, acc1, 5, acc2}; + cgh.parallel_for( + sycl::range<1>(10), [=](sycl::item<1> it) { + auto idx = it.get_linear_id(); + acc_wrapped.accessor1[idx] = 333; + acc_wrapped.accessor2[idx] = 666; + }); + }); + queue.wait(); + } + for (int i = 0; i < 10; i++) { + std::cout << "array1[" << i << "]=" << array1[i] << std::endl; + std::cout << "array2[" << i << "]=" << array2[i] << std::endl; + assert(array1[i] == 333); + assert(array2[i] == 666); + } + } + } + + // Several levels of wrappers for accessor. + { + sycl::queue queue; + if (!queue.is_host()) { + int array[10] = {0}; + { + sycl::buffer buf((int *)array, sycl::range<1>(10), + {cl::sycl::property::buffer::use_host_ptr()}); + queue.submit([&](sycl::handler &cgh) { + //auto acc = buf.get_access(cgh); + sycl::accessor acc(buf, cgh); + auto acc_wrapped = AccWrapper{acc}; + Wrapper1 wr1; + auto wr2 = Wrapper2{wr1, acc_wrapped}; + auto wr3 = Wrapper3{wr2}; + cgh.parallel_for( + sycl::range<1>(buf.get_count()), [=](sycl::item<1> it) { + auto idx = it.get_linear_id(); + wr3.w2.wrapped.accessor[idx] = 333; + }); + }); + queue.wait(); + } + for (int i = 0; i < 10; i++) { + std::cout << "array[" << i << "]=" << array[i] << std::endl; + assert(array[i] == 333); + } + } + } + + // Two accessors to the same buffer. + { + try { + sycl::queue queue; + int array[3] = {1, 1, 1}; + sycl::buffer buf(array, sycl::range<1>(3)); + + std::cout << "We are here 1" << std::endl; + + queue.submit([&](sycl::handler &cgh) { + //auto acc1 = buf.get_access(cgh); + sycl::accessor acc1(buf, cgh, sycl::read_only); + //auto acc2 = buf.get_access(cgh); + sycl::accessor acc2(buf, cgh); + + cgh.parallel_for( + sycl::range<1>{3}, + [=](sycl::id<1> index) { + acc2[index] = 41 + acc1[index]; + }); + }); + + //auto host_acc = buf.get_access(); + sycl::host_accessor host_acc(buf, sycl::read_only); + for (int i = 0; i != 3; ++i) + assert(host_acc[i] == 42); + + } catch (cl::sycl::exception e) { + std::cout << "SYCL exception caught: " << e.what(); + return 1; + } + } + + // Accessor with dimensionality 0. + { + try { + int data = -1; + { + sycl::buffer b(&data, sycl::range<1>(1)); + sycl::queue queue; + queue.submit([&](sycl::handler &cgh) { + sycl::accessor + B(b, cgh); // is not being simplified due to not deducable dimension + cgh.single_task([=]() { + auto B2 = B; + (int &)B2 = 399; + }); + }); + } + assert(data == 399); + } catch (sycl::exception e) { + std::cout << "SYCL exception caught: " << e.what(); + return 1; + } + } + + { + // Call every available accessor's constructor to ensure that they work with + // a buffer with a non-default allocator. + int data[] = {1, 2, 3}; + + using allocator_type = std::allocator; + + sycl::buffer buf1(&data[0], sycl::range<1>(1), + allocator_type{}); + sycl::buffer buf2(&data[1], sycl::range<1>(1), + allocator_type{}); + sycl::buffer buf3(&data[2], sycl::range<1>(1), + allocator_type{}); + + sycl::queue queue; + queue.submit([&](sycl::handler &cgh) { + sycl::accessor + acc1(buf1, cgh); // is not being simplified due to not deducable dimension + //sycl::accessor + // acc2(buf2, cgh); + sycl::accessor acc2(buf2, cgh); + //sycl::accessor + // acc3(buf3, cgh, sycl::range<1>(1)); + sycl::accessor acc3(buf3, cgh, sycl::range<1>(1)); + + cgh.single_task([=]() { + acc1 *= 2; + acc2[0] *= 2; + acc3[0] *= 2; + }); + }); + + sycl::accessor + acc4(buf1); // is not being simplified due to not deducable dimension + //sycl::accessor + // acc5(buf2); + sycl::host_accessor acc5(buf2, sycl::read_only); + //sycl::accessor + // acc6(buf3, sycl::range<1>(1)); + sycl::host_accessor acc6(buf3, sycl::range<1>(1), sycl::read_only); + + assert(acc4 == 2); + assert(acc5[0] == 4); + assert(acc6[0] == 6); + } + + // Constant buffer accessor + { + try { + int data = -1; + int cnst = 399; + + { + sycl::buffer d(&data, sycl::range<1>(1)); + sycl::buffer c(&cnst, sycl::range<1>(1)); + + sycl::queue queue; + queue.submit([&](sycl::handler &cgh) { + sycl::accessor D(d, cgh, sycl::write_only); + sycl::accessor C(c, cgh, sycl::read_constant); + + cgh.single_task([=]() { + D[0] = C[0]; + }); + }); + + assert(data == 399); + } + + } catch (sycl::exception e) { + std::cout << "SYCL exception caught: " << e.what(); + return 1; + } + } + + // Placeholder accessor + { + try { + int data = -1; + int cnst = 399; + + { + sycl::buffer d(&data, sycl::range<1>(1)); + sycl::buffer c(&cnst, sycl::range<1>(1)); + + sycl::accessor D(d, sycl::write_only); + sycl::accessor C(c, sycl::read_constant); + + sycl::queue queue; + queue.submit([&](sycl::handler &cgh) { + cgh.require(D); + cgh.require(C); + + cgh.single_task([=]() { + D[0] = C[0]; + }); + }); + + assert(data == 399); + } + + } catch (sycl::exception e) { + std::cout << "SYCL exception caught: " << e.what(); + return 1; + } + } +} From 6e37b7846199a956e3222d2fa928ecfcc60311bf Mon Sep 17 00:00:00 2001 From: iburylov Date: Mon, 8 Jun 2020 18:43:58 +0300 Subject: [PATCH 02/16] clang format fix Signed-off-by: iburylov --- sycl/include/CL/sycl/access/access.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index 2c28490c28708..9d24b8d3c92f7 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -77,7 +77,7 @@ enum class address_space : int { local_space }; -} // namespace access +} // namespace access namespace detail { From 7069dc728c8c74a6cb307c34114f67029cf5b407 Mon Sep 17 00:00:00 2001 From: iburylov Date: Tue, 9 Jun 2020 10:36:24 +0300 Subject: [PATCH 03/16] clang format fix Signed-off-by: iburylov --- sycl/include/CL/sycl/accessor.hpp | 32 +++++++++++++++---------------- 1 file changed, 16 insertions(+), 16 deletions(-) diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 9ab0b3fd1fdac..4f6cec10687bf 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -1239,41 +1239,41 @@ class accessor : template accessor(buffer, Ts...) - -> accessor; + ->accessor; template accessor(buffer, handler, Ts...) - -> accessor; + ->accessor; template accessor(buffer, Ts..., mode_tag_t, property_list = {}) - -> accessor; + ->accessor; template accessor(buffer, handler, Ts..., mode_tag_t, property_list = {}) - -> accessor; + ->accessor; template accessor(buffer, Ts..., mode_target_tag_t, property_list = {}) - -> accessor; + ->accessor; template accessor(buffer, handler, Ts..., mode_target_tag_t, property_list = {}) - -> accessor; + ->accessor; #endif @@ -1746,23 +1746,23 @@ class host_accessor template host_accessor(buffer, Ts...) - -> host_accessor; + ->host_accessor; template host_accessor(buffer, handler, Ts...) - -> host_accessor; + ->host_accessor; template host_accessor(buffer, Ts..., mode_tag_t, property_list = {}) - -> host_accessor; + ->host_accessor; template host_accessor(buffer, handler, Ts..., mode_tag_t, property_list = {}) - -> host_accessor; + ->host_accessor; #endif From 5c74c69f60222eabb7721b29ff85c014e5a5d271 Mon Sep 17 00:00:00 2001 From: iburylov Date: Tue, 9 Jun 2020 17:37:21 +0300 Subject: [PATCH 04/16] [SYCL] abi compatibility fix Signed-off-by: iburylov --- sycl/include/CL/sycl/access/access.hpp | 45 +++++++++++--------------- 1 file changed, 18 insertions(+), 27 deletions(-) diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index 9d24b8d3c92f7..dd87949dbbbf7 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -11,6 +11,7 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { +namespace access { enum class target { global_buffer = 2014, @@ -22,12 +23,6 @@ enum class target { image_array }; -// Backward compatibility namespace nesting -namespace access { -using sycl::target; -} - -namespace access { enum class mode { read = 1024, write, @@ -36,13 +31,26 @@ enum class mode { discard_read_write, atomic }; -} -using access_mode = access::mode; +enum class fence_space { + local_space, + global_space, + global_and_local +}; -namespace access { enum class placeholder { false_t, true_t }; -} + +enum class address_space : int { + private_space = 0, + global_space, + constant_space, + local_space +}; + +} // namespace access + +using access::target; +using access_mode = access::mode; #if __cplusplus > 201402L @@ -62,23 +70,6 @@ inline constexpr mode_target_tag_t #endif -namespace access { - -enum class fence_space { - local_space, - global_space, - global_and_local -}; - -enum class address_space : int { - private_space = 0, - global_space, - constant_space, - local_space -}; - -} // namespace access - namespace detail { constexpr bool isTargetHostAccess(access::target T) { From 05d6a9f5e739e11dfb8d589df3866ca8c122e810 Mon Sep 17 00:00:00 2001 From: Ilya Burylov Date: Tue, 16 Jun 2020 08:59:05 +0300 Subject: [PATCH 05/16] Update sycl/include/CL/sycl/property_list.hpp Fix comment in property list Co-authored-by: Sergey Semenov <43845535+sergey-semenov@users.noreply.github.com> --- sycl/include/CL/sycl/property_list.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/property_list.hpp b/sycl/include/CL/sycl/property_list.hpp index a7210f9199032..ec6e2ed01490d 100644 --- a/sycl/include/CL/sycl/property_list.hpp +++ b/sycl/include/CL/sycl/property_list.hpp @@ -152,7 +152,7 @@ RegisterProp(PropKind::BufferContextBound, buffer::context_bound); RegisterProp(PropKind::QueueEnableProfiling, queue::enable_profiling); RegisterProp(PropKind::InOrder, queue::in_order); -// Buffer +// Accessor RegisterProp(PropKind::NoInit, noinit); // Sentinel, needed for automatic build of tuple in property_list. From 59491ed7584e5c98858e0cca383deb83f2f3971a Mon Sep 17 00:00:00 2001 From: iburylov Date: Tue, 16 Jun 2020 10:11:10 +0300 Subject: [PATCH 06/16] merged tests into one file Signed-off-by: iburylov --- sycl/test/basic_tests/accessor/accessor.cpp | 76 ++++++++++++++++++++- 1 file changed, 74 insertions(+), 2 deletions(-) diff --git a/sycl/test/basic_tests/accessor/accessor.cpp b/sycl/test/basic_tests/accessor/accessor.cpp index a769df2f63003..ca5b793a21d57 100644 --- a/sycl/test/basic_tests/accessor/accessor.cpp +++ b/sycl/test/basic_tests/accessor/accessor.cpp @@ -1,8 +1,13 @@ // 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: 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 ----------------==// // @@ -74,9 +79,13 @@ 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); @@ -99,7 +108,11 @@ 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)); @@ -124,7 +137,11 @@ 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); @@ -147,7 +164,11 @@ 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(); }); @@ -175,7 +196,11 @@ 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(); @@ -202,14 +227,22 @@ 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); @@ -226,15 +259,23 @@ 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; @@ -250,7 +291,11 @@ 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) { @@ -280,8 +325,13 @@ 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( @@ -311,7 +361,11 @@ 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}; @@ -339,15 +393,24 @@ 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); @@ -400,13 +463,17 @@ 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; @@ -417,12 +484,17 @@ 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); From 5823419a2069eeb13bf90ce80c529a29303bd88a Mon Sep 17 00:00:00 2001 From: iburylov Date: Tue, 16 Jun 2020 11:18:15 +0300 Subject: [PATCH 07/16] added additional comments to the accessor test Signed-off-by: iburylov --- sycl/test/basic_tests/accessor/accessor.cpp | 7 +- .../accessor/accessor_simplification.cpp | 520 ------------------ 2 files changed, 5 insertions(+), 522 deletions(-) delete mode 100644 sycl/test/basic_tests/accessor/accessor_simplification.cpp diff --git a/sycl/test/basic_tests/accessor/accessor.cpp b/sycl/test/basic_tests/accessor/accessor.cpp index 21f269024817d..508a6bf3acdb5 100644 --- a/sycl/test/basic_tests/accessor/accessor.cpp +++ b/sycl/test/basic_tests/accessor/accessor.cpp @@ -84,9 +84,12 @@ int main() { assert(acc_src.get_count() == 2); assert(acc_src.get_range() == sycl::range<1>(2)); - // Make sure that operator[] is defined for both size_t and id<1>. + // operator[] overload for size_t was intentionally removed + // to remove ambiguity, when passing item to operator[]. // Implicit conversion from IdxSzT to size_t guarantees that no - // implicit conversion from size_t to id<1> will happen. + // implicit conversion from size_t to id<1> will happen, + // thus `acc_src[IdxSzT(0)]` will no longer compile. + // Replaced with acc_src[0]. assert(acc_src[0] + acc_src[IdxID1(1)] == 10); acc_dst[0] = acc_src[0] + acc_src[IdxID1(0)]; diff --git a/sycl/test/basic_tests/accessor/accessor_simplification.cpp b/sycl/test/basic_tests/accessor/accessor_simplification.cpp deleted file mode 100644 index fa978acf4be14..0000000000000 --- a/sycl/test/basic_tests/accessor/accessor_simplification.cpp +++ /dev/null @@ -1,520 +0,0 @@ -// RUN: %clangxx -fsycl -std=c++17 %s -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 -//==---------accessor_simplification.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 - -namespace sycl { -using namespace cl::sycl; -} - -struct IdxID1 { - int x; - - IdxID1(int x) : x(x) {} - operator sycl::id<1>() { return x; } -}; - -struct IdxID3 { - int x; - int y; - int z; - - IdxID3(int x, int y, int z) : x(x), y(y), z(z) {} - operator sycl::id<3>() { return sycl::id<3>(x, y, z); } -}; - -struct IdxSzT { - int x; - - IdxSzT(int x) : x(x) {} - operator size_t() { return x; } -}; - -template -struct AccWrapper { Acc accessor; }; - -template -struct AccsWrapper { - int a; - Acc1 accessor1; - int b; - Acc2 accessor2; -}; - -struct Wrapper1 { - int a; - int b; -}; - -template -struct Wrapper2 { - Wrapper1 w1; - AccWrapper wrapped; -}; - -template -struct Wrapper3 { Wrapper2 w2; }; - -int main() { - // Host accessor. - { - int src[2] = {3, 7}; - int dst[2]; - - sycl::buffer buf_src(src, sycl::range<1>(2), - {cl::sycl::property::buffer::use_host_ptr()}); - sycl::buffer buf_dst(dst, sycl::range<1>(2), - {cl::sycl::property::buffer::use_host_ptr()}); - - sycl::id<1> id1(1); - //auto acc_src = buf_src.get_access(); - sycl::host_accessor acc_src(buf_src, sycl::read_only); - //auto acc_dst = buf_dst.get_access(); - sycl::host_accessor acc_dst(buf_dst); - - assert(!acc_src.is_placeholder()); - assert(acc_src.get_size() == sizeof(src)); - assert(acc_src.get_count() == 2); - assert(acc_src.get_range() == sycl::range<1>(2)); - - // Make sure that operator[] is defined for both size_t and id<1>. - // Implicit conversion from IdxSzT to size_t guarantees that no - // implicit conversion from size_t to id<1> will happen. - assert(acc_src[IdxSzT(0)] + acc_src[IdxID1(1)] == 10); - - acc_dst[0] = acc_src[0] + acc_src[IdxID1(0)]; - acc_dst[id1] = acc_src[1] + acc_src[IdxSzT(1)]; - assert(dst[0] == 6 && dst[1] == 14); - } - - // Three-dimensional host accessor. - { - int data[24]; - for (int i = 0; i < 24; ++i) - data[i] = i; - { - sycl::buffer buf(data, sycl::range<3>(2, 3, 4)); - //auto acc = buf.get_access(); - sycl::host_accessor acc(buf); - - assert(!acc.is_placeholder()); - assert(acc.get_size() == sizeof(data)); - assert(acc.get_count() == 24); - assert(acc.get_range() == sycl::range<3>(2, 3, 4)); - - for (int i = 0; i < 2; ++i) - for (int j = 0; j < 3; ++j) - for (int k = 0; k < 4; ++k) - acc[IdxID3(i, j, k)] += acc[sycl::id<3>(i, j, k)]; - } - for (int i = 0; i < 24; ++i) { - assert(data[i] == 2 * i); - } - } - int data = 5; - // Device accessor. - { - sycl::queue Queue; - - sycl::buffer buf(&data, sycl::range<1>(1), - {cl::sycl::property::buffer::use_host_ptr()}); - - Queue.submit([&](sycl::handler &cgh) { - //auto acc = buf.get_access(cgh); - sycl::accessor acc(buf, cgh); - assert(!acc.is_placeholder()); - assert(acc.get_size() == sizeof(int)); - assert(acc.get_count() == 1); - assert(acc.get_range() == sycl::range<1>(1)); - cgh.single_task( - [=]() { acc[IdxSzT(0)] += acc[IdxID1(0)]; }); - }); - Queue.wait(); - } - assert(data == 10); - - // Device accessor with 2-dimensional subscript operators. - { - sycl::queue Queue; - if (!Queue.is_host()) { - int array[2][3] = {0}; - { - sycl::range<2> Range(2, 3); - sycl::buffer buf((int *)array, Range, - {cl::sycl::property::buffer::use_host_ptr()}); - - Queue.submit([&](sycl::handler &cgh) { - //auto acc = buf.get_access(cgh); - sycl::accessor acc(buf, cgh); - cgh.parallel_for(Range, [=](sycl::item<2> itemID) { - acc[itemID.get_id(0)][itemID.get_id(1)] += itemID.get_linear_id(); - }); - }); - Queue.wait(); - } - for (int i = 0; i < 2; i++) { - for (int j = 0; j < 3; j++) { - std::cout << "array[" << i << "][" << j << "]=" << array[i][j] - << std::endl; - assert(array[i][j] == i * 3 + j); - } - } - } - } - - // Device accessor with 3-dimensional subscript operators. - { - sycl::queue Queue; - if (!Queue.is_host()) { - int array[2][3][4] = {0}; - { - sycl::range<3> Range(2, 3, 4); - sycl::buffer buf((int *)array, Range, - {cl::sycl::property::buffer::use_host_ptr()}); - - Queue.submit([&](sycl::handler &cgh) { - //auto acc = buf.get_access(cgh); - sycl::accessor acc(buf, cgh); - 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(); - }); - }); - Queue.wait(); - } - for (int i = 0; i < 2; i++) { - for (int j = 0; j < 3; j++) { - for (int k = 0; k < 4; k++) { - std::cout << "array[" << i << "][" << j << "][" << k - << "]=" << array[i][j][k] << std::endl; - assert(array[i][j][k] == k + 4 * (j + 3 * i)); - } - } - } - } - } - - // Discard write accessor. - { - try { - sycl::queue Queue; - sycl::buffer buf(sycl::range<1>(3)); - - Queue.submit([&](sycl::handler &cgh) { - //auto dev_acc = buf.get_access(cgh); - sycl::accessor dev_acc(buf, cgh, sycl::noinit); - - cgh.parallel_for( - sycl::range<1>{3}, - [=](sycl::id<1> index) { dev_acc[index] = 42; }); - }); - - //auto host_acc = buf.get_access(); - sycl::host_accessor host_acc(buf, sycl::read_only); // TODO: it is read_write now - is there a reason to have read only accessor? - - for (int i = 0; i != 3; ++i) - assert(host_acc[i] == 42); - - } catch (cl::sycl::exception e) { - std::cout << "SYCL exception caught: " << e.what(); - return 1; - } - } - - // Discard read-write accessor. - { - try { - sycl::queue Queue; - sycl::buffer buf(sycl::range<1>(3)); - - Queue.submit([&](sycl::handler &cgh) { - //auto dev_acc = buf.get_access(cgh); - sycl::accessor dev_acc(buf, cgh, sycl::write_only); - - cgh.parallel_for( - sycl::range<1>{3}, - [=](sycl::id<1> index) { dev_acc[index] = 42; }); - }); - - //auto host_acc = buf.get_access(); - sycl::host_accessor host_acc(buf, sycl::noinit); - } catch (cl::sycl::exception e) { - std::cout << "SYCL exception caught: " << e.what(); - return 1; - } - } - - // Check that accessor is initialized when accessor is wrapped to some class. - { - sycl::queue queue; - if (!queue.is_host()) { - int array[10] = {0}; - { - sycl::buffer buf((int *)array, sycl::range<1>(10), - {cl::sycl::property::buffer::use_host_ptr()}); - queue.submit([&](sycl::handler &cgh) { - //auto acc = buf.get_access(cgh); - sycl::accessor acc(buf, cgh); - auto acc_wrapped = AccWrapper{acc}; - cgh.parallel_for( - sycl::range<1>(buf.get_count()), [=](sycl::item<1> it) { - auto idx = it.get_linear_id(); - acc_wrapped.accessor[idx] = 333; - }); - }); - queue.wait(); - } - for (int i = 0; i < 10; i++) { - std::cout << "array[" << i << "]=" << array[i] << std::endl; - assert(array[i] == 333); - } - } - } - - // Case when several accessors are wrapped to some class. Check that they are - // initialized in proper way and value is assigned. - { - sycl::queue queue; - if (!queue.is_host()) { - int array1[10] = {0}; - int array2[10] = {0}; - { - sycl::buffer buf1((int *)array1, sycl::range<1>(10), - {cl::sycl::property::buffer::use_host_ptr()}); - sycl::buffer buf2((int *)array2, sycl::range<1>(10), - {cl::sycl::property::buffer::use_host_ptr()}); - queue.submit([&](sycl::handler &cgh) { - //auto acc1 = buf1.get_access(cgh); - sycl::accessor acc1(buf1, cgh); - //auto acc2 = buf2.get_access(cgh); - sycl::accessor acc2(buf2, cgh); - auto acc_wrapped = - AccsWrapper{10, acc1, 5, acc2}; - cgh.parallel_for( - sycl::range<1>(10), [=](sycl::item<1> it) { - auto idx = it.get_linear_id(); - acc_wrapped.accessor1[idx] = 333; - acc_wrapped.accessor2[idx] = 666; - }); - }); - queue.wait(); - } - for (int i = 0; i < 10; i++) { - std::cout << "array1[" << i << "]=" << array1[i] << std::endl; - std::cout << "array2[" << i << "]=" << array2[i] << std::endl; - assert(array1[i] == 333); - assert(array2[i] == 666); - } - } - } - - // Several levels of wrappers for accessor. - { - sycl::queue queue; - if (!queue.is_host()) { - int array[10] = {0}; - { - sycl::buffer buf((int *)array, sycl::range<1>(10), - {cl::sycl::property::buffer::use_host_ptr()}); - queue.submit([&](sycl::handler &cgh) { - //auto acc = buf.get_access(cgh); - sycl::accessor acc(buf, cgh); - auto acc_wrapped = AccWrapper{acc}; - Wrapper1 wr1; - auto wr2 = Wrapper2{wr1, acc_wrapped}; - auto wr3 = Wrapper3{wr2}; - cgh.parallel_for( - sycl::range<1>(buf.get_count()), [=](sycl::item<1> it) { - auto idx = it.get_linear_id(); - wr3.w2.wrapped.accessor[idx] = 333; - }); - }); - queue.wait(); - } - for (int i = 0; i < 10; i++) { - std::cout << "array[" << i << "]=" << array[i] << std::endl; - assert(array[i] == 333); - } - } - } - - // Two accessors to the same buffer. - { - try { - sycl::queue queue; - int array[3] = {1, 1, 1}; - sycl::buffer buf(array, sycl::range<1>(3)); - - std::cout << "We are here 1" << std::endl; - - queue.submit([&](sycl::handler &cgh) { - //auto acc1 = buf.get_access(cgh); - sycl::accessor acc1(buf, cgh, sycl::read_only); - //auto acc2 = buf.get_access(cgh); - sycl::accessor acc2(buf, cgh); - - cgh.parallel_for( - sycl::range<1>{3}, - [=](sycl::id<1> index) { - acc2[index] = 41 + acc1[index]; - }); - }); - - //auto host_acc = buf.get_access(); - sycl::host_accessor host_acc(buf, sycl::read_only); - for (int i = 0; i != 3; ++i) - assert(host_acc[i] == 42); - - } catch (cl::sycl::exception e) { - std::cout << "SYCL exception caught: " << e.what(); - return 1; - } - } - - // Accessor with dimensionality 0. - { - try { - int data = -1; - { - sycl::buffer b(&data, sycl::range<1>(1)); - sycl::queue queue; - queue.submit([&](sycl::handler &cgh) { - sycl::accessor - B(b, cgh); // is not being simplified due to not deducable dimension - cgh.single_task([=]() { - auto B2 = B; - (int &)B2 = 399; - }); - }); - } - assert(data == 399); - } catch (sycl::exception e) { - std::cout << "SYCL exception caught: " << e.what(); - return 1; - } - } - - { - // Call every available accessor's constructor to ensure that they work with - // a buffer with a non-default allocator. - int data[] = {1, 2, 3}; - - using allocator_type = std::allocator; - - sycl::buffer buf1(&data[0], sycl::range<1>(1), - allocator_type{}); - sycl::buffer buf2(&data[1], sycl::range<1>(1), - allocator_type{}); - sycl::buffer buf3(&data[2], sycl::range<1>(1), - allocator_type{}); - - sycl::queue queue; - queue.submit([&](sycl::handler &cgh) { - sycl::accessor - acc1(buf1, cgh); // is not being simplified due to not deducable dimension - //sycl::accessor - // acc2(buf2, cgh); - sycl::accessor acc2(buf2, cgh); - //sycl::accessor - // acc3(buf3, cgh, sycl::range<1>(1)); - sycl::accessor acc3(buf3, cgh, sycl::range<1>(1)); - - cgh.single_task([=]() { - acc1 *= 2; - acc2[0] *= 2; - acc3[0] *= 2; - }); - }); - - sycl::accessor - acc4(buf1); // is not being simplified due to not deducable dimension - //sycl::accessor - // acc5(buf2); - sycl::host_accessor acc5(buf2, sycl::read_only); - //sycl::accessor - // acc6(buf3, sycl::range<1>(1)); - sycl::host_accessor acc6(buf3, sycl::range<1>(1), sycl::read_only); - - assert(acc4 == 2); - assert(acc5[0] == 4); - assert(acc6[0] == 6); - } - - // Constant buffer accessor - { - try { - int data = -1; - int cnst = 399; - - { - sycl::buffer d(&data, sycl::range<1>(1)); - sycl::buffer c(&cnst, sycl::range<1>(1)); - - sycl::queue queue; - queue.submit([&](sycl::handler &cgh) { - sycl::accessor D(d, cgh, sycl::write_only); - sycl::accessor C(c, cgh, sycl::read_constant); - - cgh.single_task([=]() { - D[0] = C[0]; - }); - }); - - assert(data == 399); - } - - } catch (sycl::exception e) { - std::cout << "SYCL exception caught: " << e.what(); - return 1; - } - } - - // Placeholder accessor - { - try { - int data = -1; - int cnst = 399; - - { - sycl::buffer d(&data, sycl::range<1>(1)); - sycl::buffer c(&cnst, sycl::range<1>(1)); - - sycl::accessor D(d, sycl::write_only); - sycl::accessor C(c, sycl::read_constant); - - sycl::queue queue; - queue.submit([&](sycl::handler &cgh) { - cgh.require(D); - cgh.require(C); - - cgh.single_task([=]() { - D[0] = C[0]; - }); - }); - - assert(data == 399); - } - - } catch (sycl::exception e) { - std::cout << "SYCL exception caught: " << e.what(); - return 1; - } - } -} From d0facb889c2ee59352a32cf8f350fc2fdbe7330e Mon Sep 17 00:00:00 2001 From: iburylov Date: Tue, 16 Jun 2020 11:58:50 +0300 Subject: [PATCH 08/16] added back checks missed when merged files Signed-off-by: iburylov --- sycl/test/basic_tests/accessor/accessor.cpp | 64 +++++++++++++++++++++ 1 file changed, 64 insertions(+) diff --git a/sycl/test/basic_tests/accessor/accessor.cpp b/sycl/test/basic_tests/accessor/accessor.cpp index 508a6bf3acdb5..2ff0258333411 100644 --- a/sycl/test/basic_tests/accessor/accessor.cpp +++ b/sycl/test/basic_tests/accessor/accessor.cpp @@ -496,4 +496,68 @@ int main() { assert(acc5[0] == 4); assert(acc6[0] == 6); } + +#ifdef simplification_test + // Constant buffer accessor + { + try { + int data = -1; + int cnst = 399; + + { + sycl::buffer d(&data, sycl::range<1>(1)); + sycl::buffer c(&cnst, sycl::range<1>(1)); + + sycl::queue queue; + queue.submit([&](sycl::handler &cgh) { + sycl::accessor D(d, cgh, sycl::write_only); + sycl::accessor C(c, cgh, sycl::read_constant); + + cgh.single_task([=]() { + D[0] = C[0]; + }); + }); + + assert(data == 399); + } + + } catch (sycl::exception e) { + std::cout << "SYCL exception caught: " << e.what(); + return 1; + } + } + + // Placeholder accessor + { + try { + int data = -1; + int cnst = 399; + + { + sycl::buffer d(&data, sycl::range<1>(1)); + sycl::buffer c(&cnst, sycl::range<1>(1)); + + sycl::accessor D(d, sycl::write_only); + sycl::accessor C(c, sycl::read_constant); + + sycl::queue queue; + queue.submit([&](sycl::handler &cgh) { + cgh.require(D); + cgh.require(C); + + cgh.single_task([=]() { + D[0] = C[0]; + }); + }); + + assert(data == 399); + } + + } catch (sycl::exception e) { + std::cout << "SYCL exception caught: " << e.what(); + return 1; + } + } +#endif + } From d846804516bd2a376310b9d4ad35fd6c6a3eb2a9 Mon Sep 17 00:00:00 2001 From: iburylov Date: Tue, 16 Jun 2020 12:27:44 +0300 Subject: [PATCH 09/16] clang format fix Signed-off-by: iburylov --- sycl/test/basic_tests/accessor/accessor.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/test/basic_tests/accessor/accessor.cpp b/sycl/test/basic_tests/accessor/accessor.cpp index 2ff0258333411..b7610d1f7a384 100644 --- a/sycl/test/basic_tests/accessor/accessor.cpp +++ b/sycl/test/basic_tests/accessor/accessor.cpp @@ -559,5 +559,4 @@ int main() { } } #endif - } From a44eb22ec60e402caa69ae7dea9b734c173ed5fe Mon Sep 17 00:00:00 2001 From: iburylov Date: Tue, 16 Jun 2020 18:17:49 +0300 Subject: [PATCH 10/16] made tags visible in C++11, removed placeholder ctors, extended testing Signed-off-by: iburylov --- sycl/include/CL/sycl/access/access.hpp | 20 +++++- sycl/include/CL/sycl/accessor.hpp | 75 +-------------------- sycl/include/CL/sycl/detail/common.hpp | 6 ++ sycl/include/CL/sycl/property_list.hpp | 10 +++ sycl/test/basic_tests/accessor/accessor.cpp | 22 +++++- 5 files changed, 56 insertions(+), 77 deletions(-) diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index dd87949dbbbf7..74ac7d8222855 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #pragma once +#include #include __SYCL_INLINE_NAMESPACE(cl) { @@ -52,8 +53,6 @@ enum class address_space : int { using access::target; using access_mode = access::mode; -#if __cplusplus > 201402L - template struct mode_tag_t { explicit mode_tag_t() = default; }; @@ -62,12 +61,29 @@ template struct mode_target_tag_t { explicit mode_target_tag_t() = default; }; +#if __cplusplus > 201402L + inline constexpr mode_tag_t read_only{}; inline constexpr mode_tag_t read_write{}; inline constexpr mode_tag_t write_only{}; inline constexpr mode_target_tag_t read_constant{}; +#else + +namespace { + +constexpr const auto& read_only = + sycl::detail::InlineVariableHelper>::value; +constexpr const auto& read_write = + sycl::detail::InlineVariableHelper>::value; +constexpr const auto& write_only = + sycl::detail::InlineVariableHelper>::value; +constexpr const auto& read_constant = + sycl::detail::InlineVariableHelper>::value; + +} + #endif namespace detail { diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index edf425213c0e3..8e9ebf4e2cf7d 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -1593,6 +1593,8 @@ class host_accessor // buffer | handler | range | id | | property_list // buffer | handler | range | id | mode_tag | property_list // -------+---------+-------+----+----------+-------------- + // host_accessor with handler argument will be added later + // to facilitate non-blocking accessor use case template ::value && Dims == 0>> - host_accessor(buffer &BufferRef, - handler &CommandGroupHandler, - const property_list &PropertyList = {}) - : AccessorT(BufferRef, PropertyList) { - (void)CommandGroupHandler; - } - template ()>> host_accessor(buffer &BufferRef, @@ -1627,25 +1619,6 @@ class host_accessor #endif - template ()>> - host_accessor(buffer &BufferRef, - handler &CommandGroupHandler, - const property_list &PropertyList = {}) - : AccessorT(BufferRef, PropertyList) { - (void)CommandGroupHandler; - } - -#if __cplusplus > 201402L - - template ()>> - host_accessor(buffer &BufferRef, - handler &CommandGroupHandler, mode_tag_t, - const property_list &PropertyList = {}) - : host_accessor(BufferRef, CommandGroupHandler, PropertyList) {} - -#endif template ()>> @@ -1663,27 +1636,6 @@ class host_accessor const property_list &PropertyList = {}) : host_accessor(BufferRef, AccessRange, {}, PropertyList) {} -#endif - - template ()>> - host_accessor(buffer &BufferRef, - handler &CommandGroupHandler, range AccessRange, - const property_list &PropertyList = {}) - : AccessorT(BufferRef, AccessRange, {}, PropertyList) { - (void)CommandGroupHandler; - } - -#if __cplusplus > 201402L - - template ()>> - host_accessor(buffer &BufferRef, - handler &CommandGroupHandler, range AccessRange, - mode_tag_t, const property_list &PropertyList = {}) - : host_accessor(BufferRef, CommandGroupHandler, AccessRange, {}, - PropertyList) {} - #endif template , const property_list &PropertyList = {}) : host_accessor(BufferRef, AccessRange, AccessOffset, PropertyList) {} -#endif - - template ()>> - host_accessor(buffer &BufferRef, - handler &CommandGroupHandler, range AccessRange, - id AccessOffset, - const property_list &PropertyList = {}) - : AccessorT(BufferRef, AccessRange, AccessOffset, PropertyList) { - (void)CommandGroupHandler; - } - -#if __cplusplus > 201402L - - template ()>> - host_accessor(buffer &BufferRef, - handler &CommandGroupHandler, range AccessRange, - id AccessOffset, mode_tag_t, - const property_list &PropertyList = {}) - : host_accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset, - PropertyList) {} - #endif }; diff --git a/sycl/include/CL/sycl/detail/common.hpp b/sycl/include/CL/sycl/detail/common.hpp index e5d70adb7e829..ffac3e2356872 100644 --- a/sycl/include/CL/sycl/detail/common.hpp +++ b/sycl/include/CL/sycl/detail/common.hpp @@ -307,6 +307,12 @@ using KernelSetId = size_t; constexpr KernelSetId SpvFileKSId = 0; constexpr KernelSetId LastKSId = SpvFileKSId; +template struct InlineVariableHelper { + static constexpr T value{}; +}; + +template +constexpr T InlineVariableHelper::value; } // namespace detail } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/CL/sycl/property_list.hpp b/sycl/include/CL/sycl/property_list.hpp index ec6e2ed01490d..7013dc8e4dff2 100644 --- a/sycl/include/CL/sycl/property_list.hpp +++ b/sycl/include/CL/sycl/property_list.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include #include @@ -228,6 +229,15 @@ class noinit : public detail::Prop {}; inline constexpr property::noinit noinit; +#else + +namespace { + +constexpr const auto& noinit = + sycl::detail::InlineVariableHelper::value; + +} + #endif class property_list { diff --git a/sycl/test/basic_tests/accessor/accessor.cpp b/sycl/test/basic_tests/accessor/accessor.cpp index b7610d1f7a384..af9c624f7e5ab 100644 --- a/sycl/test/basic_tests/accessor/accessor.cpp +++ b/sycl/test/basic_tests/accessor/accessor.cpp @@ -497,7 +497,6 @@ int main() { assert(acc6[0] == 6); } -#ifdef simplification_test // Constant buffer accessor { try { @@ -510,8 +509,17 @@ int main() { sycl::queue queue; queue.submit([&](sycl::handler &cgh) { +#ifdef 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]; @@ -537,8 +545,19 @@ int main() { sycl::buffer d(&data, sycl::range<1>(1)); sycl::buffer c(&cnst, sycl::range<1>(1)); +#ifdef simplification_test + sycl::accessor + D(d); + sycl::accessor + 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) { @@ -558,5 +577,4 @@ int main() { return 1; } } -#endif } From 81eb25de9b2aa2072ba5c8547206e7e57900cc5b Mon Sep 17 00:00:00 2001 From: iburylov Date: Tue, 16 Jun 2020 18:34:17 +0300 Subject: [PATCH 11/16] clang format Signed-off-by: iburylov --- sycl/include/CL/sycl/access/access.hpp | 20 ++++++++++---------- sycl/include/CL/sycl/accessor.hpp | 1 - sycl/include/CL/sycl/detail/common.hpp | 3 +-- sycl/include/CL/sycl/property_list.hpp | 4 ++-- sycl/test/basic_tests/accessor/accessor.cpp | 16 ++++++++-------- 5 files changed, 21 insertions(+), 23 deletions(-) diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index 74ac7d8222855..19f2f541429a2 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -73,16 +73,16 @@ inline constexpr mode_target_tag_t namespace { -constexpr const auto& read_only = - sycl::detail::InlineVariableHelper>::value; -constexpr const auto& read_write = - sycl::detail::InlineVariableHelper>::value; -constexpr const auto& write_only = - sycl::detail::InlineVariableHelper>::value; -constexpr const auto& read_constant = - sycl::detail::InlineVariableHelper>::value; - -} +constexpr const auto &read_only = + sycl::detail::InlineVariableHelper>::value; +constexpr const auto &read_write = + sycl::detail::InlineVariableHelper>::value; +constexpr const auto &write_only = + sycl::detail::InlineVariableHelper>::value; +constexpr const auto &read_constant = + sycl::detail::InlineVariableHelper>::value; + +} // namespace #endif diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 8e9ebf4e2cf7d..0dac1b4fb27ab 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -1619,7 +1619,6 @@ class host_accessor #endif - template ()>> host_accessor(buffer &BufferRef, diff --git a/sycl/include/CL/sycl/detail/common.hpp b/sycl/include/CL/sycl/detail/common.hpp index ffac3e2356872..11bb8f395a7c8 100644 --- a/sycl/include/CL/sycl/detail/common.hpp +++ b/sycl/include/CL/sycl/detail/common.hpp @@ -311,8 +311,7 @@ template struct InlineVariableHelper { static constexpr T value{}; }; -template -constexpr T InlineVariableHelper::value; +template constexpr T InlineVariableHelper::value; } // namespace detail } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/CL/sycl/property_list.hpp b/sycl/include/CL/sycl/property_list.hpp index 7013dc8e4dff2..90aea2d60164a 100644 --- a/sycl/include/CL/sycl/property_list.hpp +++ b/sycl/include/CL/sycl/property_list.hpp @@ -233,8 +233,8 @@ inline constexpr property::noinit noinit; namespace { -constexpr const auto& noinit = - sycl::detail::InlineVariableHelper::value; +constexpr const auto &noinit = + sycl::detail::InlineVariableHelper::value; } diff --git a/sycl/test/basic_tests/accessor/accessor.cpp b/sycl/test/basic_tests/accessor/accessor.cpp index af9c624f7e5ab..701f23839d775 100644 --- a/sycl/test/basic_tests/accessor/accessor.cpp +++ b/sycl/test/basic_tests/accessor/accessor.cpp @@ -512,10 +512,10 @@ int main() { #ifdef simplification_test sycl::accessor - D(d, cgh); + D(d, cgh); sycl::accessor - C(c, cgh); + C(c, cgh); #else sycl::accessor D(d, cgh, sycl::write_only); sycl::accessor C(c, cgh, sycl::read_constant); @@ -546,13 +546,13 @@ int main() { sycl::buffer c(&cnst, sycl::range<1>(1)); #ifdef simplification_test - sycl::accessor + sycl::accessor D(d); - sycl::accessor + sycl::accessor C(c); #else sycl::accessor D(d, sycl::write_only); From 765a23a381226aba283ab2976f0cf62a0098cc49 Mon Sep 17 00:00:00 2001 From: iburylov Date: Tue, 16 Jun 2020 18:53:42 +0300 Subject: [PATCH 12/16] a little bit more of clang-format Signed-off-by: iburylov --- sycl/include/CL/sycl/access/access.hpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index 19f2f541429a2..10101d02435f5 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -75,12 +75,12 @@ namespace { constexpr const auto &read_only = sycl::detail::InlineVariableHelper>::value; -constexpr const auto &read_write = - sycl::detail::InlineVariableHelper>::value; +constexpr const auto &read_write = sycl::detail::InlineVariableHelper< + mode_tag_t>::value; constexpr const auto &write_only = sycl::detail::InlineVariableHelper>::value; -constexpr const auto &read_constant = - sycl::detail::InlineVariableHelper>::value; +constexpr const auto &read_constant = sycl::detail::InlineVariableHelper< + mode_target_tag_t>::value; } // namespace From cab9ad4b8cb5ceef20d0f6521f0ddcf8416a9af3 Mon Sep 17 00:00:00 2001 From: iburylov Date: Wed, 17 Jun 2020 15:08:19 +0300 Subject: [PATCH 13/16] removed not needed deduction guides Signed-off-by: iburylov --- sycl/include/CL/sycl/accessor.hpp | 12 +----------- 1 file changed, 1 insertion(+), 11 deletions(-) diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 0dac1b4fb27ab..bc82bfd653b32 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -1514,7 +1514,7 @@ class accessor @@ -1662,22 +1662,12 @@ template host_accessor(buffer, Ts...) ->host_accessor; -template -host_accessor(buffer, handler, Ts...) - ->host_accessor; - template host_accessor(buffer, Ts..., mode_tag_t, property_list = {}) ->host_accessor; -template -host_accessor(buffer, handler, Ts..., - mode_tag_t, property_list = {}) - ->host_accessor; - #endif } // namespace sycl From 6f19112d024ad2118205921e8f2c4b2332ab0126 Mon Sep 17 00:00:00 2001 From: iburylov Date: Wed, 17 Jun 2020 15:12:30 +0300 Subject: [PATCH 14/16] clang format fix Signed-off-by: iburylov --- sycl/include/CL/sycl/accessor.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index bc82bfd653b32..ad8e1cf4f2571 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -1514,7 +1514,7 @@ class accessor From e3f3bb6d57cea2f12b38ea2fc1ddac4be9b47b60 Mon Sep 17 00:00:00 2001 From: iburylov Date: Wed, 17 Jun 2020 15:56:50 +0300 Subject: [PATCH 15/16] test fix Signed-off-by: iburylov --- sycl/test/basic_tests/accessor/accessor.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test/basic_tests/accessor/accessor.cpp b/sycl/test/basic_tests/accessor/accessor.cpp index 701f23839d775..ca3b29094803a 100644 --- a/sycl/test/basic_tests/accessor/accessor.cpp +++ b/sycl/test/basic_tests/accessor/accessor.cpp @@ -509,7 +509,7 @@ int main() { sycl::queue queue; queue.submit([&](sycl::handler &cgh) { -#ifdef simplification_test +#ifndef simplification_test sycl::accessor D(d, cgh); @@ -545,7 +545,7 @@ int main() { sycl::buffer d(&data, sycl::range<1>(1)); sycl::buffer c(&cnst, sycl::range<1>(1)); -#ifdef simplification_test +#ifndef simplification_test sycl::accessor From 3677bbf56988f04d53876015d0c07d015b1fcb62 Mon Sep 17 00:00:00 2001 From: iburylov Date: Wed, 17 Jun 2020 16:50:36 +0300 Subject: [PATCH 16/16] ensured data returned back to host in test Signed-off-by: iburylov --- sycl/test/basic_tests/accessor/accessor.cpp | 14 ++++++++++++-- 1 file changed, 12 insertions(+), 2 deletions(-) diff --git a/sycl/test/basic_tests/accessor/accessor.cpp b/sycl/test/basic_tests/accessor/accessor.cpp index ca3b29094803a..b5ac254e91966 100644 --- a/sycl/test/basic_tests/accessor/accessor.cpp +++ b/sycl/test/basic_tests/accessor/accessor.cpp @@ -526,7 +526,12 @@ int main() { }); }); - assert(data == 399); +#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); } } catch (sycl::exception e) { @@ -569,7 +574,12 @@ int main() { }); }); - assert(data == 399); +#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); } } catch (sycl::exception e) {