[SYCL][Doc] Release SYCL_INTEL_enqueue_barrier extension document (#1199)

Signed-off-by: michael.kinsner <[email protected]>

Author: Mike Kinsner

sycl/doc/extensions/EnqueueBarrier/enqueue_barrier.asciidoc

@@ -0,0 +1,315 @@
+= SYCL_INTEL_enqueue_barrier
+:source-highlighter: coderay
+:coderay-linenums-mode: table
+
+// This section needs to be after the document title.
+:doctype: book
+:toc2:
+:toc: left
+:encoding: utf-8
+:lang: en
+
+:blank: pass:[ +]
+
+// Set the default source code type in this document to C++,
+// for syntax highlighting purposes.  This is needed because
+// docbook uses c++ and html5 uses cpp.
+:language: {basebackend@docbook:c++:cpp}
+
+== Introduction
+IMPORTANT: This specification is a draft.
+
+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.
+
+NOTE: This document is better viewed when rendered as html with asciidoctor.  GitHub does not render image icons.
+
+This document presents a series of changes proposed for a future version of the SYCL Specification. The goal of this proposal is to provide non-blocking APIs that provide synchronization on SYCL command queue for programmers.
+
+
+== Name Strings
+
++SYCL_INTEL_enqueue_barrier+
+
+== Notice
+
+Copyright (c) 2019-2020 Intel Corporation.  All rights reserved.
+
+== Status
+
+Final Draft
+
+This is a preview extension specification, intended to provide early access to a feature for review and community feedback. When the feature matures, this specification may be released as a formal extension.
+
+Because the interfaces defined by this specification are not final and are subject to change they are not intended to be used by shipping software products.
+
+== Version
+
+Built On: {docdate} +
+Revision: 1
+
+== Contact
+Please open an issue in the https://github.com/intel/llvm/tree/sycl/sycl/doc/extensions/[extensions repository]
+
+== Dependencies
+
+This extension is written against the SYCL 1.2.1 specification, Revision v1.2.1-6.
+
+== Overview
+
+SYCL 1.2.1 defines a graph-based task execution model, based on kernels or explicit memory operations submitted to out-of-order queues. Dependencies between these kernels are represented by
+accessors that form data dependence edges in the execution graph. The USM extension <<usmlink,[1]>> doesn't have accessors, so instead solves
+this by defining `handler::depends_on` methods to specify event-based control dependencies between command groups.
+
+There are situations where defining dependencies based on events is more explicit than desired or required by an application. For instance, the user may know that a given task depends on all previously submitted tasks. Instead of explicitly adding all the required depends_on calls, the user could express this intent via a single call, making the program more concise and explicit.
+
+To simplify the interface, this extension adds an enqueued barrier that provides synchronization on SYCL command
+queues, with the following properties:
+
+1. Enqueued barriers are non-blocking from the host perspective.  The barrier is enqueued, so operates as part of the execution graph asynchronously from host program execution
+2. Command groups submitted to the same queue after the barrier is enqueued are not scheduled for execution until the barrier wait conditions are satisfied
+
+This proposal adds two new members to the `handler` class, and
+two new members to the `queue` class:
+
+[cols="70,70"]
+[grid="rows"]
+[options="header"]
+|========================================
+|*handler::barrier*|*queue::submit_barrier*
+|`void barrier()` | `event submit_barrier()`
+|`void barrier( const vector_class<event> &waitList )` | `event 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.
+
+The new barrier operations implicitly add dependence edges to the SYCL task execution graph, and do not have other side effects.
+
+Some forms of the new barrier methods return an `event`, which can be used to perform other synchronization operations (e.g. `depends_on()`).  The event from one of the enqueued barrier operations enters the `info::event_command_status::complete` state when all events that the barrier is dependent on (explicitly or implicitly) have entered the `info::event_command_status::complete` state.
+
+
+== Example Scenarios
+
+=== Scenario 1: Enqueued barrier dependent on all commands previously submitted to the same queue
+
+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::barrier()`:
+
+[source,c++,NoName,linenums]
+----
+...
+Queue.submit([&](cl::sycl::handler& cgh) {
+  // CG1
+});
+Queue.submit([&](cl::sycl::handler& cgh) {
+  // CG2
+});
+Queue.submit([&](cl::sycl::handler& cgh) {
+  // CG3
+});
+
+Queue.submit([&](cl::sycl::handler& cgh) {
+  cgh.barrier();
+});
+
+Queue.submit([&](cl::sycl::handler& cgh) {
+  // CG4
+});
+...
+----
+
+==== 2. Using `queue::submit_barrier()`:
+
+[source,c++,NoName,linenums]
+----
+...
+Queue.submit([&](cl::sycl::handler& cgh) {
+  // CG1
+});
+Queue.submit([&](cl::sycl::handler& cgh) {
+  // CG2
+});
+Queue.submit([&](cl::sycl::handler& cgh) {
+  // CG3
+});
+
+Queue.submit_barrier();
+
+Queue.submit([&](cl::sycl::handler& cgh) {
+  // CG4
+});
+...
+----
+
+
+=== Scenario 2: Enqueued barrier dependent on specific events from previously submitted commands
+
+CG3 requires CG1 (in Queue1) and CG2 (in Queue2) to have completed before it (CG3) begins execution.
+
+==== 1. Using `handler::barrier()`:
+
+[source,c++,NoName,linenums]
+----
+...
+auto event_barrier1 = Queue1.submit([&](cl::sycl::handler& cgh) {
+  // CG1
+});
+
+auto event_barrier2 = Queue2.submit([&](cl::sycl::handler& cgh) {
+  // CG2
+});
+
+Queue3.submit([&](cl::sycl::handler& cgh) {
+  cgh.barrier( vector_class<event>{event_barrier1, event_barrier2} );
+});
+
+Queue3.submit([&](cl::sycl::handler& cgh) {
+  // CG3
+});
+...
+----
+
+==== 2. Using `queue::submit_barrier()`:
+
+[source,c++,NoName,linenums]
+----
+...
+auto event_barrier1 = Queue1.submit([&](cl::sycl::handler& cgh) {
+  // CG1
+});
+
+auto event_barrier2 = Queue2.submit([&](cl::sycl::handler& cgh) {
+  // CG2
+});
+
+Queue3.submit_barrier( vector_class<event>{event_barrier1, event_barrier2} );
+
+Queue3.submit([&](cl::sycl::handler& cgh) {
+  // CG3
+});
+...
+----
+
+== Specification changes
+
+=== Modify part of Section 4.6.5.1
+
+*Change from:*
+[source,c++,NoName,linenums]
+----
+...
+template <typename T>
+event submit(T cgf, const queue &secondaryQueue);
+
+void wait();
+...
+----
+*To:*
+[source,c++,NoName,linenums]
+----
+...
+template <typename T>
+event submit(T cgf, const queue &secondaryQueue);
+
+event submit_barrier();
+
+event submit_barrier( const vector_class<event> &waitList );
+
+void wait();
+...
+----
+=== Add rows to Table 4.22
+
+[cols="70,300"]
+[grid="rows"]
+[options="header"]
+|========================================
+|*Member functions*|*Description*
+|`event submit_barrier()` | Same effect as submitting a `handler::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 submit_barrier( const vector_class<event> &waitList )` | Same effect as submitting a `handler: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.
+|========================================
+
+
+=== Modify Section 4.8.2
+
+==== Change first sentence from:
+A command group scope in SYCL, as it is defined in Section 3.4.1, consists of a single kernel or explicit memory
+operation (handler methods such as copy, update_host, fill), together with its requirements.
+
+==== To:
+
+A command group scope in SYCL, as it is defined in Section 3.4.1, consists of a single kernel, explicit memory
+operation (handler methods such as copy, update_host, fill) or barrier, together with its requirements.
+
+=== Modify part of Section 4.8.3
+
+*Change from:*
+[source,c++,NoName,linenums]
+----
+...
+template<typename T, int dim, access::mode mode, access::target tgt>
+void fill(accessor<T, dim, mode, tgt> dest, const T& src);
+
+};
+...
+----
+
+*To:*
+[source,c++,NoName,linenums]
+----
+...
+template<typename T, int dim, access::mode mode, access::target tgt>
+void fill(accessor<T, dim, mode, tgt> dest, const T& src);
+
+void barrier();
+
+void barrier( const vector_class<event> &waitList );
+
+};
+...
+----
+
+=== Add a new section between Section 4.8.6 and 4.8.7
+
+4.8.X SYCL functions for enqueued synchronization barriers
+
+Barriers may be submitted to a queue, with the effect that they prevent later operations submitted to the same queue from executing until the barrier wait conditions have been satisfied. The wait conditions can be explicitly described by `waitList` or implicitly from all previously submitted commands to the same queue. There are no constraints on the context from which queues may participate in the `waitList`. Enqueued barriers do not block host program execution, but instead form additional dependence edges with the execution task graph. 
+
+Barriers can be created by two members of the `handler` class that force synchronization on the SYCL command queue. The first variant of the `handler` barrier (`handler::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. The second variant of the `handler` barrier (`handler::barrier( const vector_class<event> &waitList )`) 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. 
+
+=== Add a new table in the new section between 4.8.6 and 4.8.7: Member functions of the handler class.
+
+[cols="70,300"]
+[grid="rows"]
+[options="header"]
+|========================================
+|*Member functions*|*Description*
+|`void 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 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.
+|========================================
+
+== References
+1. [[usmlink]]https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/USM/USM.adoc
+
+== Issues
+
+None.
+
+== Revision History
+
+[cols="5,15,15,70"]
+[grid="rows"]
+[options="header"]
+|========================================
+|Rev|Date|Author|Changes
+|1|2020-02-26|Ye Ting|*Initial public release*
+|========================================
+
+//************************************************************************
+//Other formatting suggestions:
+//
+//* Use *bold* text for host APIs, or [source] syntax highlighting.
+//* Use +mono+ text for device APIs, or [source] syntax highlighting.
+//* Use +mono+ text for extension names, types, or enum values.
+//* Use _italics_ for parameters.
+//************************************************************************