[SYCL][DOC] Spec for discard queue events

sycl/doc/extensions/DiscardQueueEvents/SYCL_EXT_ONEAPI_DISCARD_QUEUE_EVENTS.asciidoc

@@ -0,0 +1,223 @@
+= SYCL_EXT_ONEAPI_DISCARD_QUEUE_EVENTS
+: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}
+
+// This is necessary for asciidoc, but not for asciidoctor
+:cpp: C++
+
+== 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.
+
+This document describes an extension that introduces a discard_queue_events property for
+SYCL queues.  This property enables developers to inform a SYCL implementation that
+the returned event can be discarded.
+
+== Notice
+
+Copyright (c) 2021 Intel Corporation. All rights reserved.
+
+== Status
+
+Working 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
+
+== Contributors
+
+Alexander Flegontov, Intel  +
+Greg Lueck, Intel  +
+John Pennycook, Intel
+
+== Dependencies
+
+This extension is written against the SYCL 2020 specification, Revision 3.
+
+== 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_ONEAPI_DISCARD_QUEUE_EVENTS` 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
+
+This extension adds `ext::oneapi::property::queue::discard_queue_events` property for `sycl::queue`,
+by using this property the application informs a SYCL implementation that the returned event can be discarded.
+
+In cases where it will be possible, the implementation makes a decision about internal optimization -- avoiding
+the creation of low-level event from the backend. The property affects only kernel cases,
+i.e. avoid creating low-level event for submission of kernels.
+Performance improvement is the main motivation for using this property.
+
+Below is a usage example:
+[source,c++]
+----
+  sycl::property_list props{ext::oneapi::property::queue::discard_queue_events{}
+                            property::queue::in_order{}};
+  sycl::queue Queue( props );
+
+  // some USM preparations ..
+
+  sycl::event e1, e2, e3;
+
+  // returning "invalid" events from each submission function:
+  e1 = Queue.parallel_for(NDRange, [=](nd_item<1> item){ do_smth1(); });
+
+  e2 = Queue.single_task([=](){ do_smth2(); });
+
+  e3 = Queue.submit([&](handler &CGH) { CGH.parallel_for(NDRange, [=](nd_item<1> item){ do_smth3(); }); });
+
+  Queue.wait();
+----
+
+In the example above, the application doesn't use sycl events: `e1`, `e2`, `e3`
+and is waiting for the end of work by `queue::wait()`. The returned events will be
+_invalid_ events. _invalid_ event is a `sycl::event` with a status `invalid` and the event returns from each kernel
+submission when queue is created with the `discard_queue_events` property.
+See the description of behavior for this event below for details.
+
+_invalid_ event API behavior:
+[source,c++]
+----
+// the behavior is the same as for the default event, i.e. must throw an exception with the errc::invalid_object error code
+cl_event get() const
+
+// the behavior is the same as for the default event, i.e. must return true
+bool is_host() const;
+
+// the behavior is the same as for the default event, i.e. must return backend::host
+backend get_backend() const noexcept;
+
+// the behavior is the same as for the default event, i.e. must return a zero-sized vector
+std::vector<event> get_wait_list();
+
+// must throw an exception with the errc::invalid_object error code.
+void wait();
+
+// passing invalid event in the function must throw an exception with the errc::invalid_object error code.
+static void wait(const std::vector<event> &eventList);
+
+// must throw an exception with the errc::invalid_object error code.
+void wait_and_throw();
+
+// passing invalid event in the function must throw an exception with the errc::invalid_object error code.
+static void wait_and_throw(const std::vector<event> &eventList);
+
+// returns info::event_command_status::ext_oneapi_unknown status
+get_info<info::event::command_execution_status>() const;
+
+// the behavior is the same as for the default event, i.e. returns 0
+get_info<info::event::reference_count>() const;
+
+// must throw an exception with the errc::invalid_object error code.
+template <info::event_profiling param>
+typename param::return_type get_profiling_info() const;
+----
+
+The behavior when passing _invalid_ event to handler API:
+[source,c++]
+----
+// must throw an exception with the errc::invalid_object error code.
+handler::depends_on(event Event)
+
+// must throw an exception with the errc::invalid_object error code.
+handler::depends_on(const std::vector<event> &Events)
+----
+
+== Limitations
+
+The `discard_queue_events` property should be taken as a hint for SYCL implementation.
+In all kernel cases, SYCL implementation returns _invalid_ event, regardless of whether the optimization was done.
+Optimization occurs if the queue has `discard_queue_events` property and all features presented below are not met:
+
+ - Using out-of-order queue or using together with `enable_profiling` property
+ - Using https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc[fallback assert feature]
+ - Using streams, buffer/image accessors (excluding local accessors)
+
+
+See details for each below:
+
+=== Using out-of-order queue and/or using together with enable_profiling property
+
+No optimization if a queue is created with the `discard_queue_events` property and
+the property list includes the `enable_profiling` property or does not include `in_order` property.
+
+=== Using fallback assert feature
+
+If a kernel that uses the fallback assert feature is submitted to a queue created with
+the `discard_queue_events` property, the implementation must throw an exception with the errc::invalid_object error code.
+It is important to disable the logic for detecting assertion failures in kernels in principle
+since it is not enough just not to use the `assert` macro inside kernel code or compile with `NDEBUG`.
+To disable the logic, developer must compile with SYCL_DISABLE_FALLBACK_ASSERT macro defined,
+for more details see https://github.com/intel/llvm/blob/sycl/sycl/doc/PreprocessorMacros.md[document]
+
+[source,c++]
+----
+  sycl::property_list props{ext::oneapi::property::queue::discard_queue_events{},
+                            property::queue::in_order()};
+  sycl::queue Queue( props );
+  // if the submission of kernel1 is compiled with the logic for detecting assertion failures enabled,
+  // then the submission of the kernel throws an exception with the errc::invalid_object error code.
+  Queue.parallel_for<class kernel1>(NDRange, [=](nd_item<1> item){ do_smth(); });
+----
+
+=== Using streams, buffer/image accessors (excluding local accessors)
+
+No optimization if a kernel that uses stream objects, buffer or image accessors is submitted to a queue created with
+the `discard_queue_events` property. But using local accessors does not affect optimization.
+
+*NOTE : This extension is only compatible with kernels using Unified Shared Memory (USM) pointers.
+SYCL accessors are used to build a dependency graph that rely on events being created and not discarded.
+
+== Issues
+
+None.
+
+== Revision History
+
+[cols="5,15,15,70"]
+[grid="rows"]
+[options="header"]
+|========================================
+|Rev|Date|Author|Changes
+|1|2021-11-09|Alexander Flegontov |*Initial public working draft*
+|========================================

sycl/doc/extensions/README.md

@@ -49,6 +49,7 @@ DPC++ extensions status:
 | [SYCL_INTEL_bf16_conversion](Bf16Conversion/SYCL_INTEL_bf16_conversion.asciidoc)                                            | Partially supported (Level Zero: GPU)     | Currently available only on Xe HP GPU. ext_intel_bf16_conversion aspect is not supported. |
 | [Property List](PropertyList/SYCL_EXT_ONEAPI_property_list.asciidoc)                                                        | Proposal                                  | |
 | [KernelProperties](KernelProperties/KernelProperties.asciidoc)                                                              | Proposal                                  | |
+| [DiscardQueueEvents](DiscardQueueEvents/SYCL_EXT_ONEAPI_DISCARD_QUEUE_EVENTS.asciidoc) | Proposal | |
 
 Legend: