From f19cdb3ebfd041b41644b5b1d59c181ca62723dc Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Wed, 30 Aug 2023 14:23:45 -0700 Subject: [PATCH 01/28] [SYCL] Proposed cache control properties for annotated_ptr. --- .../sycl_ext_intel_cache_controls.asciidoc | 201 ++++++++++++++++++ 1 file changed, 201 insertions(+) create mode 100755 sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc new file mode 100755 index 0000000000000..d068035b1477d --- /dev/null +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -0,0 +1,201 @@ += sycl_ext_intel_cache_controls + +: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++] + +// 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 + + +== Dependencies + +This extension is written against the SYCL 2020 revision 7 specification. All +references below to the "core SYCL specification" or to section numbers in the +SYCL specification refer to that revision. + +This extension depends on the following SYCL extensions: + +* link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties] +* link:../experimental/sycl_ext_oneapi_annotated_ptr.asciidoc[sycl_ext_oneapi_annotated_ptr] + +This extension also depends on the following SPIR-V extension: + +* link:../supported/sycl_ext_oneapi_myotherextension.asciidoc[SPV_INTEL_cache_controls] + + +== 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.* + + +== Backend support status + +This extension is not implemented. + +== Overview + +This extension introduces additional compile-time properties for +the proposed `sycl::ext::oneapi::experimental::annotated_ptr` class to specify +cache control information. + +The cache controls are a strong request that memory accesses through the +pointer should use instructions with the specified cache controls. +However, the implementation may choose a different cache control or none +if the requested one is unsupported or for any other reason. + + +== 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_CACHE_CONTROLS` 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. +|=== + +=== Properties + +Below is a list of new compile-time constant properties supported with +`annotated_ptr`. + +```c++ +namespace sycl::ext::intel::experimental { + +enum cache_control_types_read { + Cached, Uncached, Streaming, InvalidateAfterRead, ConstCached}; +enum cache_control_types_write { + Uncached, Streaming, WriteThrough, WriteBack}; + +struct cache_control_read { + template + using value_t = property_value>; +}; +struct cache_control_write { + template + using value_t = property_value>; +}; + +template +inline constexpr cache_control_read::value_t read_control; +template +inline constexpr cache_control_write::value_t write_control; + +template<> +struct is_property : std::true_type {}; +template<> +struct is_property : std::true_type {}; + +template +struct is_property_of< + cache_control_read, annotated_ptr> : std::true_type {}; +template +struct is_property_of< + cache_control_write, annotated_ptr> : std::true_type {}; + +template +} // namespace sycl::ext::intel::experimental +``` +-- +[options="header"] +|==== +| Property | Description +|`cache_control_read` +| +This property requests that loads from memory through the `annotated_ptr` +may cache the data at level L in the memory hierarchy. +|`cache_control_read` +| +This property requests that loads from memory through the `annotated_ptr` +should not cache the data at level L in the memory hierarchy. +|`cache_control_read` +| +This property requests that loads from memory through the `annotated_ptr` +should cache the data at cache level L. The eviction policy is to give +lower priority to data cached using this property versus the Cached +property. +|`cache_control_read` +| +This property asserts that the cache line into which data is loaded +from memory through the `annotated_ptr` will not be read again +until it is overwritten. Therefore the load operation can invalidate +the cache line and discard "dirty" data. If the assertion is violated +(i.e., the cache line is read again) then the behavior is undefined. +|`cache_control_read` +| +This property asserts that the cache line containing the data +loaded from memory through the `annotated_ptr` will not be written +until kernel execution is completed. +If the assertion is violated (the cache line is written), the behavior +is undefined. +|`cache_control_write` +| +This property requests that writes to memory through the `annotated_ptr` +should not cache the data at level L in the memory hierarchy. +|`cache_control_write` +| +This property requests that writes to memory through the `annotated_ptr` +should immediately write the data to the next-level cache after L +and mark the cache line at level L as "not dirty". +|`cache_control_write` +| +This property requests that writes to memory through the `annotated_ptr` +should write the data into the cache at level L and mark the cache line as +"dirty". Upon eviction, "dirty" data will be written into the cache at +level higher than L. +|`cache_control_write` +| +This property is the same as `WriteThrough`, but requests use of a +policy that gives lower priority to data in the cache present +via a `Streaming` cache control. +|==== +-- + +== Implementation notes + +It is intended that the SYCL cache control properties will be used by the compiler +to generate SPIR-V cache control operations. + From e80276ed2d6d0964d14d7c53e87fd2c1ee32ffdd Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Thu, 7 Sep 2023 16:14:06 -0700 Subject: [PATCH 02/28] Updated doc based on review comments. --- .../sycl_ext_intel_cache_controls.asciidoc | 147 ++++++++++++------ 1 file changed, 102 insertions(+), 45 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc index d068035b1477d..e71c7e1ab6abc 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -45,9 +45,6 @@ This extension depends on the following SYCL extensions: * link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties] * link:../experimental/sycl_ext_oneapi_annotated_ptr.asciidoc[sycl_ext_oneapi_annotated_ptr] -This extension also depends on the following SPIR-V extension: - -* link:../supported/sycl_ext_oneapi_myotherextension.asciidoc[SPV_INTEL_cache_controls] == Status @@ -59,14 +56,10 @@ incompatible ways before it is finalized. *Shipping software products should not rely on APIs defined in this specification.* -== Backend support status - -This extension is not implemented. - == Overview This extension introduces additional compile-time properties for -the proposed `sycl::ext::oneapi::experimental::annotated_ptr` class to specify +the `sycl::ext::oneapi::experimental::annotated_ptr` class to specify cache control information. The cache controls are a strong request that memory accesses through the @@ -105,88 +98,128 @@ Below is a list of new compile-time constant properties supported with ```c++ namespace sycl::ext::intel::experimental { -enum cache_control_types_read { - Cached, Uncached, Streaming, InvalidateAfterRead, ConstCached}; -enum cache_control_types_write { - Uncached, Streaming, WriteThrough, WriteBack}; +enum cache_control_read_type { + read_cached, read_uncached, read_streaming, invalidate_after_read, read_const_cached}; + +enum cache_control_write_typeType { + write_const_uncached, write_streaming, write_through, write_back}; struct cache_control_read { - template - using value_t = property_value>; + template + using value_t = property_value>; }; -struct cache_control_write { + +struct cache_control_write_type { template - using value_t = property_value>; + using value_t = property_value>; }; -template +template inline constexpr cache_control_read::value_t read_control; -template -inline constexpr cache_control_write::value_t write_control; + +template +inline constexpr cache_control_write_type::value_t write_control; template<> struct is_property : std::true_type {}; + template<> -struct is_property : std::true_type {}; +struct is_property : std::true_type {}; template -struct is_property_of< +struct is_property_key_of< cache_control_read, annotated_ptr> : std::true_type {}; -template -struct is_property_of< - cache_control_write, annotated_ptr> : std::true_type {}; template +struct is_property_key_of< + cache_control_write_type, annotated_ptr> : std::true_type {}; + +template +inline constexpr cache_control_read::value_t cache_control_read_cached; + +template +inline constexpr cache_control_read::value_t cache_control_read_uncached; + +template +inline constexpr cache_control_read::value_t cache_control_read_streaming; + +template +inline constexpr cache_control_read::value_t cache_control_invalidate_after_read; + +template +inline constexpr cache_control_read::value_t cache_control_read_const_cached; + +template +inline constexpr cache_control_write::value_t cache_control_write_const_uncached; + +template +inline constexpr cache_control_write::value_t cache_control_write_streaming; + +template +inline constexpr cache_control_write::value_t cache_control_write_through; + +template +inline constexpr cache_control_write::value_t cache_control_write_back; + } // namespace sycl::ext::intel::experimental ``` +Each of these properties takes a cache level parameter indicating which level +of the cache hierarchy is affected. Cache level 0 indicates the cache closest +to the processing unit, cache level 1 indicates the next furthest cache +level, etc. It is legal to specify a cache level that does not exist on +the target device, but the property will be ignored in this case. + +Note that a property specifies the cache behavior only for the indicated +cache level. In order to specify the behavior for multiple cache levels, +multiple properties should be specified. + +It is legal to specify several different cache control properties in the +same `annotated_ptr`. However, all instances of cache_control_read_type must +have different cache levels and all instances of cache_control_write_type +must have difference cache levels. + +The cache control properties are divided into two categories: those that +are hints and those that are assertions by the application. + +==== Cache control hints +These properties are hints requesting specific cache behavior when +loading or storing to memory through the annotated_ptr. These properties can +affect the performance of device code, but they do not change the semantics. + -- [options="header"] |==== | Property | Description -|`cache_control_read` +|`cache_control_read` | This property requests that loads from memory through the `annotated_ptr` may cache the data at level L in the memory hierarchy. -|`cache_control_read` +|`cache_control_read` | This property requests that loads from memory through the `annotated_ptr` should not cache the data at level L in the memory hierarchy. -|`cache_control_read` +|`cache_control_read` | This property requests that loads from memory through the `annotated_ptr` should cache the data at cache level L. The eviction policy is to give lower priority to data cached using this property versus the Cached property. -|`cache_control_read` -| -This property asserts that the cache line into which data is loaded -from memory through the `annotated_ptr` will not be read again -until it is overwritten. Therefore the load operation can invalidate -the cache line and discard "dirty" data. If the assertion is violated -(i.e., the cache line is read again) then the behavior is undefined. -|`cache_control_read` -| -This property asserts that the cache line containing the data -loaded from memory through the `annotated_ptr` will not be written -until kernel execution is completed. -If the assertion is violated (the cache line is written), the behavior -is undefined. -|`cache_control_write` +|`cache_control_write` | This property requests that writes to memory through the `annotated_ptr` should not cache the data at level L in the memory hierarchy. -|`cache_control_write` +|`cache_control_write` | This property requests that writes to memory through the `annotated_ptr` should immediately write the data to the next-level cache after L and mark the cache line at level L as "not dirty". -|`cache_control_write` +|`cache_control_write` | This property requests that writes to memory through the `annotated_ptr` should write the data into the cache at level L and mark the cache line as "dirty". Upon eviction, "dirty" data will be written into the cache at level higher than L. -|`cache_control_write` +|`cache_control_write` | This property is the same as `WriteThrough`, but requests use of a policy that gives lower priority to data in the cache present @@ -194,6 +227,30 @@ via a `Streaming` cache control. |==== -- +==== Assertions by the application +These properties are assertions by the application, promising that the application accesses memory in a certain way. Care must be taken when using these properties because they can lead to undefined behavior if they are misused. + +-- +[options="header"] +|==== +| Property | Description +|`cache_control_read` +| +This property asserts that the cache line into which data is loaded +from memory through the `annotated_ptr` will not be read again +until it is overwritten. Therefore the load operation can invalidate +the cache line and discard "dirty" data. If the assertion is violated +(i.e., the cache line is read again) then the behavior is undefined. +|`cache_control_read` +| +This property asserts that the cache line containing the data +loaded from memory through the `annotated_ptr` will not be written +until kernel execution is completed. +If the assertion is violated (the cache line is written), the behavior +is undefined. +|==== +-- + == Implementation notes It is intended that the SYCL cache control properties will be used by the compiler From 81d91c7fd65ae3b3bbcf521bfdcc0f9ffd8358ff Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Fri, 8 Sep 2023 13:06:39 -0700 Subject: [PATCH 03/28] Cleanup of doc. --- .../sycl_ext_intel_cache_controls.asciidoc | 131 ++++++++++++------ 1 file changed, 90 insertions(+), 41 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc index e71c7e1ab6abc..4d54d22fc9789 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -98,68 +98,78 @@ Below is a list of new compile-time constant properties supported with ```c++ namespace sycl::ext::intel::experimental { -enum cache_control_read_type { - read_cached, read_uncached, read_streaming, invalidate_after_read, read_const_cached}; +enum class cache_control_read_type : /* unspecified */ { + cached, + uncached, + streaming, + invalidate_after_read, + const_cached +}; -enum cache_control_write_typeType { - write_const_uncached, write_streaming, write_through, write_back}; +enum class cache_control_write_type : /* unspecified */ { + uncached, + streaming, + write_through, + write_back +}; -struct cache_control_read { - template - using value_t = property_value>; +struct cache_control_read_key { + template + using value_t = property_value>; }; -struct cache_control_write_type { - template - using value_t = property_value>; +struct cache_control_write_key { + template + using value_t = property_value>; }; template -inline constexpr cache_control_read::value_t read_control; +inline constexpr cache_control_read_key::value_t cache_control_read; template -inline constexpr cache_control_write_type::value_t write_control; +inline constexpr cache_control_write_key::value_t cache_control_write; template<> -struct is_property : std::true_type {}; +struct is_property_key : std::true_type {}; template<> -struct is_property : std::true_type {}; +struct is_property_key : std::true_type {}; template struct is_property_key_of< - cache_control_read, annotated_ptr> : std::true_type {}; + cache_control_read_key, annotated_ptr> : std::true_type {}; template struct is_property_key_of< - cache_control_write_type, annotated_ptr> : std::true_type {}; + cache_control_write_key, annotated_ptr> : std::true_type {}; template -inline constexpr cache_control_read::value_t cache_control_read_cached; +inline constexpr cache_control_read_key::value_t +cache_control_read_cached; template -inline constexpr cache_control_read::value_t cache_control_read_uncached; +inline constexpr cache_control_read_key::value_t cache_control_read_uncached; template -inline constexpr cache_control_read::value_t cache_control_read_streaming; +inline constexpr cache_control_read_key::value_t cache_control_read_streaming; template -inline constexpr cache_control_read::value_t cache_control_invalidate_after_read; +inline constexpr cache_control_read_key::value_t cache_control_invalidate_after_read; template -inline constexpr cache_control_read::value_t cache_control_read_const_cached; +inline constexpr cache_control_read_key::value_t cache_control_read_const_cached; template -inline constexpr cache_control_write::value_t cache_control_write_const_uncached; +inline constexpr cache_control_write_key::value_t cache_control_write_uncached; template -inline constexpr cache_control_write::value_t cache_control_write_streaming; +inline constexpr cache_control_write_key::value_t cache_control_write_streaming; template -inline constexpr cache_control_write::value_t cache_control_write_through; +inline constexpr cache_control_write_key::value_t cache_control_write_through; template -inline constexpr cache_control_write::value_t cache_control_write_back; +inline constexpr cache_control_write_key::value_t cache_control_write_back; } // namespace sycl::ext::intel::experimental ``` @@ -187,61 +197,100 @@ loading or storing to memory through the annotated_ptr. These properties can affect the performance of device code, but they do not change the semantics. -- -[options="header"] +[options="header", cols="2,1"] |==== | Property | Description -|`cache_control_read` +a| +[source] +---- +cache_control_read +---- | This property requests that loads from memory through the `annotated_ptr` may cache the data at level L in the memory hierarchy. -|`cache_control_read` +a| +[source] +---- +cache_control_read +---- | This property requests that loads from memory through the `annotated_ptr` should not cache the data at level L in the memory hierarchy. -|`cache_control_read` +a| +[source] +---- +cache_control_read +---- | This property requests that loads from memory through the `annotated_ptr` should cache the data at cache level L. The eviction policy is to give lower priority to data cached using this property versus the Cached property. -|`cache_control_write` +a| +[source] +---- +cache_control_write +---- | This property requests that writes to memory through the `annotated_ptr` should not cache the data at level L in the memory hierarchy. -|`cache_control_write` +a| +[source] +---- +cache_control_write +---- | This property requests that writes to memory through the `annotated_ptr` should immediately write the data to the next-level cache after L and mark the cache line at level L as "not dirty". -|`cache_control_write` +a| +[source] +---- +cache_control_write +---- | This property requests that writes to memory through the `annotated_ptr` should write the data into the cache at level L and mark the cache line as "dirty". Upon eviction, "dirty" data will be written into the cache at level higher than L. -|`cache_control_write` +a| +[source] +---- +cache_control_write +---- | -This property is the same as `WriteThrough`, but requests use of a +This property is the same as `write_through`, but requests use of a policy that gives lower priority to data in the cache present -via a `Streaming` cache control. +via a `streaming` cache control. |==== -- ==== Assertions by the application -These properties are assertions by the application, promising that the application accesses memory in a certain way. Care must be taken when using these properties because they can lead to undefined behavior if they are misused. +These properties are assertions by the application, promising that the +application accesses memory in a certain way. Care must be taken when +using these properties because they can lead to undefined behavior if +they are misused. -- -[options="header"] +[options="header", cols="3,1"] |==== | Property | Description -|`cache_control_read` +a| +[source] +---- +cache_control_read +---- | This property asserts that the cache line into which data is loaded from memory through the `annotated_ptr` will not be read again until it is overwritten. Therefore the load operation can invalidate the cache line and discard "dirty" data. If the assertion is violated (i.e., the cache line is read again) then the behavior is undefined. -|`cache_control_read` +a| +[source] +---- +cache_control_read +---- | This property asserts that the cache line containing the data loaded from memory through the `annotated_ptr` will not be written @@ -253,6 +302,6 @@ is undefined. == Implementation notes -It is intended that the SYCL cache control properties will be used by the compiler -to generate SPIR-V cache control operations. +It is intended that the SYCL cache control properties will be used by the +compiler to generate SPIR-V cache control operations. From 8725dfc46960ce9c726837a36aff2825ba768b1c Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Mon, 11 Sep 2023 10:50:39 -0700 Subject: [PATCH 04/28] Formatting changes. --- .../sycl_ext_intel_cache_controls.asciidoc | 30 ++++++++++--------- 1 file changed, 16 insertions(+), 14 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc index 4d54d22fc9789..7f6766f930d63 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -126,7 +126,7 @@ struct cache_control_write_key { template inline constexpr cache_control_read_key::value_t cache_control_read; -template +template inline constexpr cache_control_write_key::value_t cache_control_write; template<> @@ -184,8 +184,8 @@ cache level. In order to specify the behavior for multiple cache levels, multiple properties should be specified. It is legal to specify several different cache control properties in the -same `annotated_ptr`. However, all instances of cache_control_read_type must -have different cache levels and all instances of cache_control_write_type +same `annotated_ptr`. However, all instances of `cache_control_read_type` must +have different cache levels and all instances of `cache_control_write_type` must have difference cache levels. The cache control properties are divided into two categories: those that @@ -193,7 +193,7 @@ are hints and those that are assertions by the application. ==== Cache control hints These properties are hints requesting specific cache behavior when -loading or storing to memory through the annotated_ptr. These properties can +loading or storing to memory through the `annotated_ptr`. These properties can affect the performance of device code, but they do not change the semantics. -- @@ -207,7 +207,7 @@ cache_control_read ---- | This property requests that loads from memory through the `annotated_ptr` -may cache the data at level L in the memory hierarchy. +may cache the data at level `L` in the memory hierarchy. a| [source] ---- @@ -215,7 +215,7 @@ cache_control_read ---- | This property requests that loads from memory through the `annotated_ptr` -should not cache the data at level L in the memory hierarchy. +should not cache the data at level `L` in the memory hierarchy. a| [source] ---- @@ -223,8 +223,8 @@ cache_control_read ---- | This property requests that loads from memory through the `annotated_ptr` -should cache the data at cache level L. The eviction policy is to give -lower priority to data cached using this property versus the Cached +should cache the data at cache level `L`. The eviction policy is to give +lower priority to data cached using this property versus the `cached` property. a| [source] @@ -233,7 +233,7 @@ cache_control_write ---- | This property requests that writes to memory through the `annotated_ptr` -should not cache the data at level L in the memory hierarchy. +should not cache the data at level `L` in the memory hierarchy. a| [source] ---- @@ -241,8 +241,8 @@ cache_control_write ---- | This property requests that writes to memory through the `annotated_ptr` -should immediately write the data to the next-level cache after L -and mark the cache line at level L as "not dirty". +should immediately write the data to the next-level cache after `L` +and mark the cache line at level `L` as "not dirty". a| [source] ---- @@ -250,9 +250,9 @@ cache_control_write ---- | This property requests that writes to memory through the `annotated_ptr` -should write the data into the cache at level L and mark the cache line as +should write the data into the cache at level `L` and mark the cache line as "dirty". Upon eviction, "dirty" data will be written into the cache at -level higher than L. +level higher than `L`. a| [source] ---- @@ -303,5 +303,7 @@ is undefined. == Implementation notes It is intended that the SYCL cache control properties will be used by the -compiler to generate SPIR-V cache control operations. +compiler to generate SPIR-V cache control operations. Alternatively, the +properties could be implemented by generating intrinsic function calls +that match the cache control types. From b994128bd29912f8fe8d14e96068ae5314c71ca3 Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Tue, 17 Oct 2023 19:14:51 -0700 Subject: [PATCH 05/28] [SYCL] Support for cache control properties on annotated_ptr. --- .../lib/SYCLLowerIR/CompileTimeProperties.def | 11 + .../SYCLLowerIR/CompileTimePropertiesPass.cpp | 203 ++++++++++++++---- .../experimental/cache_control_properties.hpp | 165 ++++++++++++++ .../annotated_arg/annotated_arg.hpp | 1 + .../annotated_ptr/annotated_ptr.hpp | 22 +- .../sycl/ext/oneapi/properties/property.hpp | 11 +- 6 files changed, 372 insertions(+), 41 deletions(-) create mode 100755 sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp diff --git a/llvm/lib/SYCLLowerIR/CompileTimeProperties.def b/llvm/lib/SYCLLowerIR/CompileTimeProperties.def index 926adbdd1128d..f0012f7bc38a0 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimeProperties.def +++ b/llvm/lib/SYCLLowerIR/CompileTimeProperties.def @@ -33,3 +33,14 @@ SYCL_COMPILE_TIME_PROPERTY("sycl-stable", 6183, DecorValueTy::boolean) SYCL_COMPILE_TIME_PROPERTY("sycl-strict", 19, DecorValueTy::boolean) SYCL_COMPILE_TIME_PROPERTY("sycl-latency-anchor-id", 6172, DecorValueTy::string) SYCL_COMPILE_TIME_PROPERTY("sycl-latency-constraint", 6173, DecorValueTy::string) + +// The corresponding SPIR-V OpCodes for cache control properties +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-cached", 6442, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-uncached", 6442, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-streaming", 6442, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-invalidate-after-read", 6442, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-const-cached", 6442, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-write-uncached", 6443, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-write-streaming", 6443, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-write-through", 6443, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-write-back", 6443, DecorValueTy::uint32) diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index 32bcd95bb27ba..f173fa6590e6d 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -22,6 +22,9 @@ #include "llvm/IR/Operator.h" #include "llvm/TargetParser/Triple.h" +#include +#include "llvm/Support/raw_ostream.h" + using namespace llvm; namespace { @@ -42,6 +45,9 @@ constexpr uint32_t SPIRV_HOST_ACCESS_DEFAULT_VALUE = 2; // Read/Write constexpr uint32_t SPIRV_INITIATION_INTERVAL_DECOR = 5917; constexpr uint32_t SPIRV_PIPELINE_ENABLE_DECOR = 5919; +constexpr uint32_t SPIRV_CACHE_CONTROL_READ_DECOR = 6442; +constexpr uint32_t SPIRV_CACHE_CONTROL_WRITE_DECOR = 6443; + enum class DecorValueTy { uint32, boolean, @@ -82,6 +88,62 @@ MDNode *buildSpirvDecorMetadata(LLVMContext &Ctx, uint32_t OpCode, return MDNode::get(Ctx, MD); } +/// Builds a metadata node for a SPIR-V decoration for cache controls +/// where decoration code and value are both uint32_t integers. +/// The value encodes a cache level and a cache control type. +/// +/// @param Ctx [in] the LLVM Context. +/// @param OpCode [in] the SPIR-V OpCode code. +/// @param Value [in] the SPIR-V decoration value. +/// +/// @returns a pointer to the metadata node created for the required decoration +/// and its values. +MDNode *buildSpirvDecorCacheProp(LLVMContext &Ctx, StringRef Name, + uint32_t OpCode, uint32_t CacheLevel) { + enum class LoadCachePropINTEL { + Uncached = 0, + Cached = 1, + Streaming = 2, + InvalidateAfterRead = 3, + ConstCached = 4 + }; + enum class StoreCachePropINTEL { + Uncached = 0, + WriteThrough = 1, + WriteBack = 2, + Streaming = 3 + }; + uint32_t CacheProp; + if (Name == "sycl-cache-read-uncached") + CacheProp = static_cast(LoadCachePropINTEL::Uncached); + else if (Name == "sycl-cache-read-cached") + CacheProp = static_cast(LoadCachePropINTEL::Cached); + else if (Name == "sycl-cache-read-streaming") + CacheProp = static_cast(LoadCachePropINTEL::Streaming); + else if (Name == "sycl-cache-read-invalidate-after-read") + CacheProp = static_cast(LoadCachePropINTEL::InvalidateAfterRead); + else if (Name == "sycl-cache-read-const-cached") + CacheProp = static_cast(LoadCachePropINTEL::ConstCached); + else if (Name == "sycl-cache-write-uncached") + CacheProp = static_cast(StoreCachePropINTEL::Uncached); + else if (Name == "sycl-cache-write-through") + CacheProp = static_cast(StoreCachePropINTEL::WriteThrough); + else if (Name == "sycl-cache-write-back") + CacheProp = static_cast(StoreCachePropINTEL::WriteBack); + else if (Name == "sycl-cache-write-streaming") + CacheProp = static_cast(StoreCachePropINTEL::Streaming); + + auto *Ty = Type::getInt32Ty(Ctx); + SmallVector MD; + MD.push_back(ConstantAsMetadata::get( + Constant::getIntegerValue(Ty, APInt(32, OpCode)))); + MD.push_back(ConstantAsMetadata::get( + Constant::getIntegerValue(Ty, APInt(32, CacheLevel)))); + MD.push_back(ConstantAsMetadata::get( + Constant::getIntegerValue(Ty, APInt(32, CacheProp)))); + return MDNode::get(Ctx, MD); +} + /// Builds a metadata node for a SPIR-V decoration (decoration code /// is \c uint32_t integer and value is a string). /// @@ -610,9 +672,14 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation( // check alignment annotation and apply it to load/store parseAlignmentAndApply(M, IntrInst); - // Read the annotation values and create the new annotation string. + // Read the annotation values and create new annotation strings. std::string NewAnnotString = ""; auto Properties = parseSYCLPropertiesString(M, IntrInst); + SmallVector MDOpsCacheProp; + uint32_t CacheLevelsSpecifiedLoad = 0; + uint32_t CacheLevelsSpecifiedStore = 0; + bool CacheProp = false; + bool FPGAProp = false; for (auto &Property : Properties) { // sycl-alignment is converted to align on // previous parseAlignmentAndApply(), dropping here @@ -624,51 +691,113 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation( continue; uint32_t DecorCode = DecorIt->second.Code; - // Expected format is '{X}' or '{X:Y}' where X is decoration ID and - // Y is the value if present. It encloses Y in " to ensure that - // string values are handled correctly. Note that " around values are - // always valid, even if the decoration parameters are not strings. - NewAnnotString += "{" + std::to_string(DecorCode); - if (Property.second) - NewAnnotString += ":\"" + Property.second->str() + "\""; - NewAnnotString += "}"; + // Handle cache control properties + if ((*Property.first).starts_with("sycl-cache-")) { + CacheProp = true; + auto DecorValue = Property.second; + uint32_t AttrVal; + DecorValue->getAsInteger(0, AttrVal); + // Check that a particular cache level is specified only once for + // load/store. + if ((*Property.first).starts_with("sycl-cache-read-")) { + assert( + (AttrVal & CacheLevelsSpecifiedLoad) == 0 && + "Conflicting Read cache control specified in pointer annotation"); + CacheLevelsSpecifiedLoad |= AttrVal; + } else { + assert( + (AttrVal & CacheLevelsSpecifiedStore) == 0 && + "Conflicting Write cache control specified in pointer annotation"); + CacheLevelsSpecifiedStore |= AttrVal; + } + + // Format is: + // !Annot = !{!CC1, !CC2} + // !CC1 = !{i32 Load/Store, i32 Level, i32 Control} + // !CC2 = !{i32 Load/Store, i32 Level, i32 Control} + LLVMContext &Ctx = M.getContext(); + uint32_t CacheLevel = 0; + while (AttrVal) { + // The attribute value encodes cache levels as L1->bit0, L2->bit1, + // L3->bit2 and L4->bit3. The SPIR-V encoding uses numbers 0..3. + if (AttrVal & 1) + MDOpsCacheProp.push_back(buildSpirvDecorCacheProp( + Ctx, *Property.first, DecorCode, CacheLevel)); + ++CacheLevel; + AttrVal >>= 1; + } + } else { + FPGAProp = true; + // Expected format is '{X}' or '{X:Y}' where X is decoration ID and + // Y is the value if present. It encloses Y in " to ensure that + // string values are handled correctly. Note that " around values are + // always valid, even if the decoration parameters are not strings. + NewAnnotString += "{" + std::to_string(DecorCode); + if (Property.second) + NewAnnotString += ":\"" + Property.second->str() + "\""; + NewAnnotString += "}"; + } } - // If the new annotation string is empty there is no reason to keep it, so - // replace it with the first operand and mark it for removal. - if (NewAnnotString.empty()) { + // If there are no other annotations (except "alignment") then there is no + // reason to keep the original intrinsic, so replace it with the first operand + // and mark it for removal. + if (!CacheProp && !FPGAProp) { IntrInst->replaceAllUsesWith(IntrInst->getOperand(0)); RemovableAnnotations.push_back(IntrInst); return true; } - // Either reuse a previously generated one or create a new global variable - // with the new annotation string. - GlobalVariable *NewAnnotStringGV = nullptr; - auto ExistingNewAnnotStringIt = ReusableAnnotStrings.find(NewAnnotString); - if (ExistingNewAnnotStringIt != ReusableAnnotStrings.end()) { - NewAnnotStringGV = ExistingNewAnnotStringIt->second; - } else { - Constant *NewAnnotStringData = - ConstantDataArray::getString(M.getContext(), NewAnnotString); - NewAnnotStringGV = new GlobalVariable( - M, NewAnnotStringData->getType(), true, GlobalValue::PrivateLinkage, - NewAnnotStringData, ".str", nullptr, llvm::GlobalValue::NotThreadLocal, - IntrAnnotStringArg->getType()->getPointerAddressSpace()); - NewAnnotStringGV->setSection(AnnotStrArgGV->getSection()); - NewAnnotStringGV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global); - ReusableAnnotStrings.insert({NewAnnotString, NewAnnotStringGV}); + if (FPGAProp) { + // Either reuse a previously generated one or create a new global variable + // with the new annotation string. + GlobalVariable *NewAnnotStringGV = nullptr; + auto ExistingNewAnnotStringIt = ReusableAnnotStrings.find(NewAnnotString); + if (ExistingNewAnnotStringIt != ReusableAnnotStrings.end()) { + NewAnnotStringGV = ExistingNewAnnotStringIt->second; + } else { + Constant *NewAnnotStringData = + ConstantDataArray::getString(M.getContext(), NewAnnotString); + NewAnnotStringGV = new GlobalVariable( + M, NewAnnotStringData->getType(), true, GlobalValue::PrivateLinkage, + NewAnnotStringData, ".str", nullptr, + llvm::GlobalValue::NotThreadLocal, + IntrAnnotStringArg->getType()->getPointerAddressSpace()); + NewAnnotStringGV->setSection(AnnotStrArgGV->getSection()); + NewAnnotStringGV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global); + ReusableAnnotStrings.insert({NewAnnotString, NewAnnotStringGV}); + } + + // Replace the annotation string with a bitcast of the new global variable. + IntrInst->setArgOperand( + 1, ConstantExpr::getBitCast(NewAnnotStringGV, + IntrAnnotStringArg->getType())); + + // The values are now in the annotation string, so we can remove the + // original annotation value. + PointerType *Arg4PtrTy = + cast(IntrInst->getArgOperand(4)->getType()); + IntrInst->setArgOperand(4, ConstantPointerNull::get(Arg4PtrTy)); } - // Replace the annotation string with a bitcast of the new global variable. - IntrInst->setArgOperand( - 1, ConstantExpr::getBitCast(NewAnnotStringGV, - IntrAnnotStringArg->getType())); + if (CacheProp) { + LLVMContext &Ctx = M.getContext(); + unsigned MDKindID = Ctx.getMDKindID(SPIRV_DECOR_MD_KIND); + if (!FPGAProp) { + // If there are no annotations other than cache controls we can apply the + // controls to the pointer and remove the intrinsic. + auto PtrInstr = cast(IntrInst->getArgOperand(0)); + PtrInstr->setMetadata(MDKindID, MDTuple::get(Ctx, MDOpsCacheProp)); + // Replace all uses of IntrInst with first operand + IntrInst->replaceAllUsesWith(PtrInstr); + // Delete the original IntrInst + RemovableAnnotations.push_back(IntrInst); + } else { + // If there were FPGA annotations then we retain the original intrinsic + // and apply the cache control properties to its result. + IntrInst->setMetadata(MDKindID, MDTuple::get(Ctx, MDOpsCacheProp)); + } + } - // The values are not in the annotation string, so we can remove the original - // annotation value. - PointerType *Arg4PtrTy = - cast(IntrInst->getArgOperand(4)->getType()); - IntrInst->setArgOperand(4, ConstantPointerNull::get(Arg4PtrTy)); return true; } diff --git a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp new file mode 100755 index 0000000000000..561db9ac6d02b --- /dev/null +++ b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp @@ -0,0 +1,165 @@ +//==--------- SYCL annotated_arg/ptr properties for caching control --------==// +// +// 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 +#include +#include +#include + +namespace sycl { +inline namespace _V1 { +namespace ext { +namespace intel { +namespace experimental { + +enum class level : std::uint16_t { L1=1, L2, L3, L4 }; + +template +using property_value = + sycl::ext::oneapi::experimental::property_value; + +#define __SYCL_CACHE_CONTROL_M1(P) \ + struct P { \ + template \ + using value_t = property_value...>; \ + }; + +__SYCL_CACHE_CONTROL_M1(cache_control_read_cached_key) +__SYCL_CACHE_CONTROL_M1(cache_control_read_uncached_key) +__SYCL_CACHE_CONTROL_M1(cache_control_read_streaming_key) +__SYCL_CACHE_CONTROL_M1(cache_control_invalidate_after_read_key) +__SYCL_CACHE_CONTROL_M1(cache_control_read_const_cached_key) +__SYCL_CACHE_CONTROL_M1(cache_control_write_uncached_key) +__SYCL_CACHE_CONTROL_M1(cache_control_write_streaming_key) +__SYCL_CACHE_CONTROL_M1(cache_control_write_through_key) +__SYCL_CACHE_CONTROL_M1(cache_control_write_back_key) + +} // namespace experimental +} // namespace intel + +namespace oneapi { +namespace experimental { + +template class annotated_arg; +template class annotated_ptr; + +#define __SYCL_CACHE_CONTROL_M2(P) \ + using P = intel::experimental::P; \ + template <> struct is_property_key

