Merge remote-tracking branch 'origin/sycl' into llvmspirv_pulldown

Author: sys-ce-bb

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -62,7 +62,10 @@ def AspectExt_oneapi_mipmap : Aspect<"ext_oneapi_mipmap">;
 def AspectExt_oneapi_mipmap_anisotropy : Aspect<"ext_oneapi_mipmap_anisotropy">;
 def AspectExt_oneapi_mipmap_level_reference : Aspect<"ext_oneapi_mipmap_level_reference">;
 def AspectExt_intel_esimd : Aspect<"ext_intel_esimd">;
-def AspectExt_oneapi_non_uniform_groups : Aspect<"ext_oneapi_non_uniform_groups">;
+def AspectExt_oneapi_ballot_group : Aspect<"ext_oneapi_ballot_group">;
+def AspectExt_oneapi_fixed_size_group : Aspect<"ext_oneapi_fixed_size_group">;
+def AspectExt_oneapi_opportunistic_group : Aspect<"ext_oneapi_opportunistic_group">;
+def AspectExt_oneapi_tangle_group : Aspect<"ext_oneapi_tangle_group">;
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
 def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">;
@@ -111,7 +114,8 @@ def : TargetInfo<"__TestAspectList",
     AspectExt_oneapi_bindless_images_shared_usm, AspectExt_oneapi_bindless_images_1d_usm, AspectExt_oneapi_bindless_images_2d_usm,
     AspectExt_oneapi_interop_memory_import, AspectExt_oneapi_interop_memory_export,
     AspectExt_oneapi_interop_semaphore_import, AspectExt_oneapi_interop_semaphore_export,
-    AspectExt_oneapi_mipmap, AspectExt_oneapi_mipmap_anisotropy, AspectExt_oneapi_mipmap_level_reference, AspectExt_intel_esimd, AspectExt_oneapi_non_uniform_groups],
+    AspectExt_oneapi_mipmap, AspectExt_oneapi_mipmap_anisotropy, AspectExt_oneapi_mipmap_level_reference, AspectExt_intel_esimd,
+    AspectExt_oneapi_ballot_group, AspectExt_oneapi_fixed_size_group, AspectExt_oneapi_opportunistic_group, AspectExt_oneapi_tangle_group],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.

llvm/lib/SYCLLowerIR/CompileTimeProperties.def

