Skip to content

Commit

Permalink
Add changes based on implementation feedback.
Browse files Browse the repository at this point in the history 
* Add fix for typo.
* Alter async_work_group functions overloads to correctly allow src
parameter to be a multi_ptr<T> and convert to multi_ptr<const T>.
* Switch get_pointer from using accessor_ptr as this includes SYCL 2020
change which aliases to multi_ptr<const T> when the accessor is
read_only.
* Split accessor::get_pointer into two definitions; one for
device/global_buffer which returns a multi_ptr and one for host_task
which returns a raw pointer.

Co-authored-by: Greg Lueck <gregory.m.lueck@intel.com>
  • Loading branch information
@AerialMantis @gmlueck
AerialMantis and gmlueck committed Jun 30, 2023
1 parent ad4a037 commit 99c0b3b
Show file tree
Hide file tree
Showing 5 changed files with 145 additions and 100 deletions.
170 changes: 98 additions & 72 deletions adoc/chapters/programming_interface.adoc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6962,17 +6962,32 @@ accessor was constructed. For other accessors, returns the default constructed
a@
[source]
----
accessor_ptr<access::decorated::legacy> get_pointer() const noexcept
global_ptr<access::decorated::legacy> get_pointer() const noexcept
----
a@ Returns a [code]#multi_ptr# to the start of this accessor's underlying
buffer, even if this is a <<ranged-accessor>> whose range does not start
at the beginning of the buffer. The return value is unspecified if the
accessor is empty.
a@ Available only when [code]#(AccessTarget == target::device)#.

Returns a [code]#multi_ptr# to the start of this accessor's underlying buffer,
even if this is a <<ranged-accessor>> whose range does not start at the
beginning of the buffer. The return value is unspecified if the accessor is
empty.

This function may only be called from within a <<command>>.

Deprecated in SYCL 2020. Use [code]#get_multi_ptr# instead.

a@
[source]
----
std::add_pointer_t<value_type> get_pointer() const noexcept
----
a@ Available only when [code]#(AccessTarget == target::host_task)#.

Returns a pointer to the start of this accessor's underlying buffer, even if
this is a <<ranged-accessor>> whose range does not start at the beginning of
the buffer. The return value is unspecified if the accessor is empty.

This function may only be called from within a <<command>>.

a@
[source]
----
Expand Down Expand Up @@ -8087,7 +8102,7 @@ a@
----
std::add_pointer_t<value_type> get_pointer() const noexcept
----
a@ Returns a [code]#multi_ptr# to the start of this accessor's underlying
a@ Returns a pointer to the start of this accessor's underlying
buffer, even if this is a <<ranged-accessor>> whose range does not start
at the beginning of the buffer. The return value is unspecified if the
accessor is empty.
Expand Down Expand Up @@ -8299,7 +8314,7 @@ void swap(local_accessor& other);
@
[source]
----
accessor_ptr<access::decorated::legacy> get_pointer() const noexcept
local_ptr<DataT> get_pointer() const noexcept
----
a@ Returns a [code]#multi_ptr# to the start of this accessor's local memory
region which corresponds to the calling work-group. The return value is
Expand Down Expand Up @@ -11913,56 +11928,70 @@ device_event async_work_group_copy(global_ptr<DataT> dest,
a@
[source]
----
template <typename DataT>
device_event async_work_group_copy(decorated_local_ptr<DataT> dest,
decorated_global_ptr<const DataT> src,
template <typename DestDataT, typename SrcDataT>
device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
decorated_global_ptr<SrcDataT> src,
size_t numElements) const
----
a@ Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements# from the source pointer [code]#src# to destination
pointer [code]#dest# and returns a SYCL [code]#device_event#
which can be used to wait on the completion of the copy.
a@ Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true)

Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements#
from the source pointer [code]#src# to destination pointer [code]#dest# and
returns a SYCL [code]#device_event# which can be used to wait on the completion
of the copy.

a@
[source]
----
template <typename DataT>
device_event async_work_group_copy(decorated_global_ptr<DataT> dest,
decorated_local_ptr<const DataT> src,
template <typename DestDataT, typename SrcDataT>
device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
decorated_local_ptr<SrcDataT> src,
size_t numElements) const
----
a@ Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements# from the source pointer [code]#src# to destination
pointer [code]#dest# and returns a SYCL [code]#device_event#
which can be used to wait on the completion of the copy.
a@ Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true)

Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements#
from the source pointer [code]#src# to destination pointer [code]#dest# and
returns a SYCL [code]#device_event# which can be used to wait on the
completion of the copy.