: std::true_type {}; \ + template \ + struct is_property_key_of> \ + : std::true_type {}; \ + template \ + struct is_property_key_of> \ + : std::true_type {}; + +__SYCL_CACHE_CONTROL_M2(cache_control_read_cached_key) +__SYCL_CACHE_CONTROL_M2(cache_control_read_uncached_key) +__SYCL_CACHE_CONTROL_M2(cache_control_read_streaming_key) +__SYCL_CACHE_CONTROL_M2(cache_control_invalidate_after_read_key) +__SYCL_CACHE_CONTROL_M2(cache_control_read_const_cached_key) +__SYCL_CACHE_CONTROL_M2(cache_control_write_uncached_key) +__SYCL_CACHE_CONTROL_M2(cache_control_write_streaming_key) +__SYCL_CACHE_CONTROL_M2(cache_control_write_through_key) +__SYCL_CACHE_CONTROL_M2(cache_control_write_back_key) + +using namespace intel::experimental; + +#define __SYCL_CACHE_CONTROL_M3(P) \ + template inline constexpr P##_key::value_t P; + +__SYCL_CACHE_CONTROL_M3(cache_control_read_cached) +__SYCL_CACHE_CONTROL_M3(cache_control_read_uncached) +__SYCL_CACHE_CONTROL_M3(cache_control_read_streaming) +__SYCL_CACHE_CONTROL_M3(cache_control_invalidate_after_read) +__SYCL_CACHE_CONTROL_M3(cache_control_read_const_cached) +__SYCL_CACHE_CONTROL_M3(cache_control_write_uncached) +__SYCL_CACHE_CONTROL_M3(cache_control_write_streaming) +__SYCL_CACHE_CONTROL_M3(cache_control_write_through) +__SYCL_CACHE_CONTROL_M3(cache_control_write_back) + +namespace detail { + +template static constexpr void checkLevels() {} +template static constexpr void checkLevels() { + static_assert(L_1 != L_2, "Duplicate cache level specification."); +} +template static constexpr void checkLevels() { + static_assert(L_1 != L_2 && L_1 != L_3 && L_2 != L_3, + "Duplicate cache level specification."); +} +template +static constexpr void checkLevels() { + static_assert(L_1 != L_2 && L_1 != L_3 && L_1 != L_4 && L_2 != L_3 && + L_2 != L_4 && L_3 != L_4, + "Duplicate cache level specification."); +} + +#define __SYCL_CACHE_CONTROL_M4(P, K, N) \ + template <> struct PropertyToKind

{ \ + static constexpr PropKind Kind = PropKind::K; \ + }; \ + template <> struct IsCompileTimeProperty

