Merge from 'sycl' to 'sycl-web'

Author: iclsrc

clang/lib/Sema/SemaDeclAttr.cpp

@@ -8032,10 +8032,12 @@ void Sema::CheckSYCLAddIRAttributesFunctionAttrConflicts(Decl *D) {
   // "sycl-single-task" is present on all single_task invocations, implicitly
   // added by the SYCL headers. It can only conflict with max_global_work_dim,
   // but the value will be the same so there is no need for a warning.
-  if (NumArgsWithoutFilter == 2 &&
-      AddIRFuncAttr->getAttributeNameValuePairs(Context)[0].first ==
-          "sycl-single-task")
-    return;
+  if (NumArgsWithoutFilter == 2) {
+    auto NameValuePairs = AddIRFuncAttr->getAttributeNameValuePairs(Context);
+    if (NameValuePairs.size() > 0 &&
+        NameValuePairs[0].first == "sycl-single-task")
+      return;
+  }
 
   // If there are potentially conflicting attributes, we issue a warning.
   for (const auto *Attr : std::vector<AttributeCommonInfo *>{

sycl/doc/extensions/experimental/sycl_ext_intel_cache_config.asciidoc

@@ -0,0 +1,183 @@
+= sycl_ext_intel_cache_config
+: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++]
+
+: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}
+
+== Notice
+
+[%hardbreaks]
+Copyright (c) 2023-2023 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
+
+== Contributors
+
+Greg Lueck, Intel +
+John Pennycook, Intel +
+Artur Gainullin, Intel
+
+== Dependencies
+
+This extension is written against the SYCL 2020 specification, Revision 6 and
+the following extensions:
+
+- link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties]
+- link:sycl_ext_oneapi_kernel_properties.asciidoc[sycl_ext_oneapi_kernel_properties]
+
+== Status
+
+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.*
+
+== Overview
+
+There are devices where the same hardware resources are used for shared local
+memory (SLM) and L1 data cache. Developers may want more L1 data cache or SLM based
+on their application. This extension adds runtime kernel property `cache_config`
+which provides a way to set the preferred cache configuration for a kernel.
+
+=== 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_INTEL_CACHE_CONFIG` 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.
+|===
+
+=== Introduction
+
+This extension introduces new kernel property that can be applied to kernels
+using the mechanism defined in link:sycl_ext_oneapi_kernel_properties.asciidoc[sycl_ext_oneapi_kernel_properties].
+
+=== Cache Config Property
+
+```c++
+namespace sycl::ext::intel::experimental {
+
+
+enum class cache_config_enum : /*unspecified*/ { large_slm, large_data };
+
+inline constexpr cache_config_enum large_slm =
+    cache_config_enum::large_slm;
+inline constexpr cache_config_enum large_data =
+    cache_config_enum::large_data;
+
+struct cache_config {
+  cache_config(cache_config_enum v) : value(v) {}
+  cache_config_enum value;
+};
+
+} // namespace sycl::ext::intel::experimental
+```
+
+The `cache_config` property provides a way to set the preferred cache
+configuration for a kernel. The following values are supported:
+
+ * `large_slm`: Prefer having larger shared local memory and smaller L1 data cache.
+                In this case driver will ensure that all workgroups will have enough
+                SLM to run.
+
+ * `large_data`: Prefer having larger L1 data cache and smaller shared local memory.
+                 In this case SLM size may be shrinked (which may result in workgroups
+                 spawning as there will be not enough SLM to handle multiple workgroups)
+                 but L1 data cache will be bigger. There may be rare use cases when this
+                 is beneficial.
+
+These property may be passed to any kernel invocation function (e.g.
+`parallel_for`) via the properties parameter.  At most, only one of these
+values may be passed to any single kernel invocation function.
+
+Backends that do not support this extension may accept and ignore this
+property.
+
+=== Adding a Property List to a Kernel Launch
+
+A simple example of using this extension  is shown below.
+
+The example assumes that the kernel will benefit from large SLM and hence uses the property
+`cache_config_large_slm`:
+
+```c++
+using namespace sycl::ext::intel::experimental;
+{
+  ...
+  properties kernel_properties{cache_config{large_slm}};
+
+  q.single_task(kernel_properties, [=] {
+    *a = *b + *c;
+  }).wait();
+}
+```
+
+=== Embedding Property into a Kernel
+
+The example below shows how the kernel from the previous section could be
+rewritten to leverage an embedded property list (see link:sycl_ext_oneapi_kernel_properties.asciidoc#embedding-properties-into-a-kernel[embedding-properties-into-a-kernel]):
+
+```c++
+using namespace sycl::ext::intel::experimental;
+struct KernelFunctor {
+
+  KernelFunctor(int* a, int* b, int* c) : a(a), b(b), c(c) {}
+
+  void operator()() const {
+    *a = *b + *c;
+  }
+
+  auto get(properties_tag) const {
+    return properties{cache_config{large_slm}};
+  }
+
+  int* a;
+  int* b;
+  int* c;
+};
+
+...
+
+q.single_task(KernelFunctor{a, b, c}).wait();
+```
+
+== Revision History
+
+[cols="5,15,15,70"]
+[grid="rows"]
+[options="header"]
+|========================================
+|Rev|Date|Author|Changes
+|1|2022-03-01|Artur Gainullin|*Initial public working draft*
+|========================================

sycl/include/sycl/detail/cg.hpp

@@ -144,6 +144,7 @@ class CGExecKernel : public CG {
   detail::OSModuleHandle MOSModuleHandle;
   std::vector<std::shared_ptr<detail::stream_impl>> MStreams;
   std::vector<std::shared_ptr<const void>> MAuxiliaryResources;
+  RT::PiKernelCacheConfig MKernelCacheConfig;
 
   CGExecKernel(NDRDescT NDRDesc, std::unique_ptr<HostKernelBase> HKernel,
                std::shared_ptr<detail::kernel_impl> SyclKernel,
@@ -157,7 +158,8 @@ class CGExecKernel : public CG {
                detail::OSModuleHandle OSModuleHandle,
                std::vector<std::shared_ptr<detail::stream_impl>> Streams,
                std::vector<std::shared_ptr<const void>> AuxiliaryResources,
-               CGTYPE Type, detail::code_location loc = {})
+               CGTYPE Type, RT::PiKernelCacheConfig KernelCacheConfig,
+               detail::code_location loc = {})
       : CG(Type, std::move(ArgsStorage), std::move(AccStorage),
            std::move(SharedPtrStorage), std::move(Requirements),
            std::move(Events), std::move(loc)),
@@ -166,7 +168,8 @@ class CGExecKernel : public CG {
         MKernelBundle(std::move(KernelBundle)), MArgs(std::move(Args)),
         MKernelName(std::move(KernelName)), MOSModuleHandle(OSModuleHandle),
         MStreams(std::move(Streams)),
-        MAuxiliaryResources(std::move(AuxiliaryResources)) {
+        MAuxiliaryResources(std::move(AuxiliaryResources)),
+        MKernelCacheConfig(std::move(KernelCacheConfig)) {
     assert((getType() == RunOnHostIntel || getType() == Kernel) &&
            "Wrong type of exec kernel CG.");
   }

sycl/include/sycl/detail/pi.h

@@ -77,9 +77,12 @@
 // 12.22 Add piGetDeviceAndHostTimer to query device wall-clock timestamp
 // 12.23 Added new piextEnqueueDeviceGlobalVariableWrite and
 // piextEnqueueDeviceGlobalVariableRead functions.
+// 12.24 Added new PI_EXT_KERNEL_EXEC_INFO_CACHE_CONFIG property to the
+// _pi_kernel_exec_info. Defined _pi_kernel_cache_config enum with values of
+// the new PI_EXT_KERNEL_EXEC_INFO_CACHE_CONFIG property.
 
 #define _PI_H_VERSION_MAJOR 12
-#define _PI_H_VERSION_MINOR 23
+#define _PI_H_VERSION_MINOR 24
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -621,6 +624,15 @@ constexpr pi_queue_properties PI_EXT_ONEAPI_QUEUE_FLAG_PRIORITY_LOW = (1 << 5);
 constexpr pi_queue_properties PI_EXT_ONEAPI_QUEUE_FLAG_PRIORITY_HIGH = (1 << 6);
 // clang-format on
 
+typedef enum {
+  // No preference for SLM or data cache.
+  PI_EXT_KERNEL_EXEC_INFO_CACHE_DEFAULT = 0x0,
+  // Large SLM size.
+  PI_EXT_KERNEL_EXEC_INFO_CACHE_LARGE_SLM = 0x1,
+  // Large General Data size.
+  PI_EXT_KERNEL_EXEC_INFO_CACHE_LARGE_DATA = 0x2
+} _pi_kernel_cache_config;
+
 using pi_result = _pi_result;
 using pi_platform_info = _pi_platform_info;
 using pi_device_type = _pi_device_type;
@@ -650,6 +662,7 @@ using pi_program_build_status = _pi_program_build_status;
 using pi_program_binary_type = _pi_program_binary_type;
 using pi_kernel_info = _pi_kernel_info;
 using pi_profiling_info = _pi_profiling_info;
+using pi_kernel_cache_config = _pi_kernel_cache_config;
 
 // For compatibility with OpenCL define this not as enum.
 using pi_device_partition_property = intptr_t;
@@ -1357,7 +1370,9 @@ typedef enum {
   /// indicates that the kernel might access data through USM ptrs
   PI_USM_INDIRECT_ACCESS,
   /// provides an explicit list of pointers that the kernel will access
-  PI_USM_PTRS = 0x4203
+  PI_USM_PTRS = 0x4203,
+  /// provides the preferred cache configuration (large slm or large data)
+  PI_EXT_KERNEL_EXEC_INFO_CACHE_CONFIG = 0x4204
 } _pi_kernel_exec_info;
 
 using pi_kernel_exec_info = _pi_kernel_exec_info;

sycl/include/sycl/detail/pi.hpp

@@ -145,6 +145,7 @@ using PiMemImageInfo = ::pi_image_info;
 using PiMemObjectType = ::pi_mem_type;
 using PiMemImageChannelOrder = ::pi_image_channel_order;
 using PiMemImageChannelType = ::pi_image_channel_type;
+using PiKernelCacheConfig = ::pi_kernel_cache_config;
 
 __SYCL_EXPORT void contextSetExtendedDeleter(const sycl::context &constext,
                                              pi_context_extended_deleter func,

sycl/include/sycl/ext/intel/experimental/kernel_execution_properties.hpp

@@ -0,0 +1,67 @@
+//==----- gpu_kernel_properties.hpp - Kernel properties for Intel GPUs ---==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===--------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/ext/oneapi/properties/property.hpp>
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace ext::intel::experimental {
+
+template <typename T, typename PropertyListT> class gpu_kernel_attribute;
+
+enum class cache_config_enum : std::uint16_t { large_slm, large_data };
+
+inline constexpr cache_config_enum large_slm =
+    cache_config_enum::large_slm;
+inline constexpr cache_config_enum large_data =
+    cache_config_enum::large_data;
+
+struct cache_config {
+  cache_config(cache_config_enum v) : value(v) {}
+  cache_config_enum value;
+};
+
+using cache_config_key = cache_config;
+
+inline bool operator==(const cache_config &lhs,
+                       const cache_config &rhs) {
+  return lhs.value == rhs.value;
+}
+inline bool operator!=(const cache_config &lhs,
+                       const cache_config &rhs) {
+  return !(lhs == rhs);
+}
+
+} // namespace ext::intel::experimental
+
+namespace ext::oneapi::experimental {
+template <>
+struct is_property_key<intel::experimental::cache_config_key>
+    : std::true_type {};
+
+template <typename T, typename PropertyListT>
+struct is_property_key_of<
+    intel::experimental::cache_config_key,
+    intel::experimental::gpu_kernel_attribute<T, PropertyListT>>
+    : std::true_type {};
+
+namespace detail {
+template <> struct PropertyToKind<intel::experimental::cache_config_key> {
+  static constexpr PropKind Kind = PropKind::CacheConfig;
+};
+
+template <>
+struct IsRuntimeProperty<intel::experimental::cache_config_key>
+    : std::true_type {};
+
+} // namespace detail
+} // namespace ext::oneapi::experimental
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl

sycl/include/sycl/ext/oneapi/properties/property.hpp

@@ -187,8 +187,9 @@ enum PropKind : uint32_t {
   MaxBurst = 21,
   WaitRequest = 22,
   Alignment = 23,
+  CacheConfig = 24,
   // PropKindSize must always be the last value.
-  PropKindSize = 24,
+  PropKindSize = 25,
 };
 
 // This trait must be specialized for all properties and must have a unique