a@
[source]
----
template <typename DataT>
device_event async_work_group_copy(decorated_local_ptr<DataT> dest,
decorated_global_ptr<const DataT> src,
template <typename DestDataT, typename SrcDataT>
device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
decorated_global_ptr<SrcDataT> src,
size_t numElements, size_t srcStride) const
----
a@ Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements# from the source pointer [code]#src# to destination
pointer [code]#dest# with a source stride specified by
[code]#srcStride# and returns a SYCL [code]#device_event#
which can be used to wait on the completion of the copy.
a@ Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true)

Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements#
from the source pointer [code]#src# to destination pointer [code]#dest# with a
source stride specified by [code]#srcStride# and returns a SYCL
[code]#device_event# which can be used to wait on the completion of the copy.

a@
[source]
----
template <typename DataT>
device_event async_work_group_copy(decorated_global_ptr<DataT> dest,
decorated_local_ptr<const DataT> src,
template <typename DestDataT, SrcDataT>
device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
decorated_local_ptr<SrcDataT> src,
size_t numElements, size_t destStride) const
----
a@ Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements# from the source pointer [code]#src# to destination
pointer [code]#dest# with a destination stride specified by
[code]#destStride# and returns a SYCL [code]#device_event#
which can be used to wait on the completion of the copy.
a@ Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true)

Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements#
from the source pointer [code]#src# to destination pointer [code]#dest# with a
destination stride specified by [code]#destStride# and returns a SYCL
[code]#device_event# which can be used to wait on the completion of the copy.

a@
[source]
Expand Down Expand Up @@ -12441,56 +12470,53 @@ device_event async_work_group_copy(global_ptr<DataT> dest,
a@
[source]
----
template <typename DataT>
device_event async_work_group_copy(decorated_local_ptr<DataT> dest,
decorated_global_ptr<const DataT> src,
template <typename DestDataT, typename SrcDataT>
device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
decorated_local_ptr<SrcDataT> src,
size_t numElements) const
----
a@ Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements# from the source pointer [code]#src# to destination
pointer [code]#dest# and returns a SYCL [code]#device_event#
which can be used to wait on the completion of the copy.
a@ Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true)

a@
[source]
----
template <typename DataT>
device_event async_work_group_copy(decorated_global_ptr<DataT> dest,
decorated_local_ptr<const DataT> src,
size_t numElements) const
----
a@ Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements# from the source pointer [code]#src# to destination
pointer [code]#dest# and returns a SYCL [code]#device_event#
which can be used to wait on the completion of the copy.
Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements#
from the source pointer [code]#src# to destination pointer [code]#dest# and
returns a SYCL [code]#device_event# which can be used to wait on the
completion of the copy.

a@
[source]
----
template <typename DataT>
device_event async_work_group_copy(decorated_local_ptr<DataT> dest,
decorated_global_ptr<const DataT> src,
template <typename DestDataT, typename SrcDataT>
device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
decorated_global_ptr<SrcDataT> src,
size_t numElements, size_t srcStride) const
----
a@ Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements# from the source pointer [code]#src# to destination
pointer [code]#dest# with a source stride specified by
[code]#srcStride# and returns a SYCL [code]#device_event#
which can be used to wait on the completion of the copy.
a@ Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true)

Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements#
from the source pointer [code]#src# to destination pointer [code]#dest# with a
source stride specified by [code]#srcStride# and returns a SYCL
[code]#device_event# which can be used to wait on the completion of the copy.

a@
[source]
----
template <typename DataT>
device_event async_work_group_copy(decorated_global_ptr<DataT> dest,
decorated_local_ptr<const DataT> src,
template <typename DestDataT, SrcDataT>
device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
decorated_local_ptr<SrcDataT> src,
size_t numElements, size_t destStride) const
----
a@ Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements# from the source pointer [code]#src# to destination
pointer [code]#dest# with a destination stride specified by
[code]#destStride# and returns a SYCL [code]#device_event#
which can be used to wait on the completion of the copy.
a@ Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true)

Permitted types for [code]#DataT# are all scalar and vector types.
Asynchronously copies a number of elements specified by [code]#numElements#
from the source pointer [code]#src# to destination pointer [code]#dest# with a
destination stride specified by [code]#destStride# and returns a SYCL
[code]#device_event# which can be used to wait on the completion of the copy.

a@
[source]
Expand Down
8 changes: 6 additions & 2 deletions adoc/headers/accessorBuffer.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -184,8 +184,12 @@ class accessor {
cl::sycl::atomic<DataT, access::address_space::global_space>
operator[](id<Dimensions> index) const;

/* Deprecated in SYCL 2020 */
accessor_ptr<access::decorated::legacy> get_pointer() const noexcept;
/* Deprecated in SYCL 2020
Available only when: (AccessTarget == target::device) */
global_ptr<DataT> get_pointer() const noexcept;

/* Available only when (AccessTarget == target::host_task) */
std::add_pointer_t<value_type> get_pointer() const noexcept;

template <access::decorated IsDecorated>
accessor_ptr<IsDecorated> get_multi_ptr() const noexcept;
Expand Down
2 changes: 1 addition & 1 deletion adoc/headers/accessorLocal.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -63,7 +63,7 @@ template <typename DataT, int Dimensions = 1> class local_accessor {
reference operator[](size_t index) const;

/* Deprecated in SYCL 2020 */
accessor_ptr<access::decorated::legacy> get_pointer() const noexcept;
local_ptr<DataT> get_pointer() const noexcept;

template <access::decorated IsDecorated>
accessor_ptr<IsDecorated> get_multi_ptr() const noexcept;
Expand Down
32 changes: 20 additions & 12 deletions adoc/headers/group.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -75,25 +75,33 @@ template <int Dimensions = 1> class group {
size_t numElements,
size_t destStride) const;

template <typename DataT>
device_event async_work_group_copy(decorated_local_ptr<DataT> dest,
decorated_global_ptr<const DataT> src,
/* Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true) */
template <typename DestDataT, typename SrcDataT>
device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
decorated_global_ptr<SrcDataT> src,
size_t numElements) const;

template <typename DataT>
device_event async_work_group_copy(decorated_global_ptr<DataT> dest,
decorated_local_ptr<const DataT> src,
/* Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true) */
template <typename DestDataT, typename SrcDataT>
device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
decorated_local_ptr<SrcDataT> src,
size_t numElements) const;

template <typename DataT>
device_event async_work_group_copy(decorated_local_ptr<DataT> dest,
decorated_global_ptr<const DataT> src,
/* Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true) */
template <typename DestDataT, typename SrcDataT>
device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
decorated_global_ptr<SrcDataT> src,
size_t numElements,
size_t srcStride) const;

template <typename DataT>
device_event async_work_group_copy(decorated_global_ptr<DataT> dest,
decorated_local_ptr<const DataT> src,
/* Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true) */
template <typename DestDataT, typename SrcDataT>
device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
decorated_local_ptr<SrcDataT> src,
size_t numElements,
size_t destStride) const;

Expand Down
33 changes: 20 additions & 13 deletions adoc/headers/nditem.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -73,26 +73,33 @@ template <int Dimensions = 1> class nd_item {
size_t numElements,
size_t destStride) const;


template <typename DataT>
device_event async_work_group_copy(decorated_local_ptr<DataT> dest,
decorated_global_ptr<const DataT> src,
/* Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true) */
template <typename DestDataT, typename SrcDataT>
device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
decorated_global_ptr<SrcDataT> src,
size_t numElements) const;

template <typename DataT>
device_event async_work_group_copy(decorated_global_ptr<DataT> dest,
decorated_local_ptr<const DataT> src,
/* Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true) */
template <typename DestDataT, typename SrcDataT>
device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
decorated_local_ptr<SrcDataT> src,
size_t numElements) const;

template <typename DataT>
device_event async_work_group_copy(decorated_local_ptr<DataT> dest,
decorated_global_ptr<const DataT> src,
/* Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true) */
template <typename DestDataT, typename SrcDataT>
device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
decorated_global_ptr<SrcDataT> src,
size_t numElements,
size_t srcStride) const;

template <typename DataT>
device_event async_work_group_copy(decorated_global_ptr<DataT> dest,
decorated_local_ptr<const DataT> src,
/* Available only when: (std::is_same_v<DestDataT,
std::remove_const_t<SrcDataT>> == true) */
template <typename DestDataT, typename SrcDataT>
device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
decorated_local_ptr<SrcDataT> src,
size_t numElements,
size_t destStride) const;

Expand Down

0 comments on commit 99c0b3b

Please sign in to comment.