: std::true_type {}; \ + template struct PropertyMetaInfo> { \ + static constexpr const char *name = N; \ + static constexpr const int value = \ + (checkLevels(), ((1 << (static_cast(Ls) - 1)) | ...)); \ + }; + +__SYCL_CACHE_CONTROL_M4(cache_control_read_cached_key, CacheControlReadCached, + "sycl-cache-read-cached") +__SYCL_CACHE_CONTROL_M4(cache_control_read_uncached_key, + CacheControlReadUncached, "sycl-cache-read-uncached") +__SYCL_CACHE_CONTROL_M4(cache_control_read_streaming_key, + CacheControlReadStreaming, "sycl-cache-read-streaming") +__SYCL_CACHE_CONTROL_M4(cache_control_invalidate_after_read_key, + CacheControlReadInvalidateAfterRead, + "sycl-cache-read-invalidate-after-read") +__SYCL_CACHE_CONTROL_M4(cache_control_read_const_cached_key, + CacheControlReadConstCached, + "sycl-cache-read-const-cached") +__SYCL_CACHE_CONTROL_M4(cache_control_write_uncached_key, + CacheControlWriteUncached, "sycl-cache-write-uncached") +__SYCL_CACHE_CONTROL_M4(cache_control_write_streaming_key, + CacheControlWriteStreaming, + "sycl-cache-write-streaming") +__SYCL_CACHE_CONTROL_M4(cache_control_write_through_key, + CacheControlWriteThrough, "sycl-cache-write-through") +__SYCL_CACHE_CONTROL_M4(cache_control_write_back_key, CacheControlWriteBack, + "sycl-cache-write-back") + +} // namespace detail + +#define __SYCL_CACHE_CONTROL_M5(P) \ + template \ + struct is_valid_property> \ + : std::bool_constant::value> {}; + +__SYCL_CACHE_CONTROL_M5(cache_control_read_cached_key) +__SYCL_CACHE_CONTROL_M5(cache_control_read_uncached_key) +__SYCL_CACHE_CONTROL_M5(cache_control_read_streaming_key) +__SYCL_CACHE_CONTROL_M5(cache_control_invalidate_after_read_key) +__SYCL_CACHE_CONTROL_M5(cache_control_read_const_cached_key) +__SYCL_CACHE_CONTROL_M5(cache_control_write_uncached_key) +__SYCL_CACHE_CONTROL_M5(cache_control_write_streaming_key) +__SYCL_CACHE_CONTROL_M5(cache_control_write_through_key) +__SYCL_CACHE_CONTROL_M5(cache_control_write_back_key) + +#undef __SYCL_CACHE_CONTROL_M1 +#undef __SYCL_CACHE_CONTROL_M2 +#undef __SYCL_CACHE_CONTROL_M3 +#undef __SYCL_CACHE_CONTROL_M4 +#undef __SYCL_CACHE_CONTROL_M5 + +} // namespace experimental +} // namespace oneapi +} // namespace ext +} // namespace _V1 +} // namespace sycl diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp index 90c6c5cfbfeb2..6cf38ae0c198e 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include #include diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp index 0e4ecbdf17c54..8ef76607a9f66 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp @@ -191,9 +191,25 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr> { using property_list_t = detail::properties_t; // buffer_location and alignment are allowed for annotated_ref - using allowed_properties = - std::tuple), - decltype(ext::oneapi::experimental::alignment<0>)>; + using allowed_properties = std::tuple< + decltype(ext::intel::experimental::buffer_location<0>), + decltype(ext::oneapi::experimental::alignment<0>), + decltype(ext::oneapi::experimental::cache_control_read_cached), + decltype(ext::oneapi::experimental::cache_control_read_uncached< + level::L1>), + decltype(ext::oneapi::experimental::cache_control_read_streaming< + level::L1>), + decltype(ext::oneapi::experimental::cache_control_invalidate_after_read< + level::L1>), + decltype(ext::oneapi::experimental::cache_control_read_const_cached< + level::L1>), + decltype(ext::oneapi::experimental::cache_control_write_uncached< + level::L1>), + decltype(ext::oneapi::experimental::cache_control_write_streaming< + level::L1>), + decltype(ext::oneapi::experimental::cache_control_write_through< + level::L1>), + decltype(ext::oneapi::experimental::cache_control_write_back)>; using filtered_properties = typename PropertiesFilter::tuple; diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index e6220f8a79e3e..a7675dd092256 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -201,8 +201,17 @@ enum PropKind : uint32_t { RegisterAllocMode = 31, GRFSize = 32, GRFSizeAutomatic = 33, + CacheControlReadCached = 34, + CacheControlReadUncached = 35, + CacheControlReadStreaming = 36, + CacheControlReadInvalidateAfterRead = 37, + CacheControlReadConstCached = 38, + CacheControlWriteUncached = 39, + CacheControlWriteStreaming = 40, + CacheControlWriteThrough = 41, + CacheControlWriteBack = 42, // PropKindSize must always be the last value. - PropKindSize = 34, + PropKindSize = 42, }; // This trait must be specialized for all properties and must have a unique From e097acb8fce61e4e9a5088fe90f00ad379b557e8 Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Wed, 18 Oct 2023 15:23:36 -0700 Subject: [PATCH 06/28] Formatting change. --- llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp | 3 --- .../intel/experimental/cache_control_properties.hpp | 2 +- .../properties/properties_cache_control.cpp | 11 +++++------ 3 files changed, 6 insertions(+), 10 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index 899009b2fe9d7..080f9c2feec07 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -22,9 +22,6 @@ #include "llvm/IR/Operator.h" #include "llvm/TargetParser/Triple.h" -#include -#include "llvm/Support/raw_ostream.h" - using namespace llvm; namespace { diff --git a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp index 4b78f74eb3856..5010fa7a82aa5 100755 --- a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp @@ -19,7 +19,7 @@ namespace ext { namespace intel { namespace experimental { -enum class level : std::uint16_t { L1=0, L2, L3, L4 }; +enum class level : std::uint16_t { L1 = 0, L2, L3, L4 }; template using property_value = diff --git a/sycl/test/extensions/properties/properties_cache_control.cpp b/sycl/test/extensions/properties/properties_cache_control.cpp index 394a9d68e35d0..af95dcb4a6cc9 100755 --- a/sycl/test/extensions/properties/properties_cache_control.cpp +++ b/sycl/test/extensions/properties/properties_cache_control.cpp @@ -19,7 +19,6 @@ using annotated_ptr_store = cache_control_write_through, cache_control_write_back))>; - void cache_control_read_func() { queue q; constexpr int N = 10; @@ -36,14 +35,14 @@ void cache_control_read_func() { void cache_control_write_func() { queue q; constexpr int N = 10; - float* ArrayA = malloc_shared(N, q); - q.submit([&](handler& cgh) { + float *ArrayA = malloc_shared(N, q); + q.submit([&](handler &cgh) { cgh.parallel_for<>(range<1>(N), [=](item<1> item) { auto item_id = item.get_linear_id(); - annotated_ptr_store dst{ &ArrayA[item_id] }; + annotated_ptr_store dst{&ArrayA[item_id]}; *dst = 55.0f; - }); }); + }); } // CHECK-IR: spir_kernel{{.*}}cache_control_read_func @@ -53,7 +52,7 @@ void cache_control_write_func() { // CHECK-IR: spir_kernel{{.*}}cache_control_write_func // CHECK-IR: {{.*}}getelementptr inbounds float{{.*}}!spirv.Decorations [[WDECOR:.*]] // CHECK-IR: ret void -// + // CHECK-IR: [[RDECOR]] = !{[[RDECOR1:.*]], [[RDECOR2:.*]], [[RDECOR3:.*]], [[RDECOR4:.*]]} // CHECK-IR: [[RDECOR1]] = !{i32 6442, i32 0, i32 1} // CHECK-IR: [[RDECOR2]] = !{i32 6442, i32 1, i32 0} From f847ea308c381d9e7c83073add453bea1a010b60 Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Tue, 24 Oct 2023 15:46:59 -0700 Subject: [PATCH 07/28] Change in syntax. --- .../lib/SYCLLowerIR/CompileTimeProperties.def | 11 +- .../SYCLLowerIR/CompileTimePropertiesPass.cpp | 137 +++++---- .../sycl_ext_intel_cache_controls.asciidoc | 278 ++++-------------- .../experimental/cache_control_properties.hpp | 222 +++++++------- .../annotated_ptr/annotated_ptr.hpp | 20 +- .../sycl/ext/oneapi/properties/property.hpp | 13 +- .../properties/properties_cache_control.cpp | 29 +- 7 files changed, 278 insertions(+), 432 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/CompileTimeProperties.def b/llvm/lib/SYCLLowerIR/CompileTimeProperties.def index 72b2f23d8089f..3b6fc20eea49a 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimeProperties.def +++ b/llvm/lib/SYCLLowerIR/CompileTimeProperties.def @@ -56,12 +56,5 @@ SYCL_COMPILE_TIME_PROPERTY("sycl-bi-directional-ports-true", 5885, DecorValueTy::none) // The corresponding SPIR-V OpCodes for cache control properties -SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-cached", 6442, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-uncached", 6442, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-streaming", 6442, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-invalidate-after-read", 6442, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-const-cached", 6442, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-cache-write-uncached", 6443, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-cache-write-streaming", 6443, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-cache-write-through", 6443, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-cache-write-back", 6443, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-hint", 6442, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-write-hint", 6443, DecorValueTy::uint32) diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index 080f9c2feec07..4ab2488a70712 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -104,47 +104,79 @@ MDNode *buildSpirvDecorMetadata(LLVMContext &Ctx, uint32_t OpCode, /// where decoration code and value are both uint32_t integers. /// The value encodes a cache level and a cache control type. /// -/// @param Ctx [in] the LLVM Context. -/// @param OpCode [in] the SPIR-V OpCode code. -/// @param Value [in] the SPIR-V decoration value. +/// @param Ctx [in] the LLVM Context. +/// @param Name [in] the SPIR-V property string name. +/// @param OpCode [in] the SPIR-V opcode. +/// @param CacheMode [in] whether read or write. +/// @param CacheLevel [in] the cache level. /// /// @returns a pointer to the metadata node created for the required decoration /// and its values. MDNode *buildSpirvDecorCacheProp(LLVMContext &Ctx, StringRef Name, - uint32_t OpCode, uint32_t CacheLevel) { - enum class cache_control_read_type { - uncached = 0, - cached = 1, - streaming = 2, - invalidate_after_read = 3, - const_cached = 4 + uint32_t OpCode, uint32_t CacheMode, + uint32_t CacheLevel) { + // SPIR-V encodings of read control + enum cache_control_read_type { + read_uncached = 0, + read_cached = 1, + read_streaming = 2, + read_invalidate = 3, + read_const_cached = 4 }; - enum class cache_control_write_type { - uncached = 0, + // SPIR-V encodings of write control + enum cache_control_write_type { + write_uncached = 0, write_through = 1, write_back = 2, - streaming = 3 + write_streaming = 3 + }; + // SYCL encodings of read/write control + enum cache_mode { + uncached, + cached, + streaming, + invalidate, + const_cached, + through, + back }; + + // Map SYCL encoding to SPIR-V uint32_t CacheProp; - if (Name == "sycl-cache-read-uncached") - CacheProp = static_cast(cache_control_read_type::uncached); - else if (Name == "sycl-cache-read-cached") - CacheProp = static_cast(cache_control_read_type::cached); - else if (Name == "sycl-cache-read-streaming") - CacheProp = static_cast(cache_control_read_type::streaming); - else if (Name == "sycl-cache-read-invalidate-after-read") - CacheProp = - static_cast(cache_control_read_type::invalidate_after_read); - else if (Name == "sycl-cache-read-const-cached") - CacheProp = static_cast(cache_control_read_type::const_cached); - else if (Name == "sycl-cache-write-uncached") - CacheProp = static_cast(cache_control_write_type::uncached); - else if (Name == "sycl-cache-write-through") - CacheProp = static_cast(cache_control_write_type::write_through); - else if (Name == "sycl-cache-write-back") - CacheProp = static_cast(cache_control_write_type::write_back); - else if (Name == "sycl-cache-write-streaming") - CacheProp = static_cast(cache_control_write_type::streaming); + if (Name == "sycl-cache-read-hint") { + switch (CacheMode) { + case uncached: + CacheProp = read_uncached; + break; + case cached: + CacheProp = read_cached; + break; + case streaming: + CacheProp = read_streaming; + break; + case invalidate: + CacheProp = read_invalidate; + break; + case const_cached: + CacheProp = read_const_cached; + break; + } + } else { + switch (CacheMode) { + case uncached: + CacheProp = write_uncached; + break; + case through: + CacheProp = write_through; + break; + case back: + CacheProp = write_back; + break; + case streaming: + CacheProp = write_streaming; + break; + } + } auto *Ty = Type::getInt32Ty(Ctx); SmallVector MD; @@ -689,8 +721,6 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation( std::string NewAnnotString = ""; auto Properties = parseSYCLPropertiesString(M, IntrInst); SmallVector MDOpsCacheProp; - uint32_t CacheLevelsSpecifiedLoad = 0; - uint32_t CacheLevelsSpecifiedStore = 0; bool CacheProp = false; bool FPGAProp = false; for (const auto &[PropName, PropVal] : Properties) { @@ -710,34 +740,29 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation( auto DecorValue = PropVal; uint32_t AttrVal; DecorValue->getAsInteger(0, AttrVal); - // Check that a particular cache level is specified only once for - // load/store. - if ((*PropName).starts_with("sycl-cache-read-")) { - assert( - (AttrVal & CacheLevelsSpecifiedLoad) == 0 && - "Conflicting Read cache control specified in pointer annotation"); - CacheLevelsSpecifiedLoad |= AttrVal; - } else { - assert( - (AttrVal & CacheLevelsSpecifiedStore) == 0 && - "Conflicting Write cache control specified in pointer annotation"); - CacheLevelsSpecifiedStore |= AttrVal; - } - // Format is: - // !Annot = !{!CC1, !CC2} + // !Annot = !{!CC1, !CC2, ...} // !CC1 = !{i32 Load/Store, i32 Level, i32 Control} // !CC2 = !{i32 Load/Store, i32 Level, i32 Control} + // ... LLVMContext &Ctx = M.getContext(); - uint32_t CacheLevel = 0; + uint32_t CacheMode = 0; while (AttrVal) { - // The attribute value encodes cache levels as L1->bit0, L2->bit1, + // The attribute value encodes cache control and levels. + // Low-order to high-order nibbles represent the enumerated cache modes. + // In each nibble cache levels are encodes as L1->bit0, L2->bit1, // L3->bit2 and L4->bit3. The SPIR-V encoding uses numbers 0..3. - if (AttrVal & 1) - MDOpsCacheProp.push_back( - buildSpirvDecorCacheProp(Ctx, *PropName, DecorCode, CacheLevel)); - ++CacheLevel; - AttrVal >>= 1; + uint32_t CacheLevel = 0; + uint32_t LevelMask = AttrVal & 0xf; + while (LevelMask) { + if (LevelMask & 1) + MDOpsCacheProp.push_back(buildSpirvDecorCacheProp( + Ctx, *PropName, DecorCode, CacheMode, CacheLevel)); + ++CacheLevel; + LevelMask >>= 1; + } + ++CacheMode; + AttrVal >>= 4; } } else { FPGAProp = true; diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc index 09743c706781b..666807b16ccf7 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -96,229 +96,71 @@ Below is a list of new compile-time constant properties supported with `annotated_ptr`. ```c++ -namespace sycl::ext::intel::experimental { - -enum class level : /*unspecified*/ { - L1, - L2, - L3, - L4, -}; - -enum class cache_control_read_type : /* unspecified */ { - uncached, - cached, - streaming, - invalidate_after_read, - const_cached, -}; - -enum class cache_control_write_type : /* unspecified */ { - uncached, - write_through, - write_back, - streaming, -}; - -struct cache_control_read_uncached_key { - template - using value_t = - property_value...>; -}; - -struct cache_control_read_cached_key { - template - using value_t = - property_value...>; -}; - -struct cache_control_read_streaming_key { - template - using value_t = - property_value...>; -}; - -struct cache_control_invalidate_after_read_key { - template - using value_t = - property_value...>; +namespace sycl::ext { +namespace intel::experimental { + +using cache_level = sycl::ext::oneapi::experimental::cache_level; + +enum class cache_mode { + uncached, + cached, + streaming, + invalidate, + const_cached, + write_through, + write_back }; -struct cache_control_read_const_cached_key { - template - using value_t = - property_value...>; +struct read_hint_key { + template + using value_t = property_value; }; -struct cache_control_write_uncached_key { - template - using value_t = - property_value...>; +struct write_hint_key { + template + using value_t = property_value; }; -struct cache_control_write_through_key { - template - using value_t = - property_value...>; -}; - -struct cache_control_write_back_key { - template - using value_t = - property_value...>; -}; - -struct cache_control_write_streaming_key { - template - using value_t = - property_value...>; -}; - -template -inline constexpr cache_control_read_uncached_key::value_t cache_control_read_uncached; - -template -inline constexpr cache_control_read_cached_key::value_t cache_control_read_cached; - -template -inline constexpr cache_control_read_streaming_key::value_t cache_control_read_streaming; - -template -inline constexpr cache_control_invalidate_after_read_key::value_t cache_control_invalidate_after_read; - -template -inline constexpr cache_control_read_const_cached_key::value_t cache_control_read_const_cached; - -template -inline constexpr cache_control_write_uncached_key::value_t cache_control_write_uncached; - -template -inline constexpr cache_control_write_through_key::value_t cache_control_write_through; - -template -inline constexpr cache_control_write_back_key::value_t cache_control_write_back; +} // namespace intel::experimental -template -inline constexpr cache_control_write_streaming_key::value_t cache_control_write_streaming; +using namespace intel::experimental; +using read_hint_key = intel::experimental::read_hint_key; +template <> struct is_property_key : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; -using cache_control_read_uncached_key = intel::experimental::cache_control_read_uncached_key; -template<> -struct is_property_key : std::true_type {}; +using write_hint_key = intel::experimental::write_hint_key; +template <> struct is_property_key : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; -using cache_control_read_cached_key = intel::experimental::cache_control_read_cached_key; -template<> -struct is_property_key : std::true_type {}; +template +inline constexpr read_hint_key::value_t read_hint; -using cache_control_read_streaming_key = intel::experimental::cache_control_read_streaming_key; -template<> -struct is_property_key : std::true_type {}; - -using cache_control_invalidate_after_read_key = intel::experimental::cache_control_invalidate_after_read_key; -template<> -struct is_property_key : std::true_type {}; - -using cache_contcache_control_read_const_cached_keyrol_read_key = intel::experimental::cache_control_read_const_cached_key; -template<> -struct is_property_key : std::true_type {}; - -using cache_control_write_uncached_key = intel::experimental::cache_control_write_uncached_key; -template<> -struct is_property_key : std::true_type {}; - -using cache_control_write_through_key = intel::experimental::cache_control_write_through_key; -template<> -struct is_property_key : std::true_type {}; - -using cache_controcache_control_write_back_key = intel::experimental::cache_control_write_back_key; -template<> -struct is_property_key : std::true_type {}; - -using cache_control_write_streaming_key = intel::experimental::cache_control_write_streaming_key; -template<> -struct is_property_key : std::true_type {}; - -template -struct is_property_key_of< - cache_control_read_uncached_key, annotated_ptr> : std::true_type {}; - -template -struct is_property_key_of< - cache_control_read_cached_key, annotated_ptr> : std::true_type {}; - -template -struct is_property_key_of< - cache_control_read_streaming_key, annotated_ptr> : std::true_type {}; - -template -struct is_property_key_of< - cache_control_invalidate_after_read_key, annotated_ptr> : std::true_type {}; - -template -struct is_property_key_of< - cache_control_read_const_cached_key, annotated_ptr> : std::true_type {}; - - -template -struct is_property_key_of< - cache_control_write_uncached_key, annotated_ptr> : std::true_type {}; - -template -struct is_property_key_of< - cache_control_write_through_key, annotated_ptr> : std::true_type {}; - -template -struct is_property_key_of< - cache_control_write_back_key, annotated_ptr> : std::true_type {}; - -template -struct is_property_key_of< - cache_control_write_streaming_key, annotated_ptr> : std::true_type {}; - - -template -inline constexpr cache_control_read_uncached_key::value_t cache_control_read_uncached_key; - -template -inline constexpr cache_control_read_cached_key::value_t cache_control_read_cached_key; - -template -inline constexpr cache_control_read_streaming_key::value_t cache_control_read_streaming_key; - -template -inline constexpr cache_control_invalidate_after_read_key::value_t cache_control_invalidate_after_read_key; - -template -inline constexpr cache_control_read_const_cached_key::value_t cache_control_read_const_cached_key; - -template -inline constexpr cache_control_write_uncached_key::value_t cache_control_write_uncached_key; - -template -inline constexpr cache_control_write_through_key::value_t cache_control_write_through_key; - -template -inline constexpr cache_control_write_back_key::value_t cache_control_write_back_key; - -template -inline constexpr cache_control_write_streaming_key::value_t cache_control_write_streaming_key; +template +inline constexpr write_hint_key::value_t write_hint; } // namespace sycl::ext::intel::experimental ``` -Each of these properties takes a cache level parameter indicating which level -of the cache hierarchy is affected. Cache level 0 indicates the cache closest -to the processing unit, cache level 1 indicates the next furthest cache -level, etc. It is legal to specify a cache level that does not exist on +Each of these properties takes a cache control parameter indicating +the cache control mode and a list of cache levels the control applies to. +Cache level L1 indicates the cache closest to the processing unit, +cache level L2 indicates the next furthest cache level, etc. +It is legal to specify a cache level that does not exist on the target device, but the property will be ignored in this case. -Note that a property specifies the cache behavior only for the indicated -cache level. In order to specify the behavior for multiple cache levels, -multiple properties should be specified. - It is legal to specify several different cache control properties in the -same `annotated_ptr`. However, all instances of `cache_control_read_type` must -have different cache levels and all instances of `cache_control_write_type` -must have difference cache levels. +same `annotated_ptr`. However, at any cache level there should be +at most one `cache_mode` in `read_hint` and one `cache_mode` in `write_hint`. The cache control properties are divided into two categories: those that are hints and those that are assertions by the application. @@ -335,7 +177,7 @@ affect the performance of device code, but they do not change the semantics. a| [source] ---- -cache_control_read_uncached +read_hint> ---- | This property requests that loads from memory through the `annotated_ptr` @@ -343,7 +185,7 @@ should not cache the data at levels `Ls` in the memory hierarchy. a| [source] ---- -cache_control_read_cached +read_hint> ---- | This property requests that loads from memory through the `annotated_ptr` @@ -351,7 +193,7 @@ may cache the data at levels `Ls` in the memory hierarchy. a| [source] ---- -cache_control_read_streaming +read_hint> ---- | This property requests that loads from memory through the `annotated_ptr` @@ -361,7 +203,7 @@ property. a| [source] ---- -cache_control_write_uncached +write_hint> ---- | This property requests that writes to memory through the `annotated_ptr` @@ -369,7 +211,7 @@ should not cache the data at levels `Ls` in the memory hierarchy. a| [source] ---- -cache_control_write_through +write_hint> ---- | This property requests that writes to memory through the `annotated_ptr` @@ -378,7 +220,7 @@ and mark the cache line at levels `Ls` as "not dirty". a| [source] ---- -cache_control_write_back +write_hint> ---- | This property requests that writes to memory through the `annotated_ptr` @@ -388,7 +230,7 @@ level higher than `Ls`. a| [source] ---- -cache_control_write_streaming +write_hint> ---- | This property is the same as `write_through`, but requests use of a @@ -410,7 +252,7 @@ they are misused. a| [source] ---- -cache_control_invalidate_after_read +read_hint> ---- | This property asserts that the cache line into which data is loaded @@ -421,7 +263,7 @@ the cache line and discard "dirty" data. If the assertion is violated a| [source] ---- -cache_control_read_const_cached +read_hint> ---- | This property asserts that the cache line containing the data @@ -434,8 +276,6 @@ is undefined. == Implementation notes -It is intended that the SYCL cache control properties will be used by the -compiler to generate SPIR-V cache control operations. Alternatively, the -properties could be implemented by generating intrinsic function calls -that match the cache control types. +The SYCL cache control properties will be used by the +compiler to generate SPIR-V cache control operations. diff --git a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp index 5010fa7a82aa5..c3eacc98b9f64 100755 --- a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include #include @@ -19,27 +20,59 @@ namespace ext { namespace intel { namespace experimental { -enum class level : std::uint16_t { L1 = 0, L2, L3, L4 }; +using cache_level = sycl::ext::oneapi::experimental::cache_level; +enum class cache_mode : std::uint16_t { + uncached, + cached, + streaming, + invalidate, + const_cached, + write_through, + write_back +}; + +namespace detail { + +template static constexpr void checkLevel1() { + static_assert(count < 2, "Duplicate cache_level L1 specification"); +} +template static constexpr void checkLevel2() { + static_assert(count < 2, "Duplicate cache_level L2 specification"); +} +template static constexpr void checkLevel3() { + static_assert(count < 2, "Duplicate cache_level L3 specification"); +} +template static constexpr void checkLevel4() { + static_assert(count < 2, "Duplicate cache_level L4 specification"); +} + +} // namespace detail + +template struct cache_control { + static constexpr const int countL1 = ((Ls == cache_level::L1 ? 1 : 0) + ...); + static constexpr const int countL2 = ((Ls == cache_level::L2 ? 1 : 0) + ...); + static constexpr const int countL3 = ((Ls == cache_level::L3 ? 1 : 0) + ...); + static constexpr const int countL4 = ((Ls == cache_level::L4 ? 1 : 0) + ...); + static constexpr const uint32_t levels = ((1 << static_cast(Ls)) | ...); + static constexpr const uint32_t encoding = + (countL1, countL2, countL3, countL4, detail::checkLevel1(), + detail::checkLevel2(), detail::checkLevel3(), + detail::checkLevel4(), levels << static_cast(M) * 4); +}; template using property_value = sycl::ext::oneapi::experimental::property_value; -#define __SYCL_CACHE_CONTROL_M1(P) \ - struct P { \ - template \ - using value_t = property_value...>; \ - }; - -__SYCL_CACHE_CONTROL_M1(cache_control_read_uncached_key) -__SYCL_CACHE_CONTROL_M1(cache_control_read_cached_key) -__SYCL_CACHE_CONTROL_M1(cache_control_read_streaming_key) -__SYCL_CACHE_CONTROL_M1(cache_control_invalidate_after_read_key) -__SYCL_CACHE_CONTROL_M1(cache_control_read_const_cached_key) -__SYCL_CACHE_CONTROL_M1(cache_control_write_uncached_key) -__SYCL_CACHE_CONTROL_M1(cache_control_write_through_key) -__SYCL_CACHE_CONTROL_M1(cache_control_write_back_key) -__SYCL_CACHE_CONTROL_M1(cache_control_write_streaming_key) +struct read_hint_key { + template + using value_t = property_value; +}; + +struct write_hint_key { + template + using value_t = property_value; +}; } // namespace experimental } // namespace intel @@ -50,113 +83,82 @@ namespace experimental { template class annotated_arg; template class annotated_ptr; -#define __SYCL_CACHE_CONTROL_M2(P) \ - using P = intel::experimental::P; \ - template <> struct is_property_key

