reble
diff --git a/‎README.md
Lines changed: 21 additions & 53 deletions b/‎README.md
Lines changed: 21 additions & 53 deletions
diff --git a/‎buildbot/configure.py
Lines changed: 7 additions & 0 deletions b/‎buildbot/configure.py
Lines changed: 7 additions & 0 deletions
diff --git a/‎sycl/CMakeLists.txt
Lines changed: 7 additions & 0 deletions b/‎sycl/CMakeLists.txt
Lines changed: 7 additions & 0 deletions
diff --git a/‎sycl/doc/GetStartedGuide.md
Lines changed: 1 addition & 0 deletions b/‎sycl/doc/GetStartedGuide.md
Lines changed: 1 addition & 0 deletions
diff --git a/‎sycl/include/sycl/detail/common.hpp
Lines changed: 6 additions & 4 deletions b/‎sycl/include/sycl/detail/common.hpp
Lines changed: 6 additions & 4 deletions
diff --git a/‎sycl/include/sycl/detail/pi.def
Lines changed: 7 additions & 0 deletions b/‎sycl/include/sycl/detail/pi.def
Lines changed: 7 additions & 0 deletions
diff --git a/‎sycl/include/sycl/detail/pi.h
Lines changed: 81 additions & 1 deletion b/‎sycl/include/sycl/detail/pi.h
Lines changed: 81 additions & 1 deletion
@@ -1,72 +1,44 @@
-This is the Intel staging area for llvm.org contributions and the home for
-Intel LLVM-based projects:
+# SYCL Command Graph Extensions
 
