Skip to content

[SYCL] Align some extensions with SYCL 2020 #4432

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
Show file tree
Hide file tree
Changes from 5 commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
26 changes: 13 additions & 13 deletions sycl/doc/extensions/EnqueueBarrier/enqueue_barrier.asciidoc
Original file line number Diff line number Diff line change
Expand Up @@ -86,9 +86,9 @@ two new members to the `queue` class:
[grid="rows"]
[options="header"]
|========================================
|*handler::ext_intel_barrier*|*queue::ext_oneapi_submit_barrier*
|`void ext_intel_barrier()` | `event ext_oneapi_submit_barrier()`
|`void ext_intel_barrier( const vector_class<event> &waitList )` | `event ext_oneapi_submit_barrier( const vector_class<event> &waitList )`
|*handler::ext_oneapi_barrier*|*queue::ext_oneapi_submit_barrier*
|`void ext_oneapi_barrier()` | `event ext_oneapi_submit_barrier()`
|`void ext_oneapi_barrier( const vector_class<event> &waitList )` | `event ext_oneapi_submit_barrier( const vector_class<event> &waitList )`
|========================================

The first variant of the barrier takes no parameters, and waits for all previously submitted commands to the queue to enter the `info::event_command_status::complete` state before any command later submitted to the same queue is allowed to execute. A second variant of the barrier accepts a list of events, with the behavior that no commands submitted to the same queue after barrier submission may execute until all events in the `waitList` have entered the `info::event_command_status::complete` state. Both variants are non-blocking from the host program perspective, in that they do not wait for the barrier conditions to have been met before returning.
Expand All @@ -104,7 +104,7 @@ Some forms of the new barrier methods return an `event`, which can be used to pe

CG4 doesn't execute until all previous command groups submitted to the same queue (CG1, CG2, CG3) have entered the completed state.

==== 1. Using `handler::ext_intel_barrier()`:
==== 1. Using `handler::ext_oneapi_barrier()`:

[source,c++,NoName,linenums]
----
Expand All @@ -120,7 +120,7 @@ Queue.submit([&](cl::sycl::handler& cgh) {
});

Queue.submit([&](cl::sycl::handler& cgh) {
cgh.ext_intel_barrier();
cgh.ext_oneapi_barrier();
});

