Merge from 'sycl' to 'sycl-web' (2 commits)

  CONFLICT (content): Merge conflict in libclc/CMakeLists.txt
  CONFLICT (content): Merge conflict in libclc/clc/include/clc/clcfunc.h
  CONFLICT (content): Merge conflict in libclc/cmake/modules/AddLibclc.cmake
  CONFLICT (content): Merge conflict in libclc/generic/include/config.h
  CONFLICT (content): Merge conflict in libclc/generic/include/core/clc_core.h
  CONFLICT (content): Merge conflict in libclc/generic/include/math/math.h
  CONFLICT (content): Merge conflict in libclc/generic/lib/math/math.h
  CONFLICT (content): Merge conflict in libclc/native_cpu-unknown-linux/libspirv/math/helpers.h

Author: e-kud

libclc/CMakeLists.txt

@@ -35,7 +35,7 @@ set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS
   r600/libspirv/SOURCES;
   spirv/lib/SOURCES;
   spirv64/lib/SOURCES;
-  native_cpu-unknown-linux/libspirv/SOURCES
+  native_cpu-unknown-linux/libspirv/SOURCES;
   # CLC internal libraries
   clc/lib/generic/SOURCES;
   clc/lib/clspv/SOURCES;
@@ -534,8 +534,9 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
       LIB_FILES ${libspirv_lib_files}
       GEN_FILES ${libspirv_gen_files}
       ALIASES ${${d}_aliases}
-      GENERATE_TARGET "generate_convert_spirv.cl" "generate_convert_core.cl"
       PARENT_TARGET libspirv-builtins