: std::true_type {}; \ - template \ - struct is_property_key_of> \ - : std::true_type {}; \ - template \ - struct is_property_key_of> \ - : std::true_type {}; - -__SYCL_CACHE_CONTROL_M2(cache_control_read_uncached_key) -__SYCL_CACHE_CONTROL_M2(cache_control_read_cached_key) -__SYCL_CACHE_CONTROL_M2(cache_control_read_streaming_key) -__SYCL_CACHE_CONTROL_M2(cache_control_invalidate_after_read_key) -__SYCL_CACHE_CONTROL_M2(cache_control_read_const_cached_key) -__SYCL_CACHE_CONTROL_M2(cache_control_write_uncached_key) -__SYCL_CACHE_CONTROL_M2(cache_control_write_through_key) -__SYCL_CACHE_CONTROL_M2(cache_control_write_back_key) -__SYCL_CACHE_CONTROL_M2(cache_control_write_streaming_key) +using read_hint_key = intel::experimental::read_hint_key; +template <> struct is_property_key : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; + +using write_hint_key = intel::experimental::write_hint_key; +template <> struct is_property_key : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; using namespace intel::experimental; -#define __SYCL_CACHE_CONTROL_M3(P) \ - template inline constexpr P##_key::value_t P; +template +inline constexpr read_hint_key::value_t read_hint; -__SYCL_CACHE_CONTROL_M3(cache_control_read_uncached) -__SYCL_CACHE_CONTROL_M3(cache_control_read_cached) -__SYCL_CACHE_CONTROL_M3(cache_control_read_streaming) -__SYCL_CACHE_CONTROL_M3(cache_control_invalidate_after_read) -__SYCL_CACHE_CONTROL_M3(cache_control_read_const_cached) -__SYCL_CACHE_CONTROL_M3(cache_control_write_uncached) -__SYCL_CACHE_CONTROL_M3(cache_control_write_through) -__SYCL_CACHE_CONTROL_M3(cache_control_write_back) -__SYCL_CACHE_CONTROL_M3(cache_control_write_streaming) +template +inline constexpr write_hint_key::value_t write_hint; namespace detail { -template static constexpr void checkLevels() {} -template static constexpr void checkLevels() { - static_assert(L_1 != L_2, "Duplicate cache level specification."); -} -template static constexpr void checkLevels() { - static_assert(L_1 != L_2 && L_1 != L_3 && L_2 != L_3, - "Duplicate cache level specification."); +static constexpr int countL(int levels, int mask) { + return levels & mask ? 1 : 0; } -template -static constexpr void checkLevels() { - static_assert(L_1 != L_2 && L_1 != L_3 && L_1 != L_4 && L_2 != L_3 && - L_2 != L_4 && L_3 != L_4, - "Duplicate cache level specification."); +template +static constexpr void checkUnique() { + static_assert(countL1 < 2, "Conflicting cache_mode at L1"); + static_assert(countL2 < 2, "Conflicting cache_mode at L2"); + static_assert(countL3 < 2, "Conflicting cache_mode at L3"); + static_assert(countL4 < 2, "Conflicting cache_mode at L4"); } -#define __SYCL_CACHE_CONTROL_M4(P, K, N) \ - template <> struct PropertyToKind

{ \ - static constexpr PropKind Kind = PropKind::K; \ - }; \ - template <> struct IsCompileTimeProperty