-- [oneAPI DPC++ compiler](#oneapi-dpc-compiler)
-- [Late-outline OpenMP and OpenMP Offload](#late-outline-openmp-and-openmp-offload)
+This is the collaboration space for the oneAPI vendor Command Graph extension for SYCL2020. It provides an API for defining a graph of operations and their dependencies once and submitting this graph repeatedly for execution.
 
-For general contribution process see [CONTRIBUTING.md](./CONTRIBUTING.md)
+### Specification
 
-## oneAPI DPC++ compiler
+A draft of our Command Graph extension proposal can be found here:
+[https://github.com/intel/llvm/pull/5626](https://github.com/intel/llvm/pull/5626).
 
-[![oneAPI logo](https://spec.oneapi.io/oneapi-logo-white-scaled.jpg)](https://www.oneapi.io/)
+### Implementation
 
-[![SYCL Post Commit](https://github.com/intel/llvm/actions/workflows/sycl_post_commit.yml/badge.svg?branch=sycl)](https://github.com/intel/llvm/actions/workflows/sycl_post_commit.yml)
-[![Generate Doxygen documentation](https://github.com/intel/llvm/actions/workflows/gh_pages.yml/badge.svg?branch=sycl)](https://github.com/intel/llvm/actions/workflows/gh_pages.yml)
+Our current prototype implementation can be found here:
+[https://github.com/reble/llvm/tree/sycl-graph-develop](https://github.com/reble/llvm/tree/sycl-graph-develop).
 
-The DPC++ is a LLVM-based compiler project that implements compiler and runtime
-support for the SYCL\* language. The project is hosted in the
-[sycl](https://github.com/intel/llvm/tree/sycl) branch and is synced with the
-tip of the LLVM upstream main branch on a regular basis (revisions delay is
-usually not more than 1-2 weeks). DPC++ compiler takes everything from LLVM
-upstream as is, however some modules of LLVM might be not included in the
-default project build configuration. Additional modules can be enabled by
-modifying build framework settings.
+Limitations include:
+* LevelZero backend support only.
+* Accessors and reductions are currently not supported.
 
-The DPC++ goal is to support the latest SYCL\* standard and work on that is in
-progress. DPC++ also implements a number of extensions to the SYCL\* standard,
-which can be found in the [sycl/doc/extensions](./sycl/doc/extensions)
-directory.
+### Other Material
 
-The main purpose of this project is open source collaboration on the DPC++
-compiler implementation in LLVM across a variety of architectures, prototyping
-compiler and runtime library solutions, designing future extensions, and
-conducting experiments. As the implementation becomes more mature, we try to
-upstream as much DPC++ support to LLVM main branch as possible. See
-[SYCL upstreaming working group notes](https://github.com/intel/llvm/wiki/SYCL-upstreaming-working-group-meeting-notes)
-for more details.
+This extension was presented at the oneAPI Technical Advisory board (Sept'22 meeting). Slides: [https://github.com/oneapi-src/oneAPI-tab/blob/main/language/presentations/2022-09-28-TAB-SYCL-Graph.pdf](https://github.com/oneapi-src/oneAPI-tab/blob/main/language/presentations/2022-09-28-TAB-SYCL-Graph.pdf).
 
-Note that this project can be used as a technical foundation for some
-proprietary compiler products, which may leverage implementations from this open
-source project. One of the examples is
-[Intel(R) oneAPI DPC++ Compiler](https://www.intel.com/content/www/us/en/developer/tools/oneapi/dpc-compiler.html)
-Features parity between this project and downstream projects is not guaranteed.
+## Intel Project for LLVM\* technology
 
-Project documentation is available at:
-[DPC++ Documentation](https://intel.github.io/llvm-docs/).
+We target a contribution through the origin of this fork: [Intel staging area for llvm.org contributions](https://github.com/intel/llvm).
 
 ### How to use DPC++
 
-#### Docker containers
-
-See available containers with pre-built/pre-installed DPC++ compiler at:
-[Containers](./sycl/doc/developer/DockerBKMs.md#sycl-containers-overview)
-
 #### Releases
 
-Daily builds of the sycl branch on Linux are available at
-[releases](https://github.com/intel/llvm/releases).
-A few times a year, we publish [Release Notes](./sycl/ReleaseNotes.md) to
-highlight all important changes made in the project: features implemented and
-issues addressed. The corresponding builds can be found using
-[search](https://github.com/intel/llvm/releases?q=oneAPI+DPC%2B%2B+Compiler&expanded=true)
-in daily releases. None of the branches in the project are stable or rigorously
-tested for production quality control, so the quality of these releases is
-expected to be similar to the daily releases.
+TDB
 
 #### Build from sources
 
 See [Get Started Guide](./sycl/doc/GetStartedGuide.md).
 
+SYCL Graph support is enabled with:
+ * Configuration script: `configure.py -enable-sycl-graph`.
+ * CMake: `cmake -DSYCL_ENABLE_GRAPH`.
+
+A fallback emulation mode is used otherwise that enables the graph API but eagerly submits kernels.
 ### Report a problem
 
 Submit an [issue](https://github.com/intel/llvm/issues) or initiate a 
@@ -76,10 +48,6 @@ Submit an [issue](https://github.com/intel/llvm/issues) or initiate a
 
 See [ContributeToDPCPP](./sycl/doc/developer/ContributeToDPCPP.md).
 
-## Late-outline OpenMP\* and OpenMP\* Offload
-
-See [openmp](/tree/openmp) branch.
-
 # License
 
 See [LICENSE](./sycl/LICENSE.TXT) for details.
 
@@ -57,6 +57,8 @@ def do_configure(args):
 
     if sys.platform != "darwin":
         sycl_enabled_plugins.append("level_zero")
+        
+    sycl_enable_graph = 'OFF'
 
     # lld is needed on Windows or for the HIP plugin on AMD
     if platform.system() == 'Windows' or (args.hip and args.hip_platform == 'AMD'):
@@ -108,6 +110,9 @@ def do_configure(args):
 
     if args.use_lld:
         llvm_enable_lld = 'ON'
+    
+    if args.enable_sycl_graph:
+         sycl_enable_graph = 'ON'
 
     # CI Default conditionally appends to options, keep it at the bottom of
     # args handling
@@ -166,6 +171,7 @@ def do_configure(args):
         "-DLLVM_ENABLE_SPHINX={}".format(llvm_enable_sphinx),
         "-DBUILD_SHARED_LIBS={}".format(llvm_build_shared_libs),
         "-DSYCL_ENABLE_XPTI_TRACING={}".format(sycl_enable_xpti_tracing),
+        "-DSYCL_ENABLE_GRAPH={}".format(sycl_enable_graph),
         "-DLLVM_ENABLE_LLD={}".format(llvm_enable_lld),
         "-DXPTI_ENABLE_WERROR={}".format(xpti_enable_werror),
         "-DSYCL_CLANG_EXTRA_FLAGS={}".format(sycl_clang_extra_flags),
@@ -240,6 +246,7 @@ def main():
                         help="host LLVM target architecture, defaults to X86, multiple targets may be provided as a semi-colon separated string")
     parser.add_argument("--enable-esimd-emulator", action='store_true', help="build with ESIMD emulation support")
     parser.add_argument("--enable-all-llvm-targets", action='store_true', help="build compiler with all supported targets, it doesn't change runtime build")
+    parser.add_argument("--enable-sycl-graph", action='store_true', help="build with SYCL Graph support")
     parser.add_argument("--no-assertions", action='store_true', help="build without assertions")
     parser.add_argument("--docs", action='store_true', help="build Doxygen documentation")
     parser.add_argument("--werror", action='store_true', help="Treat warnings as errors")
 
@@ -69,6 +69,13 @@ endif()
 # of the SYCL runtime and expect enabling
 option(SYCL_ENABLE_XPTI_TRACING "Enable tracing of SYCL constructs" OFF)
 
+# Create a soft option for enabling or disabling the experimental support
+# for SYCl Graph
+option(SYCL_ENABLE_GRAPH "Enable experimental SYCL Graph support" OFF)
+if (SYCL_ENABLE_GRAPH)
+  set(SYCL_BUILD_SYCL_GRAPH ON)
+endif()
+
 if(MSVC)
   set_property(GLOBAL PROPERTY USE_FOLDERS ON)
   # Skip asynchronous C++ exceptions catching and assume "extern C" functions
 
@@ -126,6 +126,7 @@ flags can be found by launching the script with `--help`):
 * `--enable-esimd-emulator` -> enable ESIMD CPU emulation (see [ESIMD CPU emulation](#build-dpc-toolchain-with-support-for-esimd-cpu-emulation))
 * `--enable-all-llvm-targets` -> build compiler (but not a runtime) with all
   supported targets
+* `--enable-sycl-graph` -> build SYCL Graph support
 * `--shared-libs` -> Build shared libraries
 * `-t` -> Build type (Debug or Release)
 * `-o` -> Path to build directory
 
@@ -104,18 +104,20 @@ struct code_location {
 
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
 #define _CODELOCONLYPARAM(a)                                                   \
-  const detail::code_location a = detail::code_location::current()
+  const ::sycl::detail::code_location a =                                      \
+      ::sycl::detail::code_location::current()
 #define _CODELOCPARAM(a)                                                       \
-  , const detail::code_location a = detail::code_location::current()
-#define _CODELOCPARAMDEF(a) , const detail::code_location a
+  , const ::sycl::detail::code_location a =                                    \
+        ::sycl::detail::code_location::current()
+#define _CODELOCPARAMDEF(a) , const ::sycl::detail::code_location a
 
 #define _CODELOCARG(a)
 #define _CODELOCFW(a) , a
 #else
 #define _CODELOCONLYPARAM(a)
 #define _CODELOCPARAM(a)
 
-#define _CODELOCARG(a) const detail::code_location a = {}
+#define _CODELOCARG(a) const ::sycl::detail::code_location a = {}
 #define _CODELOCFW(a)
 #endif
 
 
@@ -141,6 +141,13 @@ _PI_API(piextKernelSetArgSampler)
 
 _PI_API(piextPluginGetOpaqueData)
 
+/// command-buffer Extension
+_PI_API(piextCommandBufferCreate)
+_PI_API(piextCommandBufferRetain)
+_PI_API(piextCommandBufferRelease)
+_PI_API(piextCommandBufferFinalize)
+_PI_API(piextCommandBufferNDRangeKernel)
+_PI_API(piextEnqueueCommandBuffer)
 _PI_API(piPluginGetLastError)
 
 _PI_API(piTearDown)
 
@@ -90,9 +90,10 @@
 // native handles.
 // 12.29 Support PI_EXT_PLATFORM_INFO_BACKEND query in piPlatformGetInfo
 // 12.30 Added PI_EXT_INTEL_DEVICE_INFO_MEM_CHANNEL_SUPPORT device info query.
+// 12.31 Added command-buffer extension methods
 
 #define _PI_H_VERSION_MAJOR 12
-#define _PI_H_VERSION_MINOR 30
+#define _PI_H_VERSION_MINOR 31
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -468,6 +469,7 @@ typedef enum {
   PI_COMMAND_TYPE_SVM_MEMFILL = 0x120B,
   PI_COMMAND_TYPE_SVM_MAP = 0x120C,
   PI_COMMAND_TYPE_SVM_UNMAP = 0x120D,
+  PI_COMMAND_TYPE_EXT_COMMAND_BUFFER = 0x12A8,
   PI_COMMAND_TYPE_DEVICE_GLOBAL_VARIABLE_READ = 0x418E,
   PI_COMMAND_TYPE_DEVICE_GLOBAL_VARIABLE_WRITE = 0x418F
 } _pi_command_type;
@@ -2121,6 +2123,84 @@ __SYCL_EXPORT pi_result piPluginGetBackendOption(pi_platform platform,
 __SYCL_EXPORT pi_result piGetDeviceAndHostTimer(pi_device Device,
                                                 uint64_t *DeviceTime,
                                                 uint64_t *HostTime);
+                                                
+/// Command buffer extension
+struct _pi_ext_command_buffer;
+struct _pi_ext_sync_point;
+using pi_ext_command_buffer = _pi_ext_command_buffer *;
+using pi_ext_sync_point = pi_uint32;
+
+typedef enum {
+  PI_EXT_STRUCTURE_TYPE_COMMAND_BUFFER_DESC = 0
+} pi_ext_structure_type;
+
+struct pi_ext_command_buffer_desc final {
+  pi_ext_structure_type stype;
+  const void *pNext;
+  pi_queue_properties *properties;
+};
+
+/// API to create a command-buffer.
+/// \param context The context to associate the command-buffer with.
+/// \param device The device to associate the command-buffer with.
+/// \param desc Descriptor for the new command-buffer.
+/// \param ret_command_buffer Pointer to fill with the address of the new
+/// command-buffer.
+__SYCL_EXPORT pi_result
+piextCommandBufferCreate(pi_context context, pi_device device,
+                         const pi_ext_command_buffer_desc *desc,
+                         pi_ext_command_buffer *ret_command_buffer);
+
+/// API to increment the reference count of the command-buffer
+/// \param command_buffer The command_buffer to retain.
+__SYCL_EXPORT pi_result
+piextCommandBufferRetain(pi_ext_command_buffer command_buffer);
+
+/// API to decrement the reference count of the command-buffer. After the
+/// command_buffer reference count becomes zero and has finished execution, the
+/// command-buffer is deleted. \param command_buffer The command_buffer to
+/// release.
+__SYCL_EXPORT pi_result
+piextCommandBufferRelease(pi_ext_command_buffer command_buffer);
+
+/// API to stop command-buffer recording such that no more commands can be
+/// appended, and makes the command-buffer ready to enqueue on a command-queue.
+/// \param command_buffer The command_buffer to finalize.
+__SYCL_EXPORT pi_result
+piextCommandBufferFinalize(pi_ext_command_buffer command_buffer);
+
+/// API to append a kernel execution command to the command-buffer.
+/// \param command_buffer The command-buffer to append onto.
+/// \param kernel The kernel to append.
+/// \param work_dim Dimension of the kernel execution.
+/// \param global_work_offset Offset to use when executing kernel.
+/// \param global_work_size Global work size to use when executing kernel.
+/// \param local_work_size Local work size to use when executing kernel.
+/// \param num_sync_points_in_wait_list The number of sync points in the
+/// provided wait list.
+/// \param sync_point_wait_list A list of sync points that this executions must
+/// wait on.
+/// \param sync_point The sync_point associated with this kernel execution.
+__SYCL_EXPORT pi_result piextCommandBufferNDRangeKernel(
+    pi_ext_command_buffer command_buffer, pi_kernel kernel, pi_uint32 work_dim,
+    const size_t *global_work_offset, const size_t *global_work_size,
+    const size_t *local_work_size, pi_uint32 num_sync_points_in_wait_list,
+    const pi_ext_sync_point *sync_point_wait_list,
+    pi_ext_sync_point *sync_point);
+
+/// API to submit the command-buffer to queue for execution, returns an error if
+/// command-buffer not finalized or another instance of same command-buffer
+/// currently executing.
+/// \param command_buffer The command-buffer to be submitted.
+/// \param queue The PI queue to submit on.
+/// \param num_events_in_wait_list The number of events that this execution
+/// depends on.
+/// \param event_wait_list List of pi_events to wait on.
+/// \param event The pi_event associated with this enqueue.
+__SYCL_EXPORT pi_result
+piextEnqueueCommandBuffer(pi_ext_command_buffer command_buffer, pi_queue queue,
+                          pi_uint32 num_events_in_wait_list,
+                          const pi_event *event_wait_list, pi_event *event);
 
 struct _pi_plugin {
   // PI version supported by host passed to the plugin. The Plugin