[SYCL] Kernel property to control Cache/SLM size on GPU (#8597)

When L1 cache & SLM share same physical memory module, developers may
want more L1 or SLM based on their application. So, the purpose of the
introduced kernel property is to give developers flexibility to tune the
division.

E2E test: intel/llvm-test-suite#1687

Author: againull

clang/lib/Sema/SemaDeclAttr.cpp

@@ -8025,10 +8025,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 : 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;
+};
+
+} // 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