+      # Link in the CLC builtins and internalize their symbols
+      INTERNAL_LINK_DEPENDENCIES $<TARGET_PROPERTY:builtins.link.clc-${arch_suffix},TARGET_FILE>
     )
 
     add_libclc_builtin_set(
@@ -547,7 +548,6 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
       LIB_FILES ${opencl_lib_files}
       GEN_FILES ${opencl_gen_files}
       ALIASES ${${d}_aliases}
-      GENERATE_TARGET "generate_convert_clc.cl"
       PARENT_TARGET libopencl-builtins
       # Link in the CLC builtins and internalize their symbols
       INTERNAL_LINK_DEPENDENCIES $<TARGET_PROPERTY:builtins.link.clc-${arch_suffix},TARGET_FILE>

libclc/clc/include/clc/clcfunc.h

@@ -3,6 +3,11 @@
 
 #define _CLC_OVERLOAD __attribute__((overloadable))
 #define _CLC_DECL
+#define _CLC_INLINE __attribute__((always_inline)) inline
+#define _CLC_CONVERGENT __attribute__((convergent))
+#define _CLC_PURE __attribute__((pure))
+#define _CLC_CONSTFN __attribute__((const))
+
 // avoid inlines for SPIR-V related targets since we'll optimise later in the
 // chain
 #if defined(CLC_SPIRV) || defined(CLC_SPIRV64)

libclc/cmake/modules/AddLibclc.cmake

@@ -361,7 +361,6 @@ function(add_libclc_builtin_set)
 
   set( builtins_link_lib_tgt builtins.link.${ARG_ARCH_SUFFIX} )
 
-
   if( NOT ARG_INTERNAL_LINK_DEPENDENCIES )
     link_bc(
       TARGET ${builtins_link_lib_tgt}

libclc/generic/include/config.h

@@ -20,7 +20,7 @@
  * THE SOFTWARE.
  */
 
-#include "clc/clcfunc.h"
+#include <clc/clcfunc.h>
 
 _CLC_DECL bool __clc_subnormals_disabled();
 _CLC_DECL bool __clc_fp16_subnormals_supported();

libclc/generic/include/core/clc_core.h

@@ -30,8 +30,8 @@
 
 #include <as_type.h>
 #include <clc/clcfunc.h>
-#include <macros.h>
 #include <clc/clctypes.h>
+#include <macros.h>
 
 #include <clc/float/definitions.h>
 #include <clc/integer/definitions.h>

libclc/generic/include/math/math.h

@@ -11,7 +11,7 @@
 
 #include "as_type.h"
 #include "config.h"
-#include "clc/clcfunc.h"
+#include <clc/clcfunc.h>
 
 #define SNAN 0x001
 #define QNAN 0x002

libclc/generic/lib/math/math.h

@@ -23,9 +23,9 @@
 #ifndef __CLC_MATH_H_
 #define __CLC_MATH_H_
 
-#include "clc/clcfunc.h"
 #include "as_type.h"
 #include "config.h"
+#include <clc/clcfunc.h>
 
 #define SNAN 0x001
 #define QNAN 0x002

libclc/native_cpu-unknown-linux/libspirv/math/helpers.h

@@ -1,5 +1,5 @@
-#include "clc/clcfunc.h"
-#include "clc/clctypes.h"
+#include <clc/clcfunc.h>
+#include <clc/clctypes.h>
 
 #ifdef cl_khr_fp16
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable

sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc

@@ -233,6 +233,9 @@ namespace sycl::ext::oneapi::experimental {
 template <typename CommandGroupFunc>
 void submit(sycl::queue q, CommandGroupFunc&& cgf);
 
+template <typename CommandGroupFunc, typename Properties>
+void submit(sycl::queue q, Properties properties, CommandGroupFunc&& cgf);
+
 }
 ----
 !====
@@ -250,6 +253,10 @@ namespace sycl::ext::oneapi::experimental {
 template <typename CommandGroupFunc>
 sycl::event submit_with_event(sycl::queue q, CommandGroupFunc&& cgf);
 
+template <typename CommandGroupFunc, typename Properties>
+sycl::event submit_with_event(sycl::queue q, Properties properties,
+                              CommandGroupFunc&& cgf);
+
 }
 ----
 !====

sycl/doc/extensions/proposed/sycl_ext_intel_event_mode.asciidoc

@@ -0,0 +1,165 @@
+= sycl_ext_intel_event_mode
+
+: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
+:dpcpp: pass:[DPC++]
+:endnote: —{nbsp}end{nbsp}note
+
+// 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}
+
+:common_ref_sem: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:reference-semantics
+
+== Notice
+
+[%hardbreaks]
+Copyright (C) 2024 Intel Corporation.  All rights reserved.
+
+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.
+
+
+== Contact
+
+To report problems with this extension, please open a new issue at:
+
+https://github.com/intel/llvm/issues
+
+
+== Dependencies
+
+This extension is written against the SYCL 2020 revision 9 specification.  All
+references below to the "core SYCL specification" or to section numbers in the
+SYCL specification refer to that revision.
+
+This extension also depends on the following other SYCL extensions:
+
+* link:../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc[
+  sycl_ext_oneapi_enqueue_functions]
+* link:../experimental/sycl_ext_oneapi_properties.asciidoc[
+  sycl_ext_oneapi_properties]
+
+
+== Status
+
+This is a proposed extension specification, intended to gather community
+feedback.  Interfaces defined in this specification may not be implemented yet
+or may be in a preliminary state.  The specification itself may also change in
+incompatible ways before it is finalized.  *Shipping software products should
+not rely on APIs defined in this specification.*
+
+
+== Overview
+
+On some backends, calling `wait()` on an `event` will synchronize using a
+busy-waiting implementation. Though this comes at a low latency for the
+synchronization of the event, it has the downside of consuming high amounts of
+CPU time for no meaningful work.  This extension introduces a new property for
+SYCL commands that allow users to pick modes for the associated events, one of
+these modes being a "low-power" event. These new low-power events will, if
+possible, yield the thread that the `wait()` member function is called on and
+only wake up occasionally to check if the event has finished. This reduces the
+time the CPU spends checking finish condition of the wait, at the cost of
+latency.
+
+
+== Specification
+
+=== Feature test macro
+
+This extension provides a feature-test macro as described in the core SYCL
+specification.  An implementation supporting this extension must predefine the
+macro `SYCL_EXT_INTEL_EVENT_MODE` 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 features the implementation
+supports.
+
+[%header,cols="1,5"]
+|===
+|Value
+|Description
+
+|1
+|The APIs of this experimental extension are not versioned, so the
+ feature-test macro always has this value.
+|===
+
+
+=== Event mode property
+
+This extension adds a new property `event_mode` which can be used with the
+`submit_with_event` free function from
+link:../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc[sycl_ext_oneapi_enqueue_functions],
+allowing the user some control over how the resulting event is created and
+managed.
+
+```
+namespace sycl::ext::intel::experimental {
+
+enum class event_mode_enum { none, low_power };
+
+struct event_mode {
+  event_mode(event_mode_enum mode);
+
+  event_mode_enum value;
+};
+
+using event_mode_key = event_mode;
+
+} // namespace sycl::ext::intel::experimental
+```
+
+
+=== Low power event mode
+
+Passing the `event_mode` property with `event_mode_enum::low_power` to
+`submit_with_event` will act as a hint to the `event` created from the
+corresponding commands to do low-power synchronization. If the backend is able
+to handle low-power events, calling `event::wait()` or `event::wait_and_throw()`
+will cause the thread to yield and only do occasional wake-ups to check the
+event progress.
+
+[_Note:_ The low-power event mode currently only has an effect on `barrier` and
+`partial_barrier` commands enqueued on queues that return
+`backend::ext_oneapi_level_zero` from `queue::get_backend()`.
+_{endnote}_]
+
+
+=== New property usage example
+
+As an example of how to use the new `event_mode` property using the
+`event_mode_enum::low_power` mode, see the following code:
+
+```
+#include <sycl/sycl.hpp>
+
+namespace oneapiex = sycl::ext::oneapi::experimental;
+namespace intelex = sycl::ext::intel::experimental;
+
+int main() {
+  sycl::queue Q;
+
+  // Submit some work to the queue.
+  oneapiex::submit(Q, [&](sycl::handler &CGH) {...});
+
+  // Submit a command with the low-power event mode.
+  oneapiex::properties Props{intelex::event_mode{intelex::event_mode_enum::low_power}};
+  sycl::event E = oneapiex::submit_with_event(Q, Props, [&](sycl::handler &CGH) {
+    ...
+  });
+
+  // Waiting for the resulting event will use low-power waiting if possible.
+  E.wait();
+}
+```

sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp

@@ -95,21 +95,38 @@ void submit_impl(queue &Q, CommandGroupFunc &&CGF,
 }
 } // namespace detail
 
-template <typename CommandGroupFunc>
-void submit(queue Q, CommandGroupFunc &&CGF,
+template <typename CommandGroupFunc, typename PropertiesT>
+void submit(queue Q, PropertiesT Props, CommandGroupFunc &&CGF,
             const sycl::detail::code_location &CodeLoc =
                 sycl::detail::code_location::current()) {
+  std::ignore = Props;
   sycl::ext::oneapi::experimental::detail::submit_impl(
       Q, std::forward<CommandGroupFunc>(CGF), CodeLoc);
 }
 
 template <typename CommandGroupFunc>
-event submit_with_event(queue Q, CommandGroupFunc &&CGF,
+void submit(queue Q, CommandGroupFunc &&CGF,
+            const sycl::detail::code_location &CodeLoc =
+                sycl::detail::code_location::current()) {
+  submit(Q, empty_properties_t{}, std::forward<CommandGroupFunc>(CGF), CodeLoc);
+}
+
+template <typename CommandGroupFunc, typename PropertiesT>
+event submit_with_event(queue Q, PropertiesT Props, CommandGroupFunc &&CGF,
                         const sycl::detail::code_location &CodeLoc =
                             sycl::detail::code_location::current()) {
+  std::ignore = Props;
   return Q.submit(std::forward<CommandGroupFunc>(CGF), CodeLoc);
 }
 
+template <typename CommandGroupFunc>
+event submit_with_event(queue Q, CommandGroupFunc &&CGF,
+                        const sycl::detail::code_location &CodeLoc =
+                            sycl::detail::code_location::current()) {
+  return submit_with_event(Q, empty_properties_t{},
+                           std::forward<CommandGroupFunc>(CGF), CodeLoc);
+}
+
 template <typename KernelName = sycl::detail::auto_name, typename KernelType>
 void single_task(handler &CGH, const KernelType &KernelObj) {
   CGH.single_task<KernelName>(KernelObj);