@@ -42,7 +42,7 @@ SYCL_COMPILE_TIME_PROPERTY("sycl-latency-constraint", 6173,
 SYCL_COMPILE_TIME_PROPERTY("sycl-datapath", 5825, DecorValueTy::none)
 SYCL_COMPILE_TIME_PROPERTY("sycl-resource", 5826, DecorValueTy::string)
 SYCL_COMPILE_TIME_PROPERTY("sycl-num-banks", 5827, DecorValueTy::uint32)
-SYCL_COMPILE_TIME_PROPERTY("sycl-ram-stitching", 5836, DecorValueTy::boolean)
+SYCL_COMPILE_TIME_PROPERTY("sycl-ram-stitching", 5836, DecorValueTy::uint32)
 SYCL_COMPILE_TIME_PROPERTY("sycl-max-private-copies", 5829,
                            DecorValueTy::uint32)
 SYCL_COMPILE_TIME_PROPERTY("sycl-num-replicates", 5832, DecorValueTy::uint32)

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -36,6 +36,7 @@ static const char *LegalSYCLFunctions[] = {
     "^sycl::_V1::accessor<.+>::~accessor",
     "^sycl::_V1::accessor<.+>::getQualifiedPtr",
     "^sycl::_V1::accessor<.+>::__init_esimd",
+    "^sycl::_V1::address_space_cast",
     "^sycl::_V1::local_accessor<.+>::local_accessor",
     "^sycl::_V1::local_accessor<.+>::__init_esimd",
     "^sycl::_V1::local_accessor<.+>::get_pointer",

sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc

@@ -37,7 +37,7 @@ https://github.com/intel/llvm/issues
 
 == Dependencies
 
-This extension is written against the SYCL 2020 revision 7 specification.
+This extension is written against the SYCL 2020 revision 8 specification.
 All references below to the "core SYCL specification" or to section numbers in
 the SYCL specification refer to that revision.
 
@@ -215,13 +215,13 @@ a!
 ----
 namespace sycl::ext::oneapi::experimental {
 
-template<typename PropertyListT = properties<>>                       (1)
+template<typename PropertyListT = empty_properties_t>                 (1)
 kernel_bundle<bundle_state::executable> build(
   const kernel_bundle<bundle_state::ext_oneapi_source>& sourceBundle,
   const std::vector<device> &devs,
   PropertyListT props = {})
 
-template<typename PropertyListT = properties<>>                       (2)
+template<typename PropertyListT = empty_properties_t>                 (2)
 kernel_bundle<bundle_state::executable> build(
   const kernel_bundle<bundle_state::ext_oneapi_source>& sourceBundle,
   PropertyListT props = {})
@@ -230,13 +230,11 @@ kernel_bundle<bundle_state::executable> build(
 ----
 !====
 
-Overload (1):
-
 _Constraints:_ Available only when `PropertyListT` is an instance of
 `sycl::ext::oneapi::experimental::properties` which contains no properties
 other than those listed below in the section "New properties".
 
-_Effects:_ The source code from `sourceBundle` is translated into one or more
+_Effects (1):_ The source code from `sourceBundle` is translated into one or more
 device images of state `bundle_state::executable`, and a new kernel bundle is
 created to contain these device images.
 The new bundle represents all of the kernels in `sourceBundle` that are
@@ -247,6 +245,8 @@ Any remaining kernels (those that are not compatible with any of the devices in
 The new bundle has the same associated context as `sourceBundle`, and the new
 bundle's set of associated devices is `devs` (with duplicate devices removed).
 
+_Effects (2)_: Equivalent to `build(sourceBundle, ctxt.get_devices(), props)`.
+
 _Returns:_ The newly created kernel bundle, which has `executable` state.
 
 _Throws:_
@@ -274,12 +274,6 @@ source code used to create the kernel bundle being printed to the terminal.
 In situations where this is undesirable, developers must ensure that the
 exception is caught and handled appropriately.
 _{endnote}_]
-
-Overload (2):
-
-_Constraints:_ Same as overload (1).
-
-_Effects:_ Equivalent to `build(sourceBundle, ctxt.get_devices(), props)`.
 |====
 
 === New properties
@@ -333,7 +327,7 @@ namespace sycl::ext::oneapi::experimental {
 
 struct save_log {
   std::string *log;
-  save_log(std::string *to);
+  save_log(std::string *to);  (1)
 };
 using save_log_key = save_log;
 
@@ -359,7 +353,7 @@ provided in the log.
 In general, the log information is intended for human consumption, and the
 format may not be stable across implementations of this extension.
 
-_Effects:_ Constructs a `save_log` property with a pointer to a `std::string`.
+_Effects (1):_ Constructs a `save_log` property with a pointer to a `std::string`.
 When the `build` function completes successfully, this string will contain the
 log.
 
@@ -426,10 +420,7 @@ template <bundle_state State>
 class kernel_bundle {
   // ...
 
-  // Available only if bundle_state is not bundle_state::ext_oneapi_source
   bool ext_oneapi_has_kernel(const std::string &name);
-
-  // Available only if bundle_state is bundle_state::executable
   kernel ext_oneapi_get_kernel(const std::string &name);
 };
 

sycl/doc/extensions/proposed/sycl_ext_oneapi_non_uniform_groups.asciidoc renamed to sycl/doc/extensions/experimental/sycl_ext_oneapi_non_uniform_groups.asciidoc

@@ -44,17 +44,18 @@ SYCL specification refer to that revision.
 
 == 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.*
+This is an experimental extension specification, intended to provide early
+access to features and gather community feedback.  Interfaces defined in this
+specification are implemented in {dpcpp}, but they are not finalized and may
+change incompatibly in future versions of {dpcpp} without prior notice.
+*Shipping software products should not rely on APIs defined in this
+specification.*
 
 
 == Backend support status
 
-The APIs in this extension may be used only on a device that has
-`aspect::ext_oneapi_non_uniform_groups`. The application must check that the
+The APIs in this extension may be used only on a device that has one or more of
+the xref:ext-aspects[extension aspects]. The application must check that the
 device has this aspect before submitting a kernel using any of the APIs in this
 extension. If the application fails to do this, the implementation throws a
 synchronous exception with the `errc::kernel_not_supported` error code when the
@@ -105,20 +106,33 @@ implementation supports.
  feature-test macro always has this value.
 |===
 
+[#ext-aspects]
 === Extension to `enum class aspect`
 
 [source]
 ----
 namespace sycl {
 enum class aspect {
   ...
-  ext_oneapi_non_uniform_groups
+  ext_oneapi_ballot_group
+  ext_oneapi_fixed_size_group
+  ext_oneapi_opportunistic_group
+  ext_oneapi_tangle_group
 }
 }
 ----
 
-If a SYCL device has the `ext_oneapi_non_uniform_groups` aspect,
-then it supports the non-uniform groups described in the next sections.
+If a SYCL device has these aspects, that device supports the non-uniform groups
+as follows:
+
+[%header,cols="2,3"]
+|===
+| Aspect | Supported group
+| `ext_oneapi_ballot_group` | `ballot_group`
+| `ext_oneapi_fixed_size_group` | `fixed_size_group`
+| `ext_oneapi_opportunistic_group` | `opportunistic_group`
+| `ext_oneapi_tangle_group` | `tangle_group`
+|===
 
 === Control Flow
 

sycl/include/sycl/device_aspect_macros.hpp

@@ -278,9 +278,24 @@
 #define __SYCL_ALL_DEVICES_HAVE_ext_intel_esimd__ 0
 #endif
 
-#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_non_uniform_groups__
-// __SYCL_ASPECT(ext_oneapi_non_uniform_groups, 54)
-#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_non_uniform_groups__ 0
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_ballot_group__
+// __SYCL_ASPECT(ext_oneapi_ballot_group, 54)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_ballot_group__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_fixed_size_group__
+// __SYCL_ASPECT(ext_oneapi_fixed_size_group, 55)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_fixed_size_group__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_opportunistic_group__
+// __SYCL_ASPECT(ext_oneapi_opportunistic_group, 56)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_opportunistic_group__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_tangle_group__
+// __SYCL_ASPECT(ext_oneapi_tangle_group, 57)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_tangle_group__ 0
 #endif
 
 #ifndef __SYCL_ANY_DEVICE_HAS_host__
@@ -553,7 +568,22 @@
 #define __SYCL_ANY_DEVICE_HAS_ext_intel_esimd__ 0
 #endif
 
-#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_non_uniform_groups__
-// __SYCL_ASPECT(ext_oneapi_non_uniform_groups, 54)
-#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_non_uniform_groups__ 0
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_ballot_group__
+// __SYCL_ASPECT(ext_oneapi_ballot_group, 54)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_ballot_group__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_fixed_size_group__
+// __SYCL_ASPECT(ext_oneapi_fixed_size_group, 55)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_fixed_size_group__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_opportunistic_group__
+// __SYCL_ASPECT(ext_oneapi_opportunistic_group, 56)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_opportunistic_group__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_tangle_group__
+// __SYCL_ASPECT(ext_oneapi_tangle_group, 57)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_tangle_group__ 0
 #endif