Queue.submit([&](cl::sycl::handler& cgh) {
Expand Down Expand Up @@ -157,7 +157,7 @@ Queue.submit([&](cl::sycl::handler& cgh) {

CG3 requires CG1 (in Queue1) and CG2 (in Queue2) to have completed before it (CG3) begins execution.

==== 1. Using `handler::ext_intel_barrier()`:
==== 1. Using `handler::ext_oneapi_barrier()`:

[source,c++,NoName,linenums]
----
Expand All @@ -171,7 +171,7 @@ auto event_barrier2 = Queue2.submit([&](cl::sycl::handler& cgh) {
});

Queue3.submit([&](cl::sycl::handler& cgh) {
cgh.ext_intel_barrier( vector_class<event>{event_barrier1, event_barrier2} );
cgh.ext_oneapi_barrier( vector_class<event>{event_barrier1, event_barrier2} );
});

Queue3.submit([&](cl::sycl::handler& cgh) {
Expand Down Expand Up @@ -236,8 +236,8 @@ void wait();
[options="header"]
|========================================
|*Member functions*|*Description*
|`event ext_oneapi_submit_barrier()` | Same effect as submitting a `handler::ext_intel_barrier()` within a command group to this `queue`. The returned event enters the `info::event_command_status::complete` state when all events that the barrier is dependent on (implicitly from all previously submitted commands to the same queue) have entered the `info::event_command_status::complete` state.
|`event ext_oneapi_submit_barrier( const vector_class<event> &waitList )` | Same effect as submitting a `handler:ext_intel_barrier( const vector_class<event> &waitList )` within a command group to this `queue`. The returned event enters the `info::event_command_status::complete` state when all events that the barrier is dependent on (explicitly from `waitList`) have entered the `info::event_command_status::complete` state.
|`event ext_oneapi_submit_barrier()` | Same effect as submitting a `handler::ext_oneapi_barrier()` within a command group to this `queue`. The returned event enters the `info::event_command_status::complete` state when all events that the barrier is dependent on (implicitly from all previously submitted commands to the same queue) have entered the `info::event_command_status::complete` state.
|`event ext_oneapi_submit_barrier( const vector_class<event> &waitList )` | Same effect as submitting a `handler:ext_oneapi_barrier( const vector_class<event> &waitList )` within a command group to this `queue`. The returned event enters the `info::event_command_status::complete` state when all events that the barrier is dependent on (explicitly from `waitList`) have entered the `info::event_command_status::complete` state.
|========================================


Expand Down Expand Up @@ -273,9 +273,9 @@ void fill(void *ptr, const T &pattern, size_t count);
template <typename T>
void fill(void *ptr, const T &pattern, size_t count);

void ext_intel_barrier();
void ext_oneapi_barrier();

void ext_intel_barrier( const vector_class<event> &waitList );
void ext_oneapi_barrier( const vector_class<event> &waitList );

};
...
Expand All @@ -296,8 +296,8 @@ Barriers can be created by two members of the `handler` class that force synchro
[options="header"]
|========================================
|*Member functions*|*Description*
|`void ext_intel_barrier()` | Prevents any commands submitted afterward to this queue from executing until all commands previously submitted to this queue have entered the `info::event_command_status::complete` state.
|`void ext_intel_barrier( const vector_class<event> &waitList` ) | Prevents any commands submitted afterward to this queue from executing until all events in `waitList` have entered the `info::event_command_status::complete` state. If `waitList` is empty, then the barrier has no effect.
|`void ext_oneapi_barrier()` | Prevents any commands submitted afterward to this queue from executing until all commands previously submitted to this queue have entered the `info::event_command_status::complete` state.
|`void ext_oneapi_barrier( const vector_class<event> &waitList` ) | Prevents any commands submitted afterward to this queue from executing until all events in `waitList` have entered the `info::event_command_status::complete` state. If `waitList` is empty, then the barrier has no effect.
|========================================

== Issues
Expand Down
36 changes: 9 additions & 27 deletions sycl/doc/extensions/MemChannel/MemChannel.asciidoc
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
= SYCL_EXT_INTEL_MEM_CHANNEL_PROPERTY
= SYCL_INTEL_mem_channel_property

== Introduction
NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by permission by Khronos.
Expand All @@ -23,30 +23,14 @@ Because the interfaces defined by this specification are not final and are subje
== Version

Built On: {docdate} +
Revision: 2
Revision: 1

== Dependencies

This extension is written against the SYCL 2020 specification, Revision 3.
This extension is written against the SYCL 2020 provisional specification, Revision 1.

The use of this extension requires a target that supports cl_intel_mem_channel_property or equivalent if OpenCL is used as the underlying device runtime.

== Feature Test Macro

This extension provides a feature-test macro as described in the core SYCL
specification section 6.3.3 "Feature test macros". Therefore, an
implementation supporting this extension must predefine the macro
`SYCL_EXT_INTEL_MEM_CHANNEL_PROPERTY` to one of the values defined in the table below.
Applications can test for the existence of this macro to determine if the
implementation supports this feature, or applications can test the macro's
value to determine which of the extension's APIs the implementation supports.

[%header,cols="1,5"]
|===
|Value |Description
|1 |Initial extension version. Base features are supported.
|===

== Overview

On some targets manual assignment of buffers to memory regions can improve memory bandwidth. This extension adds a buffer property to indicate in which memory channel a particular buffer should be allocated. This information is an optimization hint to the runtime and thus it is legal to ignore.
Expand All @@ -61,27 +45,27 @@ Add a new property to Table 4.33: Properties supported by the SYCL buffer class
[options="header"]
|===
| Property | Description
| property::buffer::ext_intel_mem_channel | The `ext_intel_mem_channel` property is a hint to the SYCL runtime that the buffer should be stored in a particular memory channel provided to the property.
| property::buffer::mem_channel | The `mem_channel` property is a hint to the SYCL runtime that the buffer should be stored in a particular memory channel provided to the property.
|===
--

Add a new constructor to Table 41: Constructors of the buffer property classes as follows:
Add a new constructor to Table 4.34: Constructors of the buffer property classes as follows:

--
[options="header"]
|===
| Constructor | Description
| property::buffer::ext_intel_mem_channel::ext_intel_mem_channel(cl_uint channel) | Constructs a SYCL `ext_intel_mem_channel` property instance with the specified channel ID. The range of valid values depends on the target and is implementation defined. Invalid values do not need to result in an error as the property is only a hint.
| property::buffer::mem_channel::mem_channel(cl_uint channel) | Constructs a SYCL `mem_channel` property instance with the specified channel ID. The range of valid values depends on the target and is implementation defined. Invalid values do not need to result in an error as the property is only a hint.
|===
--

Add a new member function to Table 42: Member functions of the buffer property classes as follows:
Add a new member function to Table 4.35: Member functions of the buffer property classes as follows:

--
[options="header"]
|===
| Member function | Description
| cl_uint property::buffer::ext_intel_mem_channel::get_channel() const | Returns the cl_uint which was specified when constructing this SYCL `ext_intel_mem_channel` property.
| cl_uint property::buffer::mem_channel::get_channel() const | Returns the cl_uint which was specified when constructing this SYCL `mem_channel` property.
|===
--

Expand All @@ -103,7 +87,7 @@ enum class aspect {
} // namespace sycl
```

Add an entry for the new aspect to Table 26: Device aspects defined by the core SYCL specification:
Add an entry for the new aspect to Table 4.20: Device aspects defined by the core SYCL specification:

--
[options="header"]
Expand All @@ -123,6 +107,4 @@ Add an entry for the new aspect to Table 26: Device aspects defined by the core
|========================================
|Rev|Date|Author|Changes
|1|2020-10-26|Joe Garvey|*Initial public draft*
|2|2021-08-30|Dmitry Vodopyanov|*Updated according to SYCL 2020 reqs for extensions*

|========================================
8 changes: 4 additions & 4 deletions sycl/doc/extensions/SYCL_ONEAPI_dot_accumulate.asciidoc
Original file line number Diff line number Diff line change
Expand Up @@ -46,7 +46,7 @@ Because the interfaces defined by this specification are not final and are subje
== Version

Built On: {docdate} +
Revision: B
Revision: 3

== Contact

Expand Down Expand Up @@ -135,9 +135,9 @@ None.
[options="header"]
|========================================
|Rev|Date|Author|Changes
|A|2019-12-13|Ben Ashbaugh|*Initial draft*
|B|2019-12-18|Ben Ashbaugh|Switched to standard C++ fixed width types.
|C|2020-10-26|Rajiv Deodhar|Added int32 types.
|1|2019-12-13|Ben Ashbaugh|*Initial draft*
|2|2019-12-18|Ben Ashbaugh|Switched to standard C++ fixed width types.
|3|2020-10-26|Rajiv Deodhar|Added int32 types.
|========================================

//************************************************************************
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ Because the interfaces defined by this specification are not final and are subje
== Version

Built On: {docdate} +
Revision: B
Revision: 2

== Dependencies

Expand Down Expand Up @@ -134,6 +134,6 @@ using host_ptr = multi_ptr<ElementType, access::address_space::ext_intel_global_
[options="header"]
|========================================
|Rev|Date|Author|Changes
|A|2020-06-18|Joe Garvey|Initial public draft
|B|2021-08-30|Dmitry Vodopyanov|Updated according to SYCL 2020 reqs for extensions
|1|2020-06-18|Joe Garvey|Initial public draft
|2|2021-08-30|Dmitry Vodopyanov|Updated according to SYCL 2020 reqs for extensions
|========================================
Original file line number Diff line number Diff line change
Expand Up @@ -31,7 +31,7 @@ Because the interfaces defined by this specification are not final and are subje
== Version

Built On: {docdate} +
Revision: A
Revision: 1

== Contact
Michael Kinsner, Intel (michael 'dot' kinsner 'at' intel 'dot' com)
Expand Down Expand Up @@ -95,5 +95,5 @@ It also notifies the SYCL runtime to store the given accessor in that memory. |
[options="header"]
|========================================
|Rev|Date|Author|Changes
|A|2020-09-08|Joe Garvey|*Initial public draft*
|1|2020-09-08|Joe Garvey|*Initial public draft*
|========================================
1 change: 0 additions & 1 deletion sycl/include/CL/sycl/feature_test.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,6 @@ namespace sycl {
#endif
#define SYCL_EXT_INTEL_BF16_CONVERSION 1
#define SYCL_EXT_ONEAPI_ENQUEUE_BARRIER 1
#define SYCL_EXT_INTEL_MEM_CHANNEL_PROPERTY 1
#define SYCL_EXT_INTEL_USM_ADDRESS_SPACES 1
#define SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO 1

Expand Down
14 changes: 6 additions & 8 deletions sycl/include/CL/sycl/handler.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -2286,35 +2286,33 @@ class __SYCL_EXPORT handler {
/// Prevents any commands submitted afterward to this queue from executing
/// until all commands previously submitted to this queue have entered the
/// complete state.
void ext_intel_barrier() {
void ext_oneapi_barrier() {
throwIfActionIsCreated();
setType(detail::CG::Barrier);
}

/// Prevents any commands submitted afterward to this queue from executing
/// until all commands previously submitted to this queue have entered the
/// complete state.
__SYCL2020_DEPRECATED("use 'ext_intel_barrier' instead")
void barrier() { ext_intel_barrier(); }
__SYCL2020_DEPRECATED("use 'ext_oneapi_barrier' instead")
void barrier() { ext_oneapi_barrier(); }

/// Prevents any commands submitted afterward to this queue from executing
/// until all events in WaitList have entered the complete state. If WaitList
/// is empty, then the barrier has no effect.
///
/// \param WaitList is a vector of valid SYCL events that need to complete
/// before barrier command can be executed.
void ext_intel_barrier(const std::vector<event> &WaitList);
void ext_oneapi_barrier(const std::vector<event> &WaitList);

/// Prevents any commands submitted afterward to this queue from executing
/// until all events in WaitList have entered the complete state. If WaitList
/// is empty, then the barrier has no effect.
///
/// \param WaitList is a vector of valid SYCL events that need to complete
/// before barrier command can be executed.
__SYCL2020_DEPRECATED("use 'ext_intel_barrier' instead")
void barrier(const std::vector<event> &WaitList) {
ext_intel_barrier(WaitList);
}
__SYCL2020_DEPRECATED("use 'ext_oneapi_barrier' instead")
void barrier(const std::vector<event> &WaitList);

/// Copies data from one memory region to another, both pointed by
/// USM pointers.
Expand Down
12 changes: 3 additions & 9 deletions sycl/include/CL/sycl/properties/buffer_properties.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -40,22 +40,16 @@ class context_bound
sycl::context MCtx;
};

class ext_intel_mem_channel : public detail::PropertyWithData<
detail::PropWithDataKind::BufferMemChannel> {
class mem_channel : public detail::PropertyWithData<
detail::PropWithDataKind::BufferMemChannel> {
public:
ext_intel_mem_channel(uint32_t Channel) : MChannel(Channel) {}
mem_channel(uint32_t Channel) : MChannel(Channel) {}
uint32_t get_channel() const { return MChannel; }

private:
uint32_t MChannel;
};

class __SYCL2020_DEPRECATED("use 'ext_intel_mem_channel' instead") mem_channel
: public ext_intel_mem_channel {
public:
mem_channel(uint32_t Channel) : ext_intel_mem_channel(Channel) {}
};

} // namespace buffer
} // namespace property

Expand Down
4 changes: 2 additions & 2 deletions sycl/include/CL/sycl/queue.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -314,7 +314,7 @@ class __SYCL_EXPORT queue {
/// group is being enqueued on.
event ext_oneapi_submit_barrier(_CODELOCONLYPARAM(&CodeLoc)) {
return submit(
[=](handler &CGH) { CGH.ext_intel_barrier(); } _CODELOCFW(CodeLoc));
[=](handler &CGH) { CGH.ext_oneapi_barrier(); } _CODELOCFW(CodeLoc));
}

/// Prevents any commands submitted afterward to this queue from executing
Expand Down Expand Up @@ -342,7 +342,7 @@ class __SYCL_EXPORT queue {
event ext_oneapi_submit_barrier(
const std::vector<event> &WaitList _CODELOCPARAM(&CodeLoc)) {
return submit([=](handler &CGH) {
CGH.ext_intel_barrier(WaitList);
CGH.ext_oneapi_barrier(WaitList);
} _CODELOCFW(CodeLoc));
}

Expand Down
Loading