: std::true_type {}; \ - template struct PropertyMetaInfo> { \ - static constexpr const char *name = N; \ - static constexpr const int value = \ - (checkLevels(), ((1 << static_cast(Ls)) | ...)); \ - }; - -__SYCL_CACHE_CONTROL_M4(cache_control_read_uncached_key, - CacheControlReadUncached, "sycl-cache-read-uncached") -__SYCL_CACHE_CONTROL_M4(cache_control_read_cached_key, CacheControlReadCached, - "sycl-cache-read-cached") -__SYCL_CACHE_CONTROL_M4(cache_control_read_streaming_key, - CacheControlReadStreaming, "sycl-cache-read-streaming") -__SYCL_CACHE_CONTROL_M4(cache_control_invalidate_after_read_key, - CacheControlReadInvalidateAfterRead, - "sycl-cache-read-invalidate-after-read") -__SYCL_CACHE_CONTROL_M4(cache_control_read_const_cached_key, - CacheControlReadConstCached, - "sycl-cache-read-const-cached") -__SYCL_CACHE_CONTROL_M4(cache_control_write_uncached_key, - CacheControlWriteUncached, "sycl-cache-write-uncached") -__SYCL_CACHE_CONTROL_M4(cache_control_write_through_key, - CacheControlWriteThrough, "sycl-cache-write-through") -__SYCL_CACHE_CONTROL_M4(cache_control_write_back_key, CacheControlWriteBack, - "sycl-cache-write-back") -__SYCL_CACHE_CONTROL_M4(cache_control_write_streaming_key, - CacheControlWriteStreaming, - "sycl-cache-write-streaming") +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::CacheControlRead; +}; +template <> struct IsCompileTimeProperty : std::true_type {}; +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-cache-read-hint"; + static constexpr const int value = + (checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), + (countL(Cs::levels, 4) + ...), + (countL(Cs::levels, 8) + ...)>(), + ((Cs::encoding) | ...)); +}; + +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::CacheControlWrite; +}; +template <> struct IsCompileTimeProperty : std::true_type {}; +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-cache-write-hint"; + static constexpr const int value = + (checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), + (countL(Cs::levels, 4) + ...), + (countL(Cs::levels, 8) + ...)>(), + ((Cs::encoding) | ...)); +}; } // namespace detail -#define __SYCL_CACHE_CONTROL_M5(P) \ - template \ - struct is_valid_property> \ - : std::bool_constant::value> {}; - -__SYCL_CACHE_CONTROL_M5(cache_control_read_uncached_key) -__SYCL_CACHE_CONTROL_M5(cache_control_read_cached_key) -__SYCL_CACHE_CONTROL_M5(cache_control_read_streaming_key) -__SYCL_CACHE_CONTROL_M5(cache_control_invalidate_after_read_key) -__SYCL_CACHE_CONTROL_M5(cache_control_read_const_cached_key) -__SYCL_CACHE_CONTROL_M5(cache_control_write_uncached_key) -__SYCL_CACHE_CONTROL_M5(cache_control_write_through_key) -__SYCL_CACHE_CONTROL_M5(cache_control_write_back_key) -__SYCL_CACHE_CONTROL_M5(cache_control_write_streaming_key) - -#undef __SYCL_CACHE_CONTROL_M1 -#undef __SYCL_CACHE_CONTROL_M2 -#undef __SYCL_CACHE_CONTROL_M3 -#undef __SYCL_CACHE_CONTROL_M4 -#undef __SYCL_CACHE_CONTROL_M5 +template +struct is_valid_property> + : std::bool_constant::value> {}; + +template +struct is_valid_property> + : std::bool_constant::value> {}; } // namespace experimental } // namespace oneapi diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp index 1eb575f328178..850dbbe764d9b 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp @@ -159,22 +159,10 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr> { using allowed_properties = std::tuple< decltype(ext::intel::experimental::buffer_location<0>), decltype(ext::oneapi::experimental::alignment<0>), - decltype(ext::oneapi::experimental::cache_control_read_cached), - decltype(ext::oneapi::experimental::cache_control_read_uncached< - level::L1>), - decltype(ext::oneapi::experimental::cache_control_read_streaming< - level::L1>), - decltype(ext::oneapi::experimental::cache_control_invalidate_after_read< - level::L1>), - decltype(ext::oneapi::experimental::cache_control_read_const_cached< - level::L1>), - decltype(ext::oneapi::experimental::cache_control_write_uncached< - level::L1>), - decltype(ext::oneapi::experimental::cache_control_write_streaming< - level::L1>), - decltype(ext::oneapi::experimental::cache_control_write_through< - level::L1>), - decltype(ext::oneapi::experimental::cache_control_write_back)>; + decltype(ext::oneapi::experimental::read_hint< + cache_control>), + decltype(ext::oneapi::experimental::write_hint< + cache_control>)>; using filtered_properties = typename PropertiesFilter::tuple; diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index 4116db5b3904c..d8f602cec8ec6 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -214,17 +214,10 @@ enum PropKind : uint32_t { ESIMDL1CacheHint = 44, ESIMDL2CacheHint = 45, ESIMDL3CacheHint = 46, - CacheControlReadCached = 47, - CacheControlReadUncached = 48, - CacheControlReadStreaming = 49, - CacheControlReadInvalidateAfterRead = 50, - CacheControlReadConstCached = 51, - CacheControlWriteUncached = 52, - CacheControlWriteStreaming = 53, - CacheControlWriteThrough = 54, - CacheControlWriteBack = 55, + CacheControlRead = 47, + CacheControlWrite = 48, // PropKindSize must always be the last value. - PropKindSize = 56, + PropKindSize = 49, }; // This trait must be specialized for all properties and must have a unique diff --git a/sycl/test/extensions/properties/properties_cache_control.cpp b/sycl/test/extensions/properties/properties_cache_control.cpp index af95dcb4a6cc9..cf6364767a9bd 100755 --- a/sycl/test/extensions/properties/properties_cache_control.cpp +++ b/sycl/test/extensions/properties/properties_cache_control.cpp @@ -8,16 +8,21 @@ using namespace sycl; using namespace ext::oneapi::experimental; using namespace ext::intel::experimental; -using annotated_ptr_load = - annotated_ptr, cache_control_read_cached, - cache_control_read_uncached, - cache_control_invalidate_after_read))>; +using annotated_ptr_load = annotated_ptr< + float, + decltype(properties( + alignment<8>, + read_hint, + cache_control, + cache_control>))>; -using annotated_ptr_store = - annotated_ptr, - cache_control_write_back))>; +using annotated_ptr_store = annotated_ptr< + float, + decltype(properties( + write_hint, + cache_control>))>; void cache_control_read_func() { queue q; @@ -54,9 +59,9 @@ void cache_control_write_func() { // CHECK-IR: ret void // CHECK-IR: [[RDECOR]] = !{[[RDECOR1:.*]], [[RDECOR2:.*]], [[RDECOR3:.*]], [[RDECOR4:.*]]} -// CHECK-IR: [[RDECOR1]] = !{i32 6442, i32 0, i32 1} -// CHECK-IR: [[RDECOR2]] = !{i32 6442, i32 1, i32 0} -// CHECK-IR: [[RDECOR3]] = !{i32 6442, i32 2, i32 0} +// CHECK-IR: [[RDECOR1]] = !{i32 6442, i32 1, i32 0} +// CHECK-IR: [[RDECOR2]] = !{i32 6442, i32 2, i32 0} +// CHECK-IR: [[RDECOR3]] = !{i32 6442, i32 0, i32 1} // CHECK-IR: [[RDECOR4]] = !{i32 6442, i32 3, i32 3} // CHECK-IR: [[WDECOR]] = !{[[WDECOR1:.*]], [[WDECOR2:.*]], [[WDECOR3:.*]]} From 0a54f88baf9a3ad1221aff578bf55eab892614ce Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Tue, 24 Oct 2023 19:05:56 -0700 Subject: [PATCH 08/28] Added missing #include. --- .../oneapi/experimental/annotated_ptr/annotated_ptr.hpp | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp index 850dbbe764d9b..e1b5268e78785 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include #include @@ -159,10 +160,8 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr> { using allowed_properties = std::tuple< decltype(ext::intel::experimental::buffer_location<0>), decltype(ext::oneapi::experimental::alignment<0>), - decltype(ext::oneapi::experimental::read_hint< - cache_control>), - decltype(ext::oneapi::experimental::write_hint< - cache_control>)>; + decltype(read_hint>), + decltype(write_hint>)>; using filtered_properties = typename PropertiesFilter::tuple; From 5d19fae469c858cd57e451b1af02bfa7669f684c Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Thu, 26 Oct 2023 12:37:30 -0700 Subject: [PATCH 09/28] Added checks for cache_modes used in read and write hints. --- .../SYCLLowerIR/CompileTimePropertiesPass.cpp | 52 +++++-------------- .../sycl_ext_intel_cache_controls.asciidoc | 6 +-- .../experimental/cache_control_properties.hpp | 35 ++++++++++--- 3 files changed, 46 insertions(+), 47 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index 4ab2488a70712..3f45f591e3381 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -131,52 +131,28 @@ MDNode *buildSpirvDecorCacheProp(LLVMContext &Ctx, StringRef Name, write_streaming = 3 }; // SYCL encodings of read/write control - enum cache_mode { + enum class cache_mode { uncached, cached, streaming, invalidate, - const_cached, - through, - back + constant, + write_through, + write_back }; + static uint32_t SPIRVReadControl[] = {read_uncached, read_cached, + read_streaming, read_invalidate, + read_const_cached}; + static uint32_t SPIRVWriteControl[] = { + write_uncached, write_uncached, write_streaming, write_uncached, + write_uncached, write_through, write_back}; // Map SYCL encoding to SPIR-V uint32_t CacheProp; - if (Name == "sycl-cache-read-hint") { - switch (CacheMode) { - case uncached: - CacheProp = read_uncached; - break; - case cached: - CacheProp = read_cached; - break; - case streaming: - CacheProp = read_streaming; - break; - case invalidate: - CacheProp = read_invalidate; - break; - case const_cached: - CacheProp = read_const_cached; - break; - } - } else { - switch (CacheMode) { - case uncached: - CacheProp = write_uncached; - break; - case through: - CacheProp = write_through; - break; - case back: - CacheProp = write_back; - break; - case streaming: - CacheProp = write_streaming; - break; - } - } + if (Name == "sycl-cache-read-hint") + CacheProp = SPIRVReadControl[CacheMode]; + else + CacheProp = SPIRVWriteControl[CacheMode]; auto *Ty = Type::getInt32Ty(Ctx); SmallVector MD; diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc index 666807b16ccf7..09ac0519f19e3 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -106,7 +106,7 @@ enum class cache_mode { cached, streaming, invalidate, - const_cached, + constant, write_through, write_back }; @@ -215,7 +215,7 @@ write_hint> ---- | This property requests that writes to memory through the `annotated_ptr` -should immediately write the data to the next-level cache after `L` +should immediately write the data to the next-level cache after `Ls` and mark the cache line at levels `Ls` as "not dirty". a| [source] @@ -263,7 +263,7 @@ the cache line and discard "dirty" data. If the assertion is violated a| [source] ---- -read_hint> +read_hint> ---- | This property asserts that the cache line containing the data diff --git a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp index c3eacc98b9f64..13f4db1ce4d51 100755 --- a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp @@ -21,12 +21,13 @@ namespace intel { namespace experimental { using cache_level = sycl::ext::oneapi::experimental::cache_level; -enum class cache_mode : std::uint16_t { + +enum class cache_mode { uncached, cached, streaming, invalidate, - const_cached, + constant, write_through, write_back }; @@ -49,6 +50,7 @@ template static constexpr void checkLevel4() { } // namespace detail template struct cache_control { + static constexpr const auto mode = M; static constexpr const int countL1 = ((Ls == cache_level::L1 ? 1 : 0) + ...); static constexpr const int countL2 = ((Ls == cache_level::L2 ? 1 : 0) + ...); static constexpr const int countL3 = ((Ls == cache_level::L3 ? 1 : 0) + ...); @@ -101,8 +103,6 @@ template struct is_property_key_of> : std::true_type {}; -using namespace intel::experimental; - template inline constexpr read_hint_key::value_t read_hint; @@ -114,6 +114,7 @@ namespace detail { static constexpr int countL(int levels, int mask) { return levels & mask ? 1 : 0; } + template static constexpr void checkUnique() { static_assert(countL1 < 2, "Conflicting cache_mode at L1"); @@ -122,6 +123,26 @@ static constexpr void checkUnique() { static_assert(countL4 < 2, "Conflicting cache_mode at L4"); } +using cache_mode = sycl::ext::intel::experimental::cache_mode; + +template static constexpr int checkReadMode() { + static_assert(M != cache_mode::write_back, + "read_hint cannot specify cache_mode::write_back"); + static_assert(M != cache_mode::write_through, + "read_hint cannot specify cache_mode::write_through"); + return 0; +} + +template static constexpr int checkWriteMode() { + static_assert(M != cache_mode::cached, + "write_hint cannot specify cache_mode::cached"); + static_assert(M != cache_mode::invalidate, + "write_hint cannot specify cache_mode::validate"); + static_assert(M != cache_mode::constant, + "write_hint cannot specify cache_mode::constant"); + return 0; +} + template <> struct PropertyToKind { static constexpr PropKind Kind = PropKind::CacheControlRead; }; @@ -130,7 +151,8 @@ template struct PropertyMetaInfo> { static constexpr const char *name = "sycl-cache-read-hint"; static constexpr const int value = - (checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), + ((checkReadMode() + ...), + checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), (countL(Cs::levels, 4) + ...), (countL(Cs::levels, 8) + ...)>(), ((Cs::encoding) | ...)); @@ -144,7 +166,8 @@ template struct PropertyMetaInfo> { static constexpr const char *name = "sycl-cache-write-hint"; static constexpr const int value = - (checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), + ((checkWriteMode() + ...), + checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), (countL(Cs::levels, 4) + ...), (countL(Cs::levels, 8) + ...)>(), ((Cs::encoding) | ...)); From b8df87e8cd80e0064f83ca4a2b61f636a361a735 Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Thu, 26 Oct 2023 14:38:42 -0700 Subject: [PATCH 10/28] Minor corrections to namespace usage. --- .../experimental/cache_control_properties.hpp | 13 +++++++------ .../experimental/annotated_ptr/annotated_ptr.hpp | 16 +++++++++++----- 2 files changed, 18 insertions(+), 11 deletions(-) diff --git a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp index 13f4db1ce4d51..401bc072de7f9 100755 --- a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp @@ -31,6 +31,7 @@ enum class cache_mode { write_through, write_back }; +using cache_mode = sycl::ext::intel::experimental::cache_mode; namespace detail { @@ -76,6 +77,12 @@ struct write_hint_key { using value_t = property_value; }; +template +inline constexpr read_hint_key::value_t read_hint; + +template +inline constexpr write_hint_key::value_t write_hint; + } // namespace experimental } // namespace intel @@ -103,12 +110,6 @@ template struct is_property_key_of> : std::true_type {}; -template -inline constexpr read_hint_key::value_t read_hint; - -template -inline constexpr write_hint_key::value_t write_hint; - namespace detail { static constexpr int countL(int levels, int mask) { diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp index e1b5268e78785..0495e5f5bd247 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp @@ -157,11 +157,17 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr> { using property_list_t = detail::properties_t; // buffer_location and alignment are allowed for annotated_ref - using allowed_properties = std::tuple< - decltype(ext::intel::experimental::buffer_location<0>), - decltype(ext::oneapi::experimental::alignment<0>), - decltype(read_hint>), - decltype(write_hint>)>; + using allowed_properties = + std::tuple), + decltype(ext::oneapi::experimental::alignment<0>), + decltype(ext::intel::experimental::read_hint< + ext::intel::experimental::cache_control< + ext::intel::experimental::cache_mode::cached, + cache_level::L1>>), + decltype(ext::intel::experimental::write_hint< + ext::intel::experimental::cache_control< + ext::intel::experimental::cache_mode::cached, + cache_level::L1>>)>; using filtered_properties = typename PropertiesFilter::tuple; From 482f9be04238e6930a45ef1ab84146bdec69b103 Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Thu, 26 Oct 2023 15:00:17 -0700 Subject: [PATCH 11/28] Update to documentation. --- .../proposed/sycl_ext_intel_cache_controls.asciidoc | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc index 09ac0519f19e3..e23e19de9e147 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -99,8 +99,6 @@ Below is a list of new compile-time constant properties supported with namespace sycl::ext { namespace intel::experimental { -using cache_level = sycl::ext::oneapi::experimental::cache_level; - enum class cache_mode { uncached, cached, @@ -110,6 +108,8 @@ enum class cache_mode { write_through, write_back }; +using cache_mode = sycl::ext::oneapi::experimental::cache_mode; +using cache_level = sycl::ext::oneapi::experimental::cache_level; struct read_hint_key { template @@ -121,6 +121,12 @@ struct write_hint_key { using value_t = property_value; }; +template +inline constexpr read_hint_key::value_t read_hint; + +template +inline constexpr write_hint_key::value_t write_hint; + } // namespace intel::experimental using namespace intel::experimental; From dc9b4ef93be8a66ce71ceaaa81b73d587b2217f3 Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Mon, 30 Oct 2023 16:38:25 -0700 Subject: [PATCH 12/28] Renamed some read_hints as read_assertions. --- .../lib/SYCLLowerIR/CompileTimeProperties.def | 4 +- .../SYCLLowerIR/CompileTimePropertiesPass.cpp | 2 +- .../sycl_ext_intel_cache_controls.asciidoc | 54 ++++++++----- .../experimental/cache_control_properties.hpp | 78 +++++++++++++----- .../annotated_ptr/annotated_ptr.hpp | 4 + .../properties/properties_cache_control.cpp | 81 +++++++++++++------ 6 files changed, 154 insertions(+), 69 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/CompileTimeProperties.def b/llvm/lib/SYCLLowerIR/CompileTimeProperties.def index 3b6fc20eea49a..5f9a6fb1de5f6 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimeProperties.def +++ b/llvm/lib/SYCLLowerIR/CompileTimeProperties.def @@ -56,5 +56,5 @@ SYCL_COMPILE_TIME_PROPERTY("sycl-bi-directional-ports-true", 5885, DecorValueTy::none) // The corresponding SPIR-V OpCodes for cache control properties -SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-hint", 6442, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-cache-write-hint", 6443, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read", 6442, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-write", 6443, DecorValueTy::uint32) diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index 3f45f591e3381..f8a49c3d7300a 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -149,7 +149,7 @@ MDNode *buildSpirvDecorCacheProp(LLVMContext &Ctx, StringRef Name, // Map SYCL encoding to SPIR-V uint32_t CacheProp; - if (Name == "sycl-cache-read-hint") + if (Name == "sycl-cache-read") CacheProp = SPIRVReadControl[CacheMode]; else CacheProp = SPIRVWriteControl[CacheMode]; diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc index e23e19de9e147..4d736f7aafac2 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -108,14 +108,21 @@ enum class cache_mode { write_through, write_back }; -using cache_mode = sycl::ext::oneapi::experimental::cache_mode; +using cache_mode = sycl::ext::intel::experimental::cache_mode; using cache_level = sycl::ext::oneapi::experimental::cache_level; +template struct cache_control {}; + struct read_hint_key { template using value_t = property_value; }; +struct read_assertion_key { + template + using value_t = property_value; +}; + struct write_hint_key { template using value_t = property_value; @@ -124,6 +131,9 @@ struct write_hint_key { template inline constexpr read_hint_key::value_t read_hint; +template +inline constexpr read_assertion_key::value_t read_assertion; + template inline constexpr write_hint_key::value_t write_hint; @@ -140,6 +150,15 @@ template struct is_property_key_of> : std::true_type {}; +using read_assertion_key = intel::experimental::read_assertion_key; +template <> struct is_property_key : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; + using write_hint_key = intel::experimental::write_hint_key; template <> struct is_property_key : std::true_type {}; template @@ -149,25 +168,16 @@ template struct is_property_key_of> : std::true_type {}; -template -inline constexpr read_hint_key::value_t read_hint; - -template -inline constexpr write_hint_key::value_t write_hint; - } // namespace sycl::ext::intel::experimental ``` -Each of these properties takes a cache control parameter indicating -the cache control mode and a list of cache levels the control applies to. -Cache level L1 indicates the cache closest to the processing unit, -cache level L2 indicates the next furthest cache level, etc. -It is legal to specify a cache level that does not exist on +Each of these properties takes a `cache_control` parameter. +The `cache_control` parameter consists of a `cache_mode` +and a list of `cache_level` parameters the mode applies to. +Cache level `L1` indicates the cache closest to the processing unit, +cache level `L2` indicates the next furthest cache level, etc. +It is legal to specify a `cache_level` that does not exist on the target device, but the property will be ignored in this case. -It is legal to specify several different cache control properties in the -same `annotated_ptr`. However, at any cache level there should be -at most one `cache_mode` in `read_hint` and one `cache_mode` in `write_hint`. - The cache control properties are divided into two categories: those that are hints and those that are assertions by the application. @@ -258,7 +268,7 @@ they are misused. a| [source] ---- -read_hint> +read_assertion> ---- | This property asserts that the cache line into which data is loaded @@ -269,7 +279,7 @@ the cache line and discard "dirty" data. If the assertion is violated a| [source] ---- -read_hint> +read_assertion> ---- | This property asserts that the cache line containing the data @@ -280,8 +290,8 @@ is undefined. |==== -- -== Implementation notes - -The SYCL cache control properties will be used by the -compiler to generate SPIR-V cache control operations. +It is legal to specify several different `cache_control` properties in the +same `annotated_ptr`. However, at any cache level there should be at most +one `cache_mode` of `read_hint` or `read_assertion` type +and at most one of `write_hint` type. diff --git a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp index 401bc072de7f9..e8fb374ed7418 100755 --- a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp @@ -20,8 +20,6 @@ namespace ext { namespace intel { namespace experimental { -using cache_level = sycl::ext::oneapi::experimental::cache_level; - enum class cache_mode { uncached, cached, @@ -32,6 +30,7 @@ enum class cache_mode { write_back }; using cache_mode = sycl::ext::intel::experimental::cache_mode; +using cache_level = sycl::ext::oneapi::experimental::cache_level; namespace detail { @@ -72,6 +71,11 @@ struct read_hint_key { using value_t = property_value; }; +struct read_assertion_key { + template + using value_t = property_value; +}; + struct write_hint_key { template using value_t = property_value; @@ -80,6 +84,9 @@ struct write_hint_key { template inline constexpr read_hint_key::value_t read_hint; +template +inline constexpr read_assertion_key::value_t read_assertion; + template inline constexpr write_hint_key::value_t write_hint; @@ -101,6 +108,15 @@ template struct is_property_key_of> : std::true_type {}; +using read_assertion_key = intel::experimental::read_assertion_key; +template <> struct is_property_key : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; + using write_hint_key = intel::experimental::write_hint_key; template <> struct is_property_key : std::true_type {}; template @@ -126,21 +142,27 @@ static constexpr void checkUnique() { using cache_mode = sycl::ext::intel::experimental::cache_mode; -template static constexpr int checkReadMode() { - static_assert(M != cache_mode::write_back, - "read_hint cannot specify cache_mode::write_back"); - static_assert(M != cache_mode::write_through, - "read_hint cannot specify cache_mode::write_through"); +template static constexpr int checkReadHint() { + static_assert(M == cache_mode::uncached || M == cache_mode::cached || + M == cache_mode::streaming, + "read_hint must specify cache_mode::uncached or " + "cache_mode::cached or cache_mode::streaming"); return 0; } -template static constexpr int checkWriteMode() { - static_assert(M != cache_mode::cached, - "write_hint cannot specify cache_mode::cached"); - static_assert(M != cache_mode::invalidate, - "write_hint cannot specify cache_mode::validate"); - static_assert(M != cache_mode::constant, - "write_hint cannot specify cache_mode::constant"); +template static constexpr int checkReadAssertion() { + static_assert( + M == cache_mode::invalidate || M == cache_mode::constant, + "read_hint must specify cache_mode::invalidate or cache_mode::constant"); + return 0; +} + +template static constexpr int checkWriteHint() { + static_assert(M == cache_mode::uncached || M == cache_mode::write_through || + M == cache_mode::write_back || M == cache_mode::streaming, + "write_hint must specify cache_mode::uncached or " + "cache_mode::write_through or " + "cache_mode::write_back or cache_mode::streaming"); return 0; } @@ -150,9 +172,25 @@ template <> struct PropertyToKind { template <> struct IsCompileTimeProperty : std::true_type {}; template struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-cache-read-hint"; + static constexpr const char *name = "sycl-cache-read"; + static constexpr const int value = + ((checkReadHint() + ...), + checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), + (countL(Cs::levels, 4) + ...), + (countL(Cs::levels, 8) + ...)>(), + ((Cs::encoding) | ...)); +}; + +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::CacheControlRead; +}; +template <> +struct IsCompileTimeProperty : std::true_type {}; +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-cache-read"; static constexpr const int value = - ((checkReadMode() + ...), + ((checkReadAssertion() + ...), checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), (countL(Cs::levels, 4) + ...), (countL(Cs::levels, 8) + ...)>(), @@ -165,9 +203,9 @@ template <> struct PropertyToKind { template <> struct IsCompileTimeProperty : std::true_type {}; template struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-cache-write-hint"; + static constexpr const char *name = "sycl-cache-write"; static constexpr const int value = - ((checkWriteMode() + ...), + ((checkWriteHint() + ...), checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), (countL(Cs::levels, 4) + ...), (countL(Cs::levels, 8) + ...)>(), @@ -180,6 +218,10 @@ template struct is_valid_property> : std::bool_constant::value> {}; +template +struct is_valid_property> + : std::bool_constant::value> {}; + template struct is_valid_property> : std::bool_constant::value> {}; diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp index 0495e5f5bd247..ea34ba28f1e82 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp @@ -164,6 +164,10 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr> { ext::intel::experimental::cache_control< ext::intel::experimental::cache_mode::cached, cache_level::L1>>), + decltype(ext::intel::experimental::read_assertion< + ext::intel::experimental::cache_control< + ext::intel::experimental::cache_mode::cached, + cache_level::L1>>), decltype(ext::intel::experimental::write_hint< ext::intel::experimental::cache_control< ext::intel::experimental::cache_mode::cached, diff --git a/sycl/test/extensions/properties/properties_cache_control.cpp b/sycl/test/extensions/properties/properties_cache_control.cpp index cf6364767a9bd..4e1b8b03e7f28 100755 --- a/sycl/test/extensions/properties/properties_cache_control.cpp +++ b/sycl/test/extensions/properties/properties_cache_control.cpp @@ -8,63 +8,92 @@ using namespace sycl; using namespace ext::oneapi::experimental; using namespace ext::intel::experimental; -using annotated_ptr_load = annotated_ptr< - float, +using load_hint = annotated_ptr< + float, decltype(properties( + alignment<8>, + read_hint, + cache_control>))>; + +using load_assertion = annotated_ptr< + int, decltype(properties( alignment<8>, - read_hint, - cache_control, - cache_control>))>; + read_assertion, + cache_control>))>; -using annotated_ptr_store = annotated_ptr< +using store_hint = annotated_ptr< float, decltype(properties( write_hint, cache_control>))>; + cache_level::L3>, + cache_control>))>; -void cache_control_read_func() { +void cache_control_read_hint_func() { queue q; constexpr int N = 10; float *ArrayA = malloc_shared(N, q); q.submit([&](handler &cgh) { cgh.parallel_for<>(range<1>(N), [=](item<1> item) { auto item_id = item.get_linear_id(); - annotated_ptr_load src{&ArrayA[item_id]}; + load_hint src{&ArrayA[item_id]}; *src = 55.0f; }); }); } -void cache_control_write_func() { +void cache_control_read_assertion_func() { + queue q; + constexpr int N = 10; + int *ArrayA = malloc_shared(N, q); + q.submit([&](handler &cgh) { + cgh.parallel_for<>(range<1>(N), [=](item<1> item) { + auto item_id = item.get_linear_id(); + load_assertion src{&ArrayA[item_id]}; + *src = 66; + }); + }); +} + +void cache_control_write_hint_func() { queue q; constexpr int N = 10; float *ArrayA = malloc_shared(N, q); q.submit([&](handler &cgh) { cgh.parallel_for<>(range<1>(N), [=](item<1> item) { auto item_id = item.get_linear_id(); - annotated_ptr_store dst{&ArrayA[item_id]}; - *dst = 55.0f; + store_hint dst{&ArrayA[item_id]}; + *dst = 77.0f; }); }); } -// CHECK-IR: spir_kernel{{.*}}cache_control_read_func -// CHECK-IR: {{.*}}getelementptr inbounds float{{.*}}!spirv.Decorations [[RDECOR:.*]] +// CHECK-IR: spir_kernel{{.*}}cache_control_read_hint_func +// CHECK-IR: {{.*}}getelementptr inbounds float{{.*}}!spirv.Decorations [[RHINT:.*]] // CHECK-IR: ret void -// CHECK-IR: spir_kernel{{.*}}cache_control_write_func -// CHECK-IR: {{.*}}getelementptr inbounds float{{.*}}!spirv.Decorations [[WDECOR:.*]] +// CHECK-IR: spir_kernel{{.*}}cache_control_read_assertion_func +// CHECK-IR: {{.*}}getelementptr inbounds i32{{.*}}!spirv.Decorations [[RASSERT:.*]] // CHECK-IR: ret void -// CHECK-IR: [[RDECOR]] = !{[[RDECOR1:.*]], [[RDECOR2:.*]], [[RDECOR3:.*]], [[RDECOR4:.*]]} -// CHECK-IR: [[RDECOR1]] = !{i32 6442, i32 1, i32 0} -// CHECK-IR: [[RDECOR2]] = !{i32 6442, i32 2, i32 0} -// CHECK-IR: [[RDECOR3]] = !{i32 6442, i32 0, i32 1} -// CHECK-IR: [[RDECOR4]] = !{i32 6442, i32 3, i32 3} +// CHECK-IR: spir_kernel{{.*}}cache_control_write_hint_func +// CHECK-IR: {{.*}}getelementptr inbounds float{{.*}}!spirv.Decorations [[WHINT:.*]] +// CHECK-IR: ret void + +// CHECK-IR: [[RHINT]] = !{[[RHINT1:.*]], [[RHINT2:.*]], [[RHINT3:.*]]} +// CHECK-IR: [[RHINT1]] = !{i32 6442, i32 1, i32 0} +// CHECK-IR: [[RHINT2]] = !{i32 6442, i32 2, i32 0} +// CHECK-IR: [[RHINT3]] = !{i32 6442, i32 0, i32 1} + +// CHECK-IR: [[RASSERT]] = !{[[RASSERT1:.*]], [[RASSERT2:.*]], [[RASSERT3:.*]]} +// CHECK-IR: [[RASSERT1]] = !{i32 6442, i32 1, i32 3} +// CHECK-IR: [[RASSERT2]] = !{i32 6442, i32 2, i32 3} +// CHECK-IR: [[RASSERT3]] = !{i32 6442, i32 0, i32 4} -// CHECK-IR: [[WDECOR]] = !{[[WDECOR1:.*]], [[WDECOR2:.*]], [[WDECOR3:.*]]} -// CHECK-IR: [[WDECOR1]] = !{i32 6443, i32 0, i32 1} -// CHECK-IR: [[WDECOR2]] = !{i32 6443, i32 1, i32 2} -// CHECK-IR: [[WDECOR3]] = !{i32 6443, i32 2, i32 2} +// CHECK-IR: [[WHINT]] = !{[[WHINT1:.*]], [[WHINT2:.*]], [[WHINT3:.*]], [[WHINT4:.*]]} +// CHECK-IR: [[WHINT1]] = !{i32 6443, i32 3, i32 3} +// CHECK-IR: [[WHINT2]] = !{i32 6443, i32 0, i32 1} +// CHECK-IR: [[WHINT3]] = !{i32 6443, i32 1, i32 2} +// CHECK-IR: [[WHINT4]] = !{i32 6443, i32 2, i32 2} From bf9fba6d3bf582f1e3f2a5e6b41886e6c983415d Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Tue, 31 Oct 2023 20:36:57 -0700 Subject: [PATCH 13/28] Changed structure of document describing cache controls and added a test for checking cache controls. --- .../sycl_ext_intel_cache_controls.asciidoc | 208 +++++--- .../experimental/cache_control_properties.hpp | 465 +++++++++--------- .../properties_cache_control_errors.cpp | 54 ++ 3 files changed, 430 insertions(+), 297 deletions(-) create mode 100755 sycl/test/extensions/properties/properties_cache_control_errors.cpp diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc index bbc96a20074db..609b2f1d16762 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -92,8 +92,29 @@ supports. === Properties -Below is a list of new compile-time constant properties supported with -`annotated_ptr`. +This extension adds three properties: `read_hint`, `write_hint`, and +`read_assertion`. The value for each of these properties is a list +of `cache_control` structures. Each `cache_control` specifies a +cache mode and a list of cache levels to which that mode applies. +To illustrate, consider the following `read_hint` property: +```c++ +read_hint, + cache_control> +``` + +This property indicates that read operations should be uncached at +level 1 and cached at levels 2 and 3. Cache level L1 indicates the +cache closest to the processing unit, cache level L2 indicates the +next furthest cache level, etc. It is legal to specify a cache_level +that does not exist on the target device, but the cache level will +be ignored in this case. + +An `annotated_ptr` can have only one instance of each of these +cache control properties: `read_hint`, `read_assertion` and `write_hint`. + +==== Cache modes +The `cache_control` structure is used by each of the three properties +to specify a cache mode and a list of cache levels to which it applies. ```c++ namespace sycl::ext { @@ -113,30 +134,57 @@ using cache_level = sycl::ext::oneapi::experimental::cache_level; template struct cache_control {}; -struct read_hint_key { - template - using value_t = property_value; -}; +} // namespace intel::experimental +``` -struct read_assertion_key { - template - using value_t = property_value; -}; +The allowed cache modes in `read_hint` are `uncached`, `cached` or `streaming`. +`write_hint` may be `uncached`, `streaming`, `write_through` or `write_back`. +`read_assertion` is either `invalidate` or `constant`. -struct write_hint_key { +==== Cache controls + +Of the cache levels specified by the cache control properties of an +`annotated_ptr`, at each level at most one cache mode of +`read_hint`/`read_assertion` type may be specified and at most one cache +mode of `write_hint` type. + +Repeating a cache level within a cache control is an error. For example: +```c++ +read_hint> +``` + +Specifying more than one cache mode from `read_hint`/`read_assertion` +or more than one `write_hint` type at a particular cache level is an error. +For example, specifying `cached` and `uncached` at level L2: +```c++ +read_hint, + cache_control> +``` +However, a cache mode from `read_hint`/`read_assertion` and +one from `write_hint` may be specified at the same level: +```c++ +read_hint>, +write_hint> +``` + +==== Read hint property +This property is a hint requesting specific cache behavior when +loading from memory through an `annotated_ptr`. This property +can affect the performance of device code, but it does not change +the semantics. + +```c++ +namespace sycl::ext { +namespace intel::experimental { + +struct read_hint_key { template - using value_t = property_value; + using value_t = property_value; }; template inline constexpr read_hint_key::value_t read_hint; -template -inline constexpr read_assertion_key::value_t read_assertion; - -template -inline constexpr write_hint_key::value_t write_hint; - } // namespace intel::experimental namespace oneapi { @@ -151,44 +199,10 @@ template struct is_property_key_of> : std::true_type {}; -using read_assertion_key = intel::experimental::read_assertion_key; -template <> struct is_property_key : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; - -using write_hint_key = intel::experimental::write_hint_key; -template <> struct is_property_key : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; - } // namespace sycl::ext::oneapi::experimental ``` -Each of these properties takes a `cache_control` parameter. -The `cache_control` parameter consists of a `cache_mode` -and a list of `cache_level` parameters the mode applies to. -Cache level `L1` indicates the cache closest to the processing unit, -cache level `L2` indicates the next furthest cache level, etc. -It is legal to specify a `cache_level` that does not exist on -the target device, but the property will be ignored in this case. - -The cache control properties are divided into two categories: those that -are hints and those that are assertions by the application. - -==== Cache control hints -These properties are hints requesting specific cache behavior when -loading or storing to memory through the `annotated_ptr`. These properties can -affect the performance of device code, but they do not change the semantics. - -- -[options="header", cols="2,1"] +[options="header", cols="3,2"] |==== | Property | Description a| @@ -217,6 +231,47 @@ This property requests that loads from memory through the `annotated_ptr` should cache the data at cache levels `Ls`. The eviction policy is to give lower priority to data cached using this property versus the `cached` property. +|==== +-- + +==== Write hint property +This property is a hint requesting specific cache behavior when +storing to memory through an `annotated_ptr`. +This property can affect the performance of device code, but it +does not change the semantics. + +```c++ +namespace sycl::ext { +namespace intel::experimental { + +struct write_hint_key { + template + using value_t = property_value; +}; + +template +inline constexpr write_hint_key::value_t write_hint; + +} // namespace intel::experimental + +namespace oneapi { +namespace experimental { + +using write_hint_key = intel::experimental::write_hint_key; +template <> struct is_property_key : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; + +} // namespace sycl::ext::oneapi::experimental +``` +-- +[options="header", cols="3,2"] +|==== +| Property | Description a| [source] ---- @@ -256,14 +311,44 @@ via a `streaming` cache control. |==== -- -==== Assertions by the application -These properties are assertions by the application, promising that the -application accesses memory in a certain way. Care must be taken when -using these properties because they can lead to undefined behavior if -they are misused. + +==== Read assertion property +This property is an assertion by the application, promising that +the application accesses memory in a certain way. +Care must be taken when using this property because it can +lead to undefined behavior if it is misused. + +```c++ +namespace sycl::ext { +namespace intel::experimental { + +struct read_assertion_key { + template + using value_t = property_value; +}; + +template +inline constexpr read_assertion_key::value_t read_assertion; + +} // namespace intel::experimental + +namespace oneapi { +namespace experimental { + +using read_assertion_key = intel::experimental::read_assertion_key; +template <> struct is_property_key : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; + +} // namespace sycl::ext::oneapi::experimental +``` -- -[options="header", cols="3,1"] +[options="header", cols="3,2"] |==== | Property | Description a| @@ -291,8 +376,3 @@ is undefined. |==== -- - -It is legal to specify several different `cache_control` properties in the -same `annotated_ptr`. However, at any cache level there should be at most -one `cache_mode` of `read_hint` or `read_assertion` type -and at most one of `write_hint` type. diff --git a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp index e8fb374ed7418..637f5efc5f5aa 100755 --- a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp @@ -1,233 +1,232 @@ -//==--------- SYCL annotated_arg/ptr properties for caching control --------==// -// -// 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 -#include -#include -#include -#include - -namespace sycl { -inline namespace _V1 { -namespace ext { -namespace intel { -namespace experimental { - -enum class cache_mode { - uncached, - cached, - streaming, - invalidate, - constant, - write_through, - write_back -}; -using cache_mode = sycl::ext::intel::experimental::cache_mode; -using cache_level = sycl::ext::oneapi::experimental::cache_level; - -namespace detail { - -template static constexpr void checkLevel1() { - static_assert(count < 2, "Duplicate cache_level L1 specification"); -} -template static constexpr void checkLevel2() { - static_assert(count < 2, "Duplicate cache_level L2 specification"); -} -template static constexpr void checkLevel3() { - static_assert(count < 2, "Duplicate cache_level L3 specification"); -} -template static constexpr void checkLevel4() { - static_assert(count < 2, "Duplicate cache_level L4 specification"); -} - -} // namespace detail - -template struct cache_control { - static constexpr const auto mode = M; - static constexpr const int countL1 = ((Ls == cache_level::L1 ? 1 : 0) + ...); - static constexpr const int countL2 = ((Ls == cache_level::L2 ? 1 : 0) + ...); - static constexpr const int countL3 = ((Ls == cache_level::L3 ? 1 : 0) + ...); - static constexpr const int countL4 = ((Ls == cache_level::L4 ? 1 : 0) + ...); - static constexpr const uint32_t levels = ((1 << static_cast(Ls)) | ...); - static constexpr const uint32_t encoding = - (countL1, countL2, countL3, countL4, detail::checkLevel1(), - detail::checkLevel2(), detail::checkLevel3(), - detail::checkLevel4(), levels << static_cast(M) * 4); -}; - -template -using property_value = - sycl::ext::oneapi::experimental::property_value; - -struct read_hint_key { - template - using value_t = property_value; -}; - -struct read_assertion_key { - template - using value_t = property_value; -}; - -struct write_hint_key { - template - using value_t = property_value; -}; - -template -inline constexpr read_hint_key::value_t read_hint; - -template -inline constexpr read_assertion_key::value_t read_assertion; - -template -inline constexpr write_hint_key::value_t write_hint; - -} // namespace experimental -} // namespace intel - -namespace oneapi { -namespace experimental { - -template class annotated_arg; -template class annotated_ptr; - -using read_hint_key = intel::experimental::read_hint_key; -template <> struct is_property_key : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; - -using read_assertion_key = intel::experimental::read_assertion_key; -template <> struct is_property_key : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; - -using write_hint_key = intel::experimental::write_hint_key; -template <> struct is_property_key : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; - -namespace detail { - -static constexpr int countL(int levels, int mask) { - return levels & mask ? 1 : 0; -} - -template -static constexpr void checkUnique() { - static_assert(countL1 < 2, "Conflicting cache_mode at L1"); - static_assert(countL2 < 2, "Conflicting cache_mode at L2"); - static_assert(countL3 < 2, "Conflicting cache_mode at L3"); - static_assert(countL4 < 2, "Conflicting cache_mode at L4"); -} - -using cache_mode = sycl::ext::intel::experimental::cache_mode; - -template static constexpr int checkReadHint() { - static_assert(M == cache_mode::uncached || M == cache_mode::cached || - M == cache_mode::streaming, - "read_hint must specify cache_mode::uncached or " - "cache_mode::cached or cache_mode::streaming"); - return 0; -} - -template static constexpr int checkReadAssertion() { - static_assert( - M == cache_mode::invalidate || M == cache_mode::constant, - "read_hint must specify cache_mode::invalidate or cache_mode::constant"); - return 0; -} - -template static constexpr int checkWriteHint() { - static_assert(M == cache_mode::uncached || M == cache_mode::write_through || - M == cache_mode::write_back || M == cache_mode::streaming, - "write_hint must specify cache_mode::uncached or " - "cache_mode::write_through or " - "cache_mode::write_back or cache_mode::streaming"); - return 0; -} - -template <> struct PropertyToKind { - static constexpr PropKind Kind = PropKind::CacheControlRead; -}; -template <> struct IsCompileTimeProperty : std::true_type {}; -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-cache-read"; - static constexpr const int value = - ((checkReadHint() + ...), - checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), - (countL(Cs::levels, 4) + ...), - (countL(Cs::levels, 8) + ...)>(), - ((Cs::encoding) | ...)); -}; - -template <> struct PropertyToKind { - static constexpr PropKind Kind = PropKind::CacheControlRead; -}; -template <> -struct IsCompileTimeProperty : std::true_type {}; -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-cache-read"; - static constexpr const int value = - ((checkReadAssertion() + ...), - checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), - (countL(Cs::levels, 4) + ...), - (countL(Cs::levels, 8) + ...)>(), - ((Cs::encoding) | ...)); -}; - -template <> struct PropertyToKind { - static constexpr PropKind Kind = PropKind::CacheControlWrite; -}; -template <> struct IsCompileTimeProperty : std::true_type {}; -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-cache-write"; - static constexpr const int value = - ((checkWriteHint() + ...), - checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), - (countL(Cs::levels, 4) + ...), - (countL(Cs::levels, 8) + ...)>(), - ((Cs::encoding) | ...)); -}; - -} // namespace detail - -template -struct is_valid_property> - : std::bool_constant::value> {}; - -template -struct is_valid_property> - : std::bool_constant::value> {}; - -template -struct is_valid_property> - : std::bool_constant::value> {}; - -} // namespace experimental -} // namespace oneapi -} // namespace ext -} // namespace _V1 -} // namespace sycl +//==--------- SYCL annotated_arg/ptr properties for caching control --------==// +// +// 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 +#include +#include +#include +#include + +namespace sycl { +inline namespace _V1 { +namespace ext { +namespace intel { +namespace experimental { + +enum class cache_mode { + uncached, + cached, + streaming, + invalidate, + constant, + write_through, + write_back +}; +using cache_mode = sycl::ext::intel::experimental::cache_mode; +using cache_level = sycl::ext::oneapi::experimental::cache_level; + +namespace detail { + +template static constexpr void checkLevel1() { + static_assert(count < 2, "Duplicate cache_level L1 specification"); +} +template static constexpr void checkLevel2() { + static_assert(count < 2, "Duplicate cache_level L2 specification"); +} +template static constexpr void checkLevel3() { + static_assert(count < 2, "Duplicate cache_level L3 specification"); +} +template static constexpr void checkLevel4() { + static_assert(count < 2, "Duplicate cache_level L4 specification"); +} + +} // namespace detail + +template struct cache_control { + static constexpr const auto mode = M; + static constexpr const int countL1 = ((Ls == cache_level::L1 ? 1 : 0) + ...); + static constexpr const int countL2 = ((Ls == cache_level::L2 ? 1 : 0) + ...); + static constexpr const int countL3 = ((Ls == cache_level::L3 ? 1 : 0) + ...); + static constexpr const int countL4 = ((Ls == cache_level::L4 ? 1 : 0) + ...); + static constexpr const uint32_t levels = ((1 << static_cast(Ls)) | ...); + static constexpr const uint32_t encoding = + (countL1, countL2, countL3, countL4, detail::checkLevel1(), + detail::checkLevel2(), detail::checkLevel3(), + detail::checkLevel4(), levels << static_cast(M) * 4); +}; + +template +using property_value = + sycl::ext::oneapi::experimental::property_value; + +struct read_hint_key { + template + using value_t = property_value; +}; + +struct read_assertion_key { + template + using value_t = property_value; +}; + +struct write_hint_key { + template + using value_t = property_value; +}; + +template +inline constexpr read_hint_key::value_t read_hint; + +template +inline constexpr read_assertion_key::value_t read_assertion; + +template +inline constexpr write_hint_key::value_t write_hint; + +} // namespace experimental +} // namespace intel + +namespace oneapi { +namespace experimental { + +template class annotated_arg; +template class annotated_ptr; + +using read_hint_key = intel::experimental::read_hint_key; +template <> struct is_property_key : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; + +using read_assertion_key = intel::experimental::read_assertion_key; +template <> struct is_property_key : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; + +using write_hint_key = intel::experimental::write_hint_key; +template <> struct is_property_key : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; + +namespace detail { + +static constexpr int countL(int levels, int mask) { + return levels & mask ? 1 : 0; +} + +template +static constexpr void checkUnique() { + static_assert(countL1 < 2, "Conflicting cache_mode at L1"); + static_assert(countL2 < 2, "Conflicting cache_mode at L2"); + static_assert(countL3 < 2, "Conflicting cache_mode at L3"); + static_assert(countL4 < 2, "Conflicting cache_mode at L4"); +} + +using cache_mode = sycl::ext::intel::experimental::cache_mode; + +template static constexpr int checkReadHint() { + static_assert( + M == cache_mode::uncached || M == cache_mode::cached || + M == cache_mode::streaming, + "read_hint must specify cache_mode uncached, cached or streaming"); + return 0; +} + +template static constexpr int checkReadAssertion() { + static_assert( + M == cache_mode::invalidate || M == cache_mode::constant, + "read_assertion must specify cache_mode invalidate or constant"); + return 0; +} + +template static constexpr int checkWriteHint() { + static_assert(M == cache_mode::uncached || M == cache_mode::write_through || + M == cache_mode::write_back || M == cache_mode::streaming, + "write_hint must specify cache_mode uncached, write_through, " + "write_back or streaming"); + return 0; +} + +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::CacheControlRead; +}; +template <> struct IsCompileTimeProperty : std::true_type {}; +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-cache-read"; + static constexpr const int value = + ((checkReadHint() + ...), + checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), + (countL(Cs::levels, 4) + ...), + (countL(Cs::levels, 8) + ...)>(), + ((Cs::encoding) | ...)); +}; + +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::CacheControlRead; +}; +template <> +struct IsCompileTimeProperty : std::true_type {}; +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-cache-read"; + static constexpr const int value = + ((checkReadAssertion() + ...), + checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), + (countL(Cs::levels, 4) + ...), + (countL(Cs::levels, 8) + ...)>(), + ((Cs::encoding) | ...)); +}; + +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::CacheControlWrite; +}; +template <> struct IsCompileTimeProperty : std::true_type {}; +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-cache-write"; + static constexpr const int value = + ((checkWriteHint() + ...), + checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), + (countL(Cs::levels, 4) + ...), + (countL(Cs::levels, 8) + ...)>(), + ((Cs::encoding) | ...)); +}; + +} // namespace detail + +template +struct is_valid_property> + : std::bool_constant::value> {}; + +template +struct is_valid_property> + : std::bool_constant::value> {}; + +template +struct is_valid_property> + : std::bool_constant::value> {}; + +} // namespace experimental +} // namespace oneapi +} // namespace ext +} // namespace _V1 +} // namespace sycl diff --git a/sycl/test/extensions/properties/properties_cache_control_errors.cpp b/sycl/test/extensions/properties/properties_cache_control_errors.cpp new file mode 100755 index 0000000000000..879a1b6a4e6d5 --- /dev/null +++ b/sycl/test/extensions/properties/properties_cache_control_errors.cpp @@ -0,0 +1,54 @@ +// RUN: %clangxx -fsycl-device-only -Xclang -fsycl-is-device -ferror-limit=0 \ +// RUN: -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s + +#include + +using namespace sycl; +using namespace ext::oneapi::experimental; +using namespace ext::intel::experimental; + +using annotated_ptr_load1 = annotated_ptr< + float, decltype(properties( + read_hint>))>; +using annotated_ptr_load2 = annotated_ptr< + float, + decltype(properties( + read_hint, + cache_control>))>; +using annotated_ptr_load3 = annotated_ptr< + float, + decltype(properties( + read_hint>))>; +using annotated_ptr_load4 = + annotated_ptr>))>; +using annotated_ptr_load5 = annotated_ptr< + float, + decltype(properties( + write_hint>))>; + +void cache_control_read_func(queue q) { + float *ArrayA = malloc_shared(10, q); + q.submit([&](handler &cgh) { + cgh.single_task<>([=]() { + // expected-error@sycl/ext/intel/experimental/cache_control_properties.hpp:* {{Duplicate cache_level L3 specification}} + annotated_ptr_load1 src1{&ArrayA[0]}; + + // expected-error@sycl/ext/intel/experimental/cache_control_properties.hpp:* {{Conflicting cache_mode at L3}} + annotated_ptr_load2 src2{&ArrayA[0]}; + + // expected-error@sycl/ext/intel/experimental/cache_control_properties.hpp:* {{read_hint must specify cache_mode uncached, cached or streaming}} + annotated_ptr_load3 src3{&ArrayA[0]}; + + // expected-error@sycl/ext/intel/experimental/cache_control_properties.hpp:* {{read_assertion must specify cache_mode invalidate or constant}} + annotated_ptr_load4 src4{&ArrayA[0]}; + + // expected-error@sycl/ext/intel/experimental/cache_control_properties.hpp:* {{write_hint must specify cache_mode uncached, write_through, write_back or streaming}} + annotated_ptr_load5 src5{&ArrayA[0]}; + }); + }); +} + From be39b74a88f89791da20a76db05d26c872014d37 Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Tue, 31 Oct 2023 20:43:44 -0700 Subject: [PATCH 14/28] Formatting change. --- .../extensions/properties/properties_cache_control_errors.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/test/extensions/properties/properties_cache_control_errors.cpp b/sycl/test/extensions/properties/properties_cache_control_errors.cpp index 879a1b6a4e6d5..b3de0a503f7f5 100755 --- a/sycl/test/extensions/properties/properties_cache_control_errors.cpp +++ b/sycl/test/extensions/properties/properties_cache_control_errors.cpp @@ -51,4 +51,3 @@ void cache_control_read_func(queue q) { }); }); } - From a0176304233b86db5aa21077623eea861bde1fb6 Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Thu, 2 Nov 2023 14:47:15 -0700 Subject: [PATCH 15/28] Removed annotated_arg changes. --- .../lib/SYCLLowerIR/CompileTimeProperties.def | 5 +- .../SYCLLowerIR/CompileTimePropertiesPass.cpp | 2 +- .../sycl_ext_intel_cache_controls.asciidoc | 41 +++++----- .../experimental/cache_control_properties.hpp | 79 +++++++++---------- .../annotated_arg/annotated_arg.hpp | 2 +- .../sycl/ext/oneapi/properties/property.hpp | 7 +- 6 files changed, 66 insertions(+), 70 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/CompileTimeProperties.def b/llvm/lib/SYCLLowerIR/CompileTimeProperties.def index 34050c2db94b5..02417d366fbff 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimeProperties.def +++ b/llvm/lib/SYCLLowerIR/CompileTimeProperties.def @@ -62,5 +62,6 @@ SYCL_COMPILE_TIME_PROPERTY("sycl-prefetch-hint", 6442, DecorValueTy::uint32) SYCL_COMPILE_TIME_PROPERTY("sycl-prefetch-hint-nt", 6442, DecorValueTy::uint32) // The corresponding SPIR-V OpCodes for cache control properties -SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read", 6442, DecorValueTy::uint32) -SYCL_COMPILE_TIME_PROPERTY("sycl-cache-write", 6443, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-hint", 6442, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-assertion", 6442, DecorValueTy::uint32) +SYCL_COMPILE_TIME_PROPERTY("sycl-cache-write-hint", 6443, DecorValueTy::uint32) diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index 15d84b6d79088..cfb03c997dbd6 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -149,7 +149,7 @@ MDNode *buildSpirvDecorCacheProp(LLVMContext &Ctx, StringRef Name, // Map SYCL encoding to SPIR-V uint32_t CacheProp; - if (Name == "sycl-cache-read") + if (Name.starts_with("sycl-cache-read")) CacheProp = SPIRVReadControl[CacheMode]; else CacheProp = SPIRVWriteControl[CacheMode]; diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc index 609b2f1d16762..5c7ae735b1e29 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -109,8 +109,12 @@ next furthest cache level, etc. It is legal to specify a cache_level that does not exist on the target device, but the cache level will be ignored in this case. -An `annotated_ptr` can have only one instance of each of these -cache control properties: `read_hint`, `read_assertion` and `write_hint`. +Note that a property list may contain at most one instance of any +particular property. For example, it is not valid for a property list +to contain multiple `read_hint` properties. In order to specify multiple +"read hint" cache controls in the same property list, use a single +`read_hint` property with several `cache_control` structures. +The same rule applies to `write_hint` and `read_assertion`. ==== Cache modes The `cache_control` structure is used by each of the three properties @@ -129,7 +133,6 @@ enum class cache_mode { write_through, write_back }; -using cache_mode = sycl::ext::intel::experimental::cache_mode; using cache_level = sycl::ext::oneapi::experimental::cache_level; template struct cache_control {}; @@ -190,14 +193,11 @@ inline constexpr read_hint_key::value_t read_hint; namespace oneapi { namespace experimental { -using read_hint_key = intel::experimental::read_hint_key; -template <> struct is_property_key : std::true_type {}; +template <> +struct is_property_key : std::true_type {}; template -struct is_property_key_of> - : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; +struct is_property_key_of> : std::true_type {}; } // namespace sycl::ext::oneapi::experimental ``` @@ -257,14 +257,12 @@ inline constexpr write_hint_key::value_t write_hint; namespace oneapi { namespace experimental { -using write_hint_key = intel::experimental::write_hint_key; -template <> struct is_property_key : std::true_type {}; -template -struct is_property_key_of> +template <> +struct is_property_key : std::true_type {}; template -struct is_property_key_of> - : std::true_type {}; +struct is_property_key_of> : std::true_type {}; } // namespace sycl::ext::oneapi::experimental ``` @@ -335,14 +333,11 @@ inline constexpr read_assertion_key::value_t read_assertion; namespace oneapi { namespace experimental { -using read_assertion_key = intel::experimental::read_assertion_key; -template <> struct is_property_key : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; +template <> +struct is_property_key : std::true_type {}; template -struct is_property_key_of> - : std::true_type {}; +struct is_property_key_of> : std::true_type {}; } // namespace sycl::ext::oneapi::experimental ``` diff --git a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp index 637f5efc5f5aa..d5a7002c43047 100755 --- a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp @@ -1,4 +1,4 @@ -//==--------- SYCL annotated_arg/ptr properties for caching control --------==// +//==--------- SYCL annotated_ptr properties for caching control ------------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -14,6 +14,8 @@ #include #include +#include + namespace sycl { inline namespace _V1 { namespace ext { @@ -29,7 +31,6 @@ enum class cache_mode { write_through, write_back }; -using cache_mode = sycl::ext::intel::experimental::cache_mode; using cache_level = sycl::ext::oneapi::experimental::cache_level; namespace detail { @@ -96,35 +97,26 @@ inline constexpr write_hint_key::value_t write_hint; namespace oneapi { namespace experimental { -template class annotated_arg; template class annotated_ptr; -using read_hint_key = intel::experimental::read_hint_key; -template <> struct is_property_key : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; +template <> +struct is_property_key : std::true_type {}; template -struct is_property_key_of> - : std::true_type {}; +struct is_property_key_of> : std::true_type {}; -using read_assertion_key = intel::experimental::read_assertion_key; -template <> struct is_property_key : std::true_type {}; -template -struct is_property_key_of> +template <> +struct is_property_key : std::true_type {}; template -struct is_property_key_of> - : std::true_type {}; +struct is_property_key_of> : std::true_type {}; -using write_hint_key = intel::experimental::write_hint_key; -template <> struct is_property_key : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; +template <> +struct is_property_key : std::true_type {}; template -struct is_property_key_of> - : std::true_type {}; +struct is_property_key_of> : std::true_type {}; namespace detail { @@ -165,13 +157,15 @@ template static constexpr int checkWriteHint() { return 0; } -template <> struct PropertyToKind { - static constexpr PropKind Kind = PropKind::CacheControlRead; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::CacheControlReadHint; }; -template <> struct IsCompileTimeProperty : std::true_type {}; +template <> +struct IsCompileTimeProperty + : std::true_type {}; template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-cache-read"; +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-cache-read-hint"; static constexpr const int value = ((checkReadHint() + ...), checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), @@ -180,14 +174,16 @@ struct PropertyMetaInfo> { ((Cs::encoding) | ...)); }; -template <> struct PropertyToKind { - static constexpr PropKind Kind = PropKind::CacheControlRead; +template <> struct PropertyToKind { + static constexpr PropKind Kind = PropKind::CacheControlReadAssertion; }; template <> -struct IsCompileTimeProperty : std::true_type {}; +struct IsCompileTimeProperty + : std::true_type {}; template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-cache-read"; +struct PropertyMetaInfo< + intel::experimental::read_assertion_key::value_t> { + static constexpr const char *name = "sycl-cache-read-assertion"; static constexpr const int value = ((checkReadAssertion() + ...), checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), @@ -196,13 +192,15 @@ struct PropertyMetaInfo> { ((Cs::encoding) | ...)); }; -template <> struct PropertyToKind { +template <> struct PropertyToKind { static constexpr PropKind Kind = PropKind::CacheControlWrite; }; -template <> struct IsCompileTimeProperty : std::true_type {}; +template <> +struct IsCompileTimeProperty + : std::true_type {}; template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-cache-write"; +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-cache-write-hint"; static constexpr const int value = ((checkWriteHint() + ...), checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), @@ -214,15 +212,16 @@ struct PropertyMetaInfo> { } // namespace detail template -struct is_valid_property> +struct is_valid_property> : std::bool_constant::value> {}; template -struct is_valid_property> +struct is_valid_property< + T, intel::experimental::read_assertion_key::value_t> : std::bool_constant::value> {}; template -struct is_valid_property> +struct is_valid_property> : std::bool_constant::value> {}; } // namespace experimental diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp index 6b9c1313a7be9..4f92cd19f28ea 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp @@ -9,7 +9,7 @@ #pragma once #include -#include +//#include #include #include #include diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index d8f602cec8ec6..8ead37e938fc0 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -214,10 +214,11 @@ enum PropKind : uint32_t { ESIMDL1CacheHint = 44, ESIMDL2CacheHint = 45, ESIMDL3CacheHint = 46, - CacheControlRead = 47, - CacheControlWrite = 48, + CacheControlReadHint = 47, + CacheControlReadAssertion = 48, + CacheControlWrite = 49, // PropKindSize must always be the last value. - PropKindSize = 49, + PropKindSize = 50, }; // This trait must be specialized for all properties and must have a unique From ac6597c1c4cca40d1319d66eed52cfa3c7231964 Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Thu, 2 Nov 2023 18:04:36 -0700 Subject: [PATCH 16/28] Formatting change. --- .../annotated_arg/annotated_arg.hpp | 103 +++++++----------- 1 file changed, 37 insertions(+), 66 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp index afc5161a3726d..ace8f59c5d8e0 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp @@ -9,7 +9,6 @@ #pragma once #include -//#include #include #include #include @@ -55,9 +54,8 @@ annotated_arg(annotated_arg, properties>) template class annotated_arg { // This should always fail when instantiating the unspecialized version. - static constexpr bool is_valid_property_list = - is_property_list::value; - static_assert(is_valid_property_list, "Property list is invalid."); + static_assert(is_property_list::value, + "Property list is invalid."); }; // Partial specialization for pointer type @@ -85,17 +83,12 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { #endif public: - static constexpr bool is_valid_property_list = - is_property_list::value; - static_assert(is_valid_property_list, "Property list is invalid."); - static constexpr bool contains_valid_properties = - check_property_list::value; - static_assert(contains_valid_properties, + static_assert(is_property_list::value, + "Property list is invalid."); + static_assert(check_property_list::value, "The property list contains invalid property."); // check the set if FPGA specificed properties are used - static constexpr bool hasValidFPGAProperties = - detail::checkValidFPGAPropertySet::value; - static_assert(hasValidFPGAProperties, + static_assert(detail::checkValidFPGAPropertySet::value, "FPGA Interface properties (i.e. awidth, dwidth, etc.)" "can only be set with BufferLocation together."); @@ -116,12 +109,11 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { template annotated_arg(T *_ptr, const PropertyValueTs &...props) noexcept : obj(global_pointer_t(_ptr)) { - static constexpr bool has_same_properties = std::is_same< - property_list_t, - detail::merged_properties_t>::value; static_assert( - has_same_properties, + std::is_same< + property_list_t, + detail::merged_properties_t>::value, "The property list must contain all properties of the input of the " "constructor"); } @@ -133,19 +125,16 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { template explicit annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { - static constexpr bool is_input_convertible = - std::is_convertible::value; - static_assert(is_input_convertible, + static_assert(std::is_convertible::value, "The underlying data type of the input annotated_arg is not " "compatible"); - static constexpr bool has_same_properties = std::is_same< - property_list_t, - detail::merged_properties_t>::value; static_assert( - has_same_properties, - "The constructed annotated_arg type must contain all the properties " - "of the input annotated_arg"); + std::is_same< + property_list_t, + detail::merged_properties_t>::value, + "The constructed annotated_arg type must contain all the properties of " + "the input annotated_arg"); } // Constructs an annotated_arg object from another annotated_arg object and a @@ -157,17 +146,13 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { const PropertyListV &proplist) noexcept : obj(other.obj) { (void)proplist; - static constexpr bool is_input_convertible = - std::is_convertible::value; - static_assert(is_input_convertible, + static_assert(std::is_convertible::value, "The underlying data type of the input annotated_arg is not " "compatible"); - static constexpr bool has_same_properties = std::is_same< - property_list_t, - detail::merged_properties_t>::value; static_assert( - has_same_properties, + std::is_same>::value, "The property list of constructed annotated_arg type must be the union " "of the input property lists"); } @@ -207,19 +192,13 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { #endif public: - static constexpr bool is_device_copyable = is_device_copyable_v; - static_assert(is_device_copyable, "Type T must be device copyable."); - static constexpr bool is_valid_property_list = - is_property_list::value; - static_assert(is_valid_property_list, "Property list is invalid."); - static constexpr bool contains_valid_properties = - check_property_list::value; - static_assert(contains_valid_properties, + static_assert(is_device_copyable_v, "Type T must be device copyable."); + static_assert(is_property_list::value, + "Property list is invalid."); + static_assert(check_property_list::value, "The property list contains invalid property."); // check the set if FPGA specificed properties are used - static constexpr bool hasValidFPGAProperties = - detail::checkValidFPGAPropertySet::value; - static_assert(hasValidFPGAProperties, + static_assert(detail::checkValidFPGAPropertySet::value, "FPGA Interface properties (i.e. awidth, dwidth, etc.)" "can only be set with BufferLocation together."); @@ -239,12 +218,11 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { // `PropertyValueTs...` must have the same property value. template annotated_arg(const T &_obj, PropertyValueTs... props) noexcept : obj(_obj) { - static constexpr bool has_same_properties = std::is_same< - property_list_t, - detail::merged_properties_t>::value; static_assert( - has_same_properties, + std::is_same< + property_list_t, + detail::merged_properties_t>::value, "The property list must contain all properties of the input of the " "constructor"); } @@ -256,19 +234,16 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { template explicit annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { - static constexpr bool is_input_convertible = - std::is_convertible::value; - static_assert(is_input_convertible, + static_assert(std::is_convertible::value, "The underlying data type of the input annotated_arg is not " "compatible"); - static constexpr bool has_same_properties = std::is_same< - property_list_t, - detail::merged_properties_t>::value; static_assert( - has_same_properties, - "The constructed annotated_arg type must contain all the properties " - "of the input annotated_arg"); + std::is_same< + property_list_t, + detail::merged_properties_t>::value, + "The constructed annotated_arg type must contain all the properties of " + "the input annotated_arg"); } // Constructs an annotated_arg object from another annotated_arg object and a @@ -280,17 +255,13 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { const PropertyListV &proplist) noexcept : obj(other.obj) { (void)proplist; - static constexpr bool is_input_convertible = - std::is_convertible::value; - static_assert(is_input_convertible, + static_assert(std::is_convertible::value, "The underlying data type of the input annotated_arg is not " "compatible"); - static constexpr bool has_same_properties = std::is_same< - property_list_t, - detail::merged_properties_t>::value; static_assert( - has_same_properties, + std::is_same>::value, "The property list of constructed annotated_arg type must be the union " "of the input property lists"); } From f638d43436a89d7bac168e9298ac76a87933fc24 Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Thu, 2 Nov 2023 18:10:05 -0700 Subject: [PATCH 17/28] Restore damaged file. --- .../annotated_arg/annotated_arg.hpp | 102 +++++++++++------- 1 file changed, 65 insertions(+), 37 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp index ace8f59c5d8e0..cb8a3229ae38a 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp @@ -54,8 +54,9 @@ annotated_arg(annotated_arg, properties>) template class annotated_arg { // This should always fail when instantiating the unspecialized version. - static_assert(is_property_list::value, - "Property list is invalid."); + static constexpr bool is_valid_property_list = + is_property_list::value; + static_assert(is_valid_property_list, "Property list is invalid."); }; // Partial specialization for pointer type @@ -83,12 +84,17 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { #endif public: - static_assert(is_property_list::value, - "Property list is invalid."); - static_assert(check_property_list::value, + static constexpr bool is_valid_property_list = + is_property_list::value; + static_assert(is_valid_property_list, "Property list is invalid."); + static constexpr bool contains_valid_properties = + check_property_list::value; + static_assert(contains_valid_properties, "The property list contains invalid property."); // check the set if FPGA specificed properties are used - static_assert(detail::checkValidFPGAPropertySet::value, + static constexpr bool hasValidFPGAProperties = + detail::checkValidFPGAPropertySet::value; + static_assert(hasValidFPGAProperties, "FPGA Interface properties (i.e. awidth, dwidth, etc.)" "can only be set with BufferLocation together."); @@ -109,11 +115,12 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { template annotated_arg(T *_ptr, const PropertyValueTs &...props) noexcept : obj(global_pointer_t(_ptr)) { + static constexpr bool has_same_properties = std::is_same< + property_list_t, + detail::merged_properties_t>::value; static_assert( - std::is_same< - property_list_t, - detail::merged_properties_t>::value, + has_same_properties, "The property list must contain all properties of the input of the " "constructor"); } @@ -125,16 +132,19 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { template explicit annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { - static_assert(std::is_convertible::value, + static constexpr bool is_input_convertible = + std::is_convertible::value; + static_assert(is_input_convertible, "The underlying data type of the input annotated_arg is not " "compatible"); + static constexpr bool has_same_properties = std::is_same< + property_list_t, + detail::merged_properties_t>::value; static_assert( - std::is_same< - property_list_t, - detail::merged_properties_t>::value, - "The constructed annotated_arg type must contain all the properties of " - "the input annotated_arg"); + has_same_properties, + "The constructed annotated_arg type must contain all the properties " + "of the input annotated_arg"); } // Constructs an annotated_arg object from another annotated_arg object and a @@ -146,13 +156,17 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { const PropertyListV &proplist) noexcept : obj(other.obj) { (void)proplist; - static_assert(std::is_convertible::value, + static constexpr bool is_input_convertible = + std::is_convertible::value; + static_assert(is_input_convertible, "The underlying data type of the input annotated_arg is not " "compatible"); + static constexpr bool has_same_properties = std::is_same< + property_list_t, + detail::merged_properties_t>::value; static_assert( - std::is_same>::value, + has_same_properties, "The property list of constructed annotated_arg type must be the union " "of the input property lists"); } @@ -192,13 +206,19 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { #endif public: - static_assert(is_device_copyable_v, "Type T must be device copyable."); - static_assert(is_property_list::value, - "Property list is invalid."); - static_assert(check_property_list::value, + static constexpr bool is_device_copyable = is_device_copyable_v; + static_assert(is_device_copyable, "Type T must be device copyable."); + static constexpr bool is_valid_property_list = + is_property_list::value; + static_assert(is_valid_property_list, "Property list is invalid."); + static constexpr bool contains_valid_properties = + check_property_list::value; + static_assert(contains_valid_properties, "The property list contains invalid property."); // check the set if FPGA specificed properties are used - static_assert(detail::checkValidFPGAPropertySet::value, + static constexpr bool hasValidFPGAProperties = + detail::checkValidFPGAPropertySet::value; + static_assert(hasValidFPGAProperties, "FPGA Interface properties (i.e. awidth, dwidth, etc.)" "can only be set with BufferLocation together."); @@ -218,11 +238,12 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { // `PropertyValueTs...` must have the same property value. template annotated_arg(const T &_obj, PropertyValueTs... props) noexcept : obj(_obj) { + static constexpr bool has_same_properties = std::is_same< + property_list_t, + detail::merged_properties_t>::value; static_assert( - std::is_same< - property_list_t, - detail::merged_properties_t>::value, + has_same_properties, "The property list must contain all properties of the input of the " "constructor"); } @@ -234,16 +255,19 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { template explicit annotated_arg(const annotated_arg &other) noexcept : obj(other.obj) { - static_assert(std::is_convertible::value, + static constexpr bool is_input_convertible = + std::is_convertible::value; + static_assert(is_input_convertible, "The underlying data type of the input annotated_arg is not " "compatible"); + static constexpr bool has_same_properties = std::is_same< + property_list_t, + detail::merged_properties_t>::value; static_assert( - std::is_same< - property_list_t, - detail::merged_properties_t>::value, - "The constructed annotated_arg type must contain all the properties of " - "the input annotated_arg"); + has_same_properties, + "The constructed annotated_arg type must contain all the properties " + "of the input annotated_arg"); } // Constructs an annotated_arg object from another annotated_arg object and a @@ -255,13 +279,17 @@ __SYCL_TYPE(annotated_arg) annotated_arg> { const PropertyListV &proplist) noexcept : obj(other.obj) { (void)proplist; - static_assert(std::is_convertible::value, + static constexpr bool is_input_convertible = + std::is_convertible::value; + static_assert(is_input_convertible, "The underlying data type of the input annotated_arg is not " "compatible"); + static constexpr bool has_same_properties = std::is_same< + property_list_t, + detail::merged_properties_t>::value; static_assert( - std::is_same>::value, + has_same_properties, "The property list of constructed annotated_arg type must be the union " "of the input property lists"); } From 6fafa9c1e8df8b915c0bbb0d026f2e97613f29cb Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Fri, 3 Nov 2023 11:07:48 -0700 Subject: [PATCH 18/28] Corrected cut-paste erros in doc. --- .../proposed/sycl_ext_intel_cache_controls.asciidoc | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc index 5c7ae735b1e29..e7cbc11c16f96 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -258,10 +258,9 @@ namespace oneapi { namespace experimental { template <> -struct is_property_key - : std::true_type {}; +struct is_property_key : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; } // namespace sycl::ext::oneapi::experimental @@ -334,9 +333,10 @@ namespace oneapi { namespace experimental { template <> -struct is_property_key : std::true_type {}; +struct is_property_key + : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; } // namespace sycl::ext::oneapi::experimental From 3c5720cd04627480b3d97e528b90b61ed0fe7696 Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Sat, 4 Nov 2023 14:26:11 -0700 Subject: [PATCH 19/28] General cleanup and enhanced a test. --- .../SYCLLowerIR/CompileTimePropertiesPass.cpp | 8 +- .../experimental/cache_control_properties.hpp | 84 +++++++++++++++++-- .../annotated_ptr/annotated_ptr.hpp | 6 ++ .../properties/properties_cache_control.cpp | 54 ++++++++---- .../properties_cache_control_errors.cpp | 8 ++ 5 files changed, 132 insertions(+), 28 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index cfb03c997dbd6..f1991a6038362 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -725,9 +725,11 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation( uint32_t CacheMode = 0; while (AttrVal) { // The attribute value encodes cache control and levels. - // Low-order to high-order nibbles represent the enumerated cache modes. - // In each nibble cache levels are encodes as L1->bit0, L2->bit1, - // L3->bit2 and L4->bit3. The SPIR-V encoding uses numbers 0..3. + // Low-order to high-order nibbles hold cache levels specified for the + // enumerated SYCL cache modes. Lowest order nibble for uncached, next + // for cached, and so on. + // In each nibble cache levels are encoded as L1=1, L2=2, L3=4 and L4=8. + // The SPIR-V encoding of cache levels L1..L4 uses values 0..3. uint32_t CacheLevel = 0; uint32_t LevelMask = AttrVal & 0xf; while (LevelMask) { diff --git a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp index d5a7002c43047..8ca303e6e7f33 100755 --- a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp @@ -120,6 +120,21 @@ struct is_property_key_of> { static constexpr const char *name = "sycl-cache-read-hint"; static constexpr const int value = ((checkReadHint() + ...), - checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), - (countL(Cs::levels, 4) + ...), - (countL(Cs::levels, 8) + ...)>(), + checkUnique<(countL(Cs::levels, L1BIT) + ...), + (countL(Cs::levels, L2BIT) + ...), + (countL(Cs::levels, L3BIT) + ...), + (countL(Cs::levels, L4BIT) + ...)>(), ((Cs::encoding) | ...)); }; @@ -186,9 +202,10 @@ struct PropertyMetaInfo< static constexpr const char *name = "sycl-cache-read-assertion"; static constexpr const int value = ((checkReadAssertion() + ...), - checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), - (countL(Cs::levels, 4) + ...), - (countL(Cs::levels, 8) + ...)>(), + checkUnique<(countL(Cs::levels, L1BIT) + ...), + (countL(Cs::levels, L2BIT) + ...), + (countL(Cs::levels, L3BIT) + ...), + (countL(Cs::levels, L4BIT) + ...)>(), ((Cs::encoding) | ...)); }; @@ -203,12 +220,61 @@ struct PropertyMetaInfo> { static constexpr const char *name = "sycl-cache-write-hint"; static constexpr const int value = ((checkWriteHint() + ...), - checkUnique<(countL(Cs::levels, 1) + ...), (countL(Cs::levels, 2) + ...), - (countL(Cs::levels, 4) + ...), - (countL(Cs::levels, 8) + ...)>(), + checkUnique<(countL(Cs::levels, L1BIT) + ...), + (countL(Cs::levels, L2BIT) + ...), + (countL(Cs::levels, L3BIT) + ...), + (countL(Cs::levels, L4BIT) + ...)>(), ((Cs::encoding) | ...)); }; +// Check read_hint and read_assertion cache levels across all properties. + +class levelCounts { +public: + int countL1; + int countL2; + int countL3; + int countL4; +}; + +static constexpr levelCounts operator+(levelCounts L1, levelCounts L2) { + return {L1.countL1 + L2.countL1, L1.countL2 + L2.countL2, + L1.countL3 + L2.countL3, L1.countL4 + L2.countL4}; +} + +// Gather all levels specified in a property into one 4-bit mask. Then, count +// how many times each level is specified. +template constexpr static levelCounts allLevels() { + constexpr const int levelUsed = + (encoding | (encoding >> CACHED) | (encoding >> STREAMING) | + (encoding >> INVALIDATE) | (encoding >> CONSTANT) | (encoding >> WT) | + (encoding >> WB)); + return {(levelUsed & L1BIT) != 0, (levelUsed & L2BIT) != 0, + (levelUsed & L3BIT) != 0, (levelUsed & L4BIT) != 0}; +} + +// Compare strings at compile time +constexpr bool compareStrs(const char *Str1, const char *Str2) { + return std::string_view(Str1) == Str2; +} + +// Check that the number of times a particular cache level is specified in +// read_hint and read_assertion properties is at most 1. +template struct checkValidCacheControlProperties { + static constexpr const levelCounts allZeros{0, 0, 0, 0}; + static constexpr levelCounts lCounts = + (((compareStrs(detail::PropertyMetaInfo::name, + "sycl-cache-read-assertion") || + compareStrs(detail::PropertyMetaInfo::name, + "sycl-cache-read-hint")) + ? allLevels::value>() + : allZeros) + + ...); + static constexpr bool value = + sizeof...(Cs) == 1 || (lCounts.countL1 < 2 && lCounts.countL2 < 2 && + lCounts.countL3 < 2 && lCounts.countL4 < 2); +}; + } // namespace detail template diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp index 193dacba77cae..e51dfe1b98837 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp @@ -222,6 +222,12 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr> { "FPGA Interface properties (i.e. awidth, dwidth, etc.)" "can only be set with BufferLocation together."); + // Check that Read cache controls do not conflict at any cache level. + static_assert( + detail::checkValidCacheControlProperties::value, + "Specify either read_hint or read_assertion at a cache level, but not " + "both"); + annotated_ptr() noexcept = default; annotated_ptr(const annotated_ptr &) = default; annotated_ptr &operator=(const annotated_ptr &) = default; diff --git a/sycl/test/extensions/properties/properties_cache_control.cpp b/sycl/test/extensions/properties/properties_cache_control.cpp index 4e1b8b03e7f28..273079334036e 100755 --- a/sycl/test/extensions/properties/properties_cache_control.cpp +++ b/sycl/test/extensions/properties/properties_cache_control.cpp @@ -1,5 +1,4 @@ -// RUN: %clangxx -D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ \ -// RUN: -fsycl-device-only -S -Xclang -emit-llvm %s -o - | \ +// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | \ // RUN: FileCheck %s --check-prefix CHECK-IR #include @@ -14,7 +13,6 @@ using load_hint = annotated_ptr< read_hint, cache_control>))>; - using load_assertion = annotated_ptr< int, decltype(properties( @@ -22,7 +20,6 @@ using load_assertion = annotated_ptr< read_assertion, cache_control>))>; - using store_hint = annotated_ptr< float, decltype(properties( @@ -30,15 +27,21 @@ using store_hint = annotated_ptr< cache_control, cache_control>))>; +using load_store_hint = annotated_ptr< + float, + decltype(properties( + read_hint>, + read_assertion>, + write_hint< + cache_control>))>; void cache_control_read_hint_func() { queue q; constexpr int N = 10; float *ArrayA = malloc_shared(N, q); q.submit([&](handler &cgh) { - cgh.parallel_for<>(range<1>(N), [=](item<1> item) { - auto item_id = item.get_linear_id(); - load_hint src{&ArrayA[item_id]}; + cgh.single_task<>([=]() { + load_hint src{&ArrayA[0]}; *src = 55.0f; }); }); @@ -49,9 +52,8 @@ void cache_control_read_assertion_func() { constexpr int N = 10; int *ArrayA = malloc_shared(N, q); q.submit([&](handler &cgh) { - cgh.parallel_for<>(range<1>(N), [=](item<1> item) { - auto item_id = item.get_linear_id(); - load_assertion src{&ArrayA[item_id]}; + cgh.single_task<>([=]() { + load_assertion src{&ArrayA[0]}; *src = 66; }); }); @@ -62,24 +64,39 @@ void cache_control_write_hint_func() { constexpr int N = 10; float *ArrayA = malloc_shared(N, q); q.submit([&](handler &cgh) { - cgh.parallel_for<>(range<1>(N), [=](item<1> item) { - auto item_id = item.get_linear_id(); - store_hint dst{&ArrayA[item_id]}; + cgh.single_task<>([=]() { + store_hint dst{&ArrayA[0]}; + *dst = 77.0f; + }); + }); +} + +void cache_control_read_write_func() { + queue q; + constexpr int N = 10; + float *ArrayA = malloc_shared(N, q); + q.submit([&](handler &cgh) { + cgh.single_task<>([=]() { + load_store_hint dst{&ArrayA[0]}; *dst = 77.0f; }); }); } // CHECK-IR: spir_kernel{{.*}}cache_control_read_hint_func -// CHECK-IR: {{.*}}getelementptr inbounds float{{.*}}!spirv.Decorations [[RHINT:.*]] +// CHECK-IR: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RHINT:.*]] // CHECK-IR: ret void // CHECK-IR: spir_kernel{{.*}}cache_control_read_assertion_func -// CHECK-IR: {{.*}}getelementptr inbounds i32{{.*}}!spirv.Decorations [[RASSERT:.*]] +// CHECK-IR: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RASSERT:.*]] // CHECK-IR: ret void // CHECK-IR: spir_kernel{{.*}}cache_control_write_hint_func -// CHECK-IR: {{.*}}getelementptr inbounds float{{.*}}!spirv.Decorations [[WHINT:.*]] +// CHECK-IR: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[WHINT:.*]] +// CHECK-IR: ret void + +// CHECK-IR: spir_kernel{{.*}}cache_control_read_write_func +// CHECK-IR: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RWHINT:.*]] // CHECK-IR: ret void // CHECK-IR: [[RHINT]] = !{[[RHINT1:.*]], [[RHINT2:.*]], [[RHINT3:.*]]} @@ -97,3 +114,8 @@ void cache_control_write_hint_func() { // CHECK-IR: [[WHINT2]] = !{i32 6443, i32 0, i32 1} // CHECK-IR: [[WHINT3]] = !{i32 6443, i32 1, i32 2} // CHECK-IR: [[WHINT4]] = !{i32 6443, i32 2, i32 2} + +// CHECK-IR: [[RWHINT]] = !{[[RWHINT1:.*]], [[RWHINT2:.*]], [[RWHINT3:.*]]} +// CHECK-IR: [[RWHINT1]] = !{i32 6442, i32 2, i32 1} +// CHECK-IR: [[RWHINT2]] = !{i32 6442, i32 3, i32 4} +// CHECK-IR: [[RWHINT3]] = !{i32 6443, i32 3, i32 1} diff --git a/sycl/test/extensions/properties/properties_cache_control_errors.cpp b/sycl/test/extensions/properties/properties_cache_control_errors.cpp index b3de0a503f7f5..c892789598151 100755 --- a/sycl/test/extensions/properties/properties_cache_control_errors.cpp +++ b/sycl/test/extensions/properties/properties_cache_control_errors.cpp @@ -29,6 +29,11 @@ using annotated_ptr_load5 = annotated_ptr< float, decltype(properties( write_hint>))>; +using annotated_ptr_load6 = annotated_ptr< + float, + decltype(properties( + read_hint>, + read_assertion>))>; void cache_control_read_func(queue q) { float *ArrayA = malloc_shared(10, q); @@ -48,6 +53,9 @@ void cache_control_read_func(queue q) { // expected-error@sycl/ext/intel/experimental/cache_control_properties.hpp:* {{write_hint must specify cache_mode uncached, write_through, write_back or streaming}} annotated_ptr_load5 src5{&ArrayA[0]}; + + // expected-error@sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp:* {{Specify either read_hint or read_assertion at a cache level, but not both}} + annotated_ptr_load6 src6{ &ArrayA[0] }; }); }); } From d745c62bbc5739c8ae2d051f85b09904b156628e Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Sat, 4 Nov 2023 14:47:30 -0700 Subject: [PATCH 20/28] Formatting change. --- .../extensions/properties/properties_cache_control_errors.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/extensions/properties/properties_cache_control_errors.cpp b/sycl/test/extensions/properties/properties_cache_control_errors.cpp index c892789598151..a23b535a6d963 100755 --- a/sycl/test/extensions/properties/properties_cache_control_errors.cpp +++ b/sycl/test/extensions/properties/properties_cache_control_errors.cpp @@ -55,7 +55,7 @@ void cache_control_read_func(queue q) { annotated_ptr_load5 src5{&ArrayA[0]}; // expected-error@sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp:* {{Specify either read_hint or read_assertion at a cache level, but not both}} - annotated_ptr_load6 src6{ &ArrayA[0] }; + annotated_ptr_load6 src6{&ArrayA[0]}; }); }); } From ea7e677afb338abaae88881af389b8ef06664fa2 Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Sat, 4 Nov 2023 17:10:26 -0700 Subject: [PATCH 21/28] Test update. --- .../experimental/cache_control_properties.hpp | 48 ------------------- .../annotated_ptr/annotated_ptr.hpp | 12 ----- .../properties_cache_control_errors.cpp | 8 ---- 3 files changed, 68 deletions(-) diff --git a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp index 8ca303e6e7f33..a85ecd1a45737 100755 --- a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp @@ -227,54 +227,6 @@ struct PropertyMetaInfo> { ((Cs::encoding) | ...)); }; -// Check read_hint and read_assertion cache levels across all properties. - -class levelCounts { -public: - int countL1; - int countL2; - int countL3; - int countL4; -}; - -static constexpr levelCounts operator+(levelCounts L1, levelCounts L2) { - return {L1.countL1 + L2.countL1, L1.countL2 + L2.countL2, - L1.countL3 + L2.countL3, L1.countL4 + L2.countL4}; -} - -// Gather all levels specified in a property into one 4-bit mask. Then, count -// how many times each level is specified. -template constexpr static levelCounts allLevels() { - constexpr const int levelUsed = - (encoding | (encoding >> CACHED) | (encoding >> STREAMING) | - (encoding >> INVALIDATE) | (encoding >> CONSTANT) | (encoding >> WT) | - (encoding >> WB)); - return {(levelUsed & L1BIT) != 0, (levelUsed & L2BIT) != 0, - (levelUsed & L3BIT) != 0, (levelUsed & L4BIT) != 0}; -} - -// Compare strings at compile time -constexpr bool compareStrs(const char *Str1, const char *Str2) { - return std::string_view(Str1) == Str2; -} - -// Check that the number of times a particular cache level is specified in -// read_hint and read_assertion properties is at most 1. -template struct checkValidCacheControlProperties { - static constexpr const levelCounts allZeros{0, 0, 0, 0}; - static constexpr levelCounts lCounts = - (((compareStrs(detail::PropertyMetaInfo::name, - "sycl-cache-read-assertion") || - compareStrs(detail::PropertyMetaInfo::name, - "sycl-cache-read-hint")) - ? allLevels::value>() - : allZeros) + - ...); - static constexpr bool value = - sizeof...(Cs) == 1 || (lCounts.countL1 < 2 && lCounts.countL2 < 2 && - lCounts.countL3 < 2 && lCounts.countL4 < 2); -}; - } // namespace detail template diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp index 1a8b1db81f9cd..193dacba77cae 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp @@ -221,18 +221,6 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr> { static_assert(hasValidFPGAProperties, "FPGA Interface properties (i.e. awidth, dwidth, etc.)" "can only be set with BufferLocation together."); - // check if conduit and register_map properties are specified together - static constexpr bool hasConduitAndRegisterMapProperties = - detail::checkHasConduitAndRegisterMap::value; - static_assert(hasConduitAndRegisterMapProperties, - "The properties conduit and register_map cannot be " - "specified at the same time."); - - // Check that Read cache controls do not conflict at any cache level. - static_assert( - detail::checkValidCacheControlProperties::value, - "Specify either read_hint or read_assertion at a cache level, but not " - "both"); annotated_ptr() noexcept = default; annotated_ptr(const annotated_ptr &) = default; diff --git a/sycl/test/extensions/properties/properties_cache_control_errors.cpp b/sycl/test/extensions/properties/properties_cache_control_errors.cpp index a23b535a6d963..b3de0a503f7f5 100755 --- a/sycl/test/extensions/properties/properties_cache_control_errors.cpp +++ b/sycl/test/extensions/properties/properties_cache_control_errors.cpp @@ -29,11 +29,6 @@ using annotated_ptr_load5 = annotated_ptr< float, decltype(properties( write_hint>))>; -using annotated_ptr_load6 = annotated_ptr< - float, - decltype(properties( - read_hint>, - read_assertion>))>; void cache_control_read_func(queue q) { float *ArrayA = malloc_shared(10, q); @@ -53,9 +48,6 @@ void cache_control_read_func(queue q) { // expected-error@sycl/ext/intel/experimental/cache_control_properties.hpp:* {{write_hint must specify cache_mode uncached, write_through, write_back or streaming}} annotated_ptr_load5 src5{&ArrayA[0]}; - - // expected-error@sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp:* {{Specify either read_hint or read_assertion at a cache level, but not both}} - annotated_ptr_load6 src6{&ArrayA[0]}; }); }); } From 5ef494d1f85018ba02be0db195084650993d00bd Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Sun, 5 Nov 2023 12:00:18 -0800 Subject: [PATCH 22/28] Fixed merge error. --- .../oneapi/experimental/annotated_ptr/annotated_ptr.hpp | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp index 193dacba77cae..d0d9f33a5459b 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp @@ -160,6 +160,7 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr> { using property_list_t = detail::properties_t; // buffer_location and alignment are allowed for annotated_ref + // Cache controls are allowed for annotated_ptr using allowed_properties = std::tuple), decltype(ext::oneapi::experimental::alignment<0>), @@ -221,6 +222,12 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr> { static_assert(hasValidFPGAProperties, "FPGA Interface properties (i.e. awidth, dwidth, etc.)" "can only be set with BufferLocation together."); + // check if conduit and register_map properties are specified together + static constexpr bool hasConduitAndRegisterMapProperties = + detail::checkHasConduitAndRegisterMap::value; + static_assert(hasConduitAndRegisterMapProperties, + "The properties conduit and register_map cannot be " + "specified at the same time."); annotated_ptr() noexcept = default; annotated_ptr(const annotated_ptr &) = default; From fc87598e84dd1ce03111661a181dcb61b2852fe5 Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Mon, 6 Nov 2023 10:21:33 -0800 Subject: [PATCH 23/28] Removed some unused declarations. --- .../intel/experimental/cache_control_properties.hpp | 12 +++--------- 1 file changed, 3 insertions(+), 9 deletions(-) diff --git a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp index a85ecd1a45737..b736c169f5b37 100755 --- a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp @@ -57,6 +57,9 @@ template struct cache_control { static constexpr const int countL3 = ((Ls == cache_level::L3 ? 1 : 0) + ...); static constexpr const int countL4 = ((Ls == cache_level::L4 ? 1 : 0) + ...); static constexpr const uint32_t levels = ((1 << static_cast(Ls)) | ...); + // Starting bit position for cache levels of a cache mode are uncached=0, + // cached=4, streaming=8, invalidate=12, constant=16, write_through=20 and + // write_back=24. The shift value is computed as cache_mode * 4. static constexpr const uint32_t encoding = (countL1, countL2, countL3, countL4, detail::checkLevel1(), detail::checkLevel2(), detail::checkLevel3(), @@ -126,15 +129,6 @@ static constexpr int L2BIT = 2; static constexpr int L3BIT = 4; static constexpr int L4BIT = 8; -// Starting bit position for cache levels of a cache mode. -static constexpr int UNCACHED = 0; -static constexpr int CACHED = 4; -static constexpr int STREAMING = 8; -static constexpr int INVALIDATE = 12; -static constexpr int CONSTANT = 16; -static constexpr int WT = 20; -static constexpr int WB = 24; - static constexpr int countL(int levels, int mask) { return levels & mask ? 1 : 0; } From abf4bcf3cfba17463f7cd091b05451c9881884f9 Mon Sep 17 00:00:00 2001 From: rdeodhar Date: Mon, 6 Nov 2023 20:50:42 -0700 Subject: [PATCH 24/28] Update sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc Co-authored-by: Alexey Bader --- .../extensions/proposed/sycl_ext_intel_cache_controls.asciidoc | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc index e7cbc11c16f96..e0893e45f31f1 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -117,6 +117,7 @@ to contain multiple `read_hint` properties. In order to specify multiple The same rule applies to `write_hint` and `read_assertion`. ==== Cache modes + The `cache_control` structure is used by each of the three properties to specify a cache mode and a list of cache levels to which it applies. From b9ae23c5d25508aef8832eeb83a4c2ef916d074d Mon Sep 17 00:00:00 2001 From: rdeodhar Date: Mon, 6 Nov 2023 20:51:08 -0700 Subject: [PATCH 25/28] Update sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc Co-authored-by: Alexey Bader --- .../extensions/proposed/sycl_ext_intel_cache_controls.asciidoc | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc index e0893e45f31f1..301587cbc7400 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -236,6 +236,7 @@ property. -- ==== Write hint property + This property is a hint requesting specific cache behavior when storing to memory through an `annotated_ptr`. This property can affect the performance of device code, but it From 904190affaa859e0f7c8a46d1091a3ccb94e83e8 Mon Sep 17 00:00:00 2001 From: rdeodhar Date: Mon, 6 Nov 2023 20:51:51 -0700 Subject: [PATCH 26/28] Update sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc Co-authored-by: Alexey Bader --- .../extensions/proposed/sycl_ext_intel_cache_controls.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc index 301587cbc7400..d0502199e102e 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -310,8 +310,8 @@ via a `streaming` cache control. |==== -- - ==== Read assertion property + This property is an assertion by the application, promising that the application accesses memory in a certain way. Care must be taken when using this property because it can From 11e7887e4c06fc052e24fd45e78261f3ae93dacd Mon Sep 17 00:00:00 2001 From: rdeodhar Date: Mon, 6 Nov 2023 20:58:21 -0700 Subject: [PATCH 27/28] Update sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc Co-authored-by: Alexey Bader --- .../extensions/proposed/sycl_ext_intel_cache_controls.asciidoc | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc index d0502199e102e..9de9df84051da 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc @@ -172,6 +172,7 @@ write_hint> ``` ==== Read hint property + This property is a hint requesting specific cache behavior when loading from memory through an `annotated_ptr`. This property can affect the performance of device code, but it does not change From 694676bd5c4eee7129dedab8a7aa6c2015396476 Mon Sep 17 00:00:00 2001 From: Rajiv Deodhar Date: Mon, 6 Nov 2023 20:09:15 -0800 Subject: [PATCH 28/28] Added comments noting need for consistency in cache_mode definition. --- llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp | 3 ++- .../sycl/ext/intel/experimental/cache_control_properties.hpp | 2 ++ 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index f1991a6038362..8c3707852456f 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -130,7 +130,8 @@ MDNode *buildSpirvDecorCacheProp(LLVMContext &Ctx, StringRef Name, write_back = 2, write_streaming = 3 }; - // SYCL encodings of read/write control + // SYCL encodings of read/write control. Definition of cache_mode should match + // definition in SYCL header file cache_control_properties.hpp. enum class cache_mode { uncached, cached, diff --git a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp index b736c169f5b37..58103d39a17a0 100755 --- a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp @@ -22,6 +22,8 @@ namespace ext { namespace intel { namespace experimental { +// SYCL encodings of read/write control. Definition of cache_mode should match +// definition in file CompileTimePropertiesPass.cpp. enum class cache_mode { uncached, cached,