address comments and disable dynamic tests on gen12 (not locally reproducible)

Author: Victor Lomuller

sycl/doc/extensions/experimental/sycl_ext_oneapi_work_group_scratch_memory.asciidoc

@@ -36,7 +36,7 @@ https://github.com/intel/llvm/issues
 
 == Dependencies
 
-This extension is written against the SYCL 2020 revision 8 specification.  All
+This extension is written against the SYCL 2020 revision 9 specification.  All
 references below to the "core SYCL specification" or to section numbers in the
 SYCL specification refer to that revision.
 
@@ -49,11 +49,12 @@ The following extensions are required:
 
 == Status
 
-This is an experimental extension specification, intended to gather community
-feedback.  Interfaces defined in this specification may not be implemented yet
-or may be in a preliminary state.  The specification itself may also change in
-incompatible ways before it is finalized.  *Shipping software products should
-not rely on APIs defined in this specification.*
+This is an experimental extension specification, intended to provide early
+access to features and gather community feedback.  Interfaces defined in this
+specification are implemented in {dpcpp}, but they are not finalized and may
+change incompatibly in future versions of {dpcpp} without prior notice.
+*Shipping software products should not rely on APIs defined in this
+specification.*
 
 
 == Overview
@@ -71,7 +72,7 @@ The behavior is similar to the usage of unbounded array with the CUDA `+__shared
 
 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_ONEAPI_WORK_GROUP_DYNAMIC` to one of the values defined in the
+macro `SYCL_EXT_ONEAPI_WORK_GROUP_SCRATCH_MEMORY` 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
@@ -95,9 +96,10 @@ to a dynamically allocated buffer in the device local memory.
 
 [source,c++]
 ----
-void* get_work_group_scratch_memory()
+namespace sycl::ext::oneapi::experimental {
+  void* get_work_group_scratch_memory();
+}
 ----
-_Constraints_: `T` must be trivially constructible and trivially destructible.
 
 _Returns_: A pointer to a dynamically allocated buffer
            in the device local memory.
@@ -141,29 +143,27 @@ device local memory required by the kernel in bytes.
 
 === Total allocation check
 
-If the total amount of device local memory requested (i.e., the sum of
-all memory requested by `local_accessor`, `group_local_memory`,
-`group_local_memory_for_overwrite`, `work_group_static` and `work_group_scratch_size`) exceeds a device's
-local memory capacity (as reported by `local_mem_size`) then the implementation
-must throw a synchronous `exception` with the `errc::memory_allocation` error
-code from the kernel invocation command (e.g. `parallel_for`).
-
-==== Usage examples
+If the total amount of device local memory requested exceeds a device's
+local memory capacity as reported by `info::device::local_mem_size`
+then the implementation must throw a synchronous exception with the
+`errc::memory_allocation` error code from the kernel invocation command
+(e.g. `parallel_for`). This check must take all APIs that allocation device
+local memory into account, whether via the `work_group_scratch_size` property
+or other APIs such as `local_accessor`.
 
-===== Allocations with size unknown at compile-time
+=== Example
 
 [source,c++]
 ----
-using namespace syclex = sycl::ext::oneapi::experimental;
-
+namespace syclex = sycl::ext::oneapi::experimental;
 
 ...
 
 q.parallel_for(sycl::nd_range<1>{N, M},
   syclex::properties{syclex::work_group_scratch_size(M * sizeof(int))},
   [=](sycl::nd_item<1> it) {
- auto ptr = get_work_group_scratch_memory();
- auto ptr2 = get_work_group_scratch_memory();
+ auto ptr = syclex::get_work_group_scratch_memory();
+ auto ptr2 = syclex::get_work_group_scratch_memory();
 });
 ----
 

sycl/doc/extensions/experimental/sycl_ext_oneapi_work_group_static.asciidoc

@@ -20,7 +20,7 @@
 == Notice
 
 [%hardbreaks]
-Copyright (C) 2024 Intel Corporation.  All rights reserved.
+Copyright (C) 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
@@ -36,23 +36,23 @@ https://github.com/intel/llvm/issues
 
 == Dependencies
 
-This extension is written against the SYCL 2020 revision 8 specification.  All
+This extension is written against the SYCL 2020 revision 9 specification.  All
 references below to the "core SYCL specification" or to section numbers in the
 SYCL specification refer to that revision.
 
 
 == Status
 
-This is an experimental extension specification, intended to gather community
-feedback.  Interfaces defined in this specification may not be implemented yet
-or may be in a preliminary state.  The specification itself may also change in
-incompatible ways before it is finalized.  *Shipping software products should
-not rely on APIs defined in this specification.*
-
+This is an experimental extension specification, intended to provide early
+access to features and gather community feedback.  Interfaces defined in this
+specification are implemented in {dpcpp}, but they are not finalized and may
+change incompatibly in future versions of {dpcpp} without prior notice.
+*Shipping software products should not rely on APIs defined in this
+specification.*
 
 == Overview
 
-This extension adds a ways to allocate device local memory, without passing a
+This extension adds a way to allocate device local memory, without passing a
 kernel argument: `work_group_static`.
 Device local memory is memory that is shared by all work-items in a work-group.
 The behavior is similar to the CUDA `+__shared__+` keyword, and the extension
@@ -111,15 +111,12 @@ public:
   const work_group_static& operator=(const T& value) const noexcept;
 
   T* operator&() const noexcept;
-
-private:
-  T storage;
 };
 
 } // namespace sycl::ext::oneapi::experimental
 ----
 
-`T` must be trivially constructible and trivially destructible.
+`T` must be a cv-unqualified trivially constructible and trivially destructible.
 
 The storage for the object is allocated in device local memory before
 calling the user's kernel lambda, and deallocated when all work-items
@@ -164,7 +161,7 @@ work_group_static& operator=(const T& value) noexcept;
 ----
 _Constraints_: Available only if `std::is_array_v<T>` is false.
 
-_Effects_: Replaces the value referenced by `*ptr` with `value`.
+_Effects_: Copy `value` into the storage of the `work_group_static` instance.
 
 _Returns_: A reference to this instance of `work_group_static`.
 
@@ -173,7 +170,7 @@ _Returns_: A reference to this instance of `work_group_static`.
 T* operator&() noexcept;
 ----
 _Returns_: A pointer to the device local memory associated with this
-instance of `work_group_static` (i.e., `ptr`).
+instance of `work_group_static`.
 
 ==== Interaction with common address space deduction rules
 
@@ -183,33 +180,39 @@ the local address space.
 
 === Total allocation check
 
-If the total amount of device local memory requested (i.e., the sum of
-all memory requested by `local_accessor`, `group_local_memory`,
-`group_local_memory_for_overwrite`, `work_group_static` and `work_group_scratch_size`) exceeds a device's
-local memory capacity (as reported by `local_mem_size`) then the implementation
-must throw a synchronous `exception` with the `errc::memory_allocation` error
-code from the kernel invocation command (e.g. `parallel_for`).
-
-==== Usage examples
+If the total amount of device local memory requested exceeds a device's
+local memory capacity as reported by `info::device::local_mem_size`
+then the implementation must throw a synchronous exception with the
+`errc::memory_allocation` error code from the kernel invocation command
+(e.g. `parallel_for`). This check must take all APIs that allocation device
+local memory into account, whether via the `work_group_scratch_size` property
+or other APIs such as `local_accessor`.
 
-===== Allocations with size known at compile-time
+=== Example
 
 [source,c++]
 ----
-using namespace syclex = sycl::ext::oneapi::experimental;
+namespace syclex = sycl::ext::oneapi::experimental;
 
 /* optional: static */ syclex::work_group_static<int> program_scope_scalar;
 /* optional: static */ syclex::work_group_static<int[16]> program_scope_array;
 
+class ClassScope {
+  static syclex::work_group_static<int> class_scope_scalar;
+};
+
+syclex::work_group_static<int> ClassScope::class_scope_scalar;
+
 void foo() {
   /* optional: static */ syclex::work_group_static<int> function_scope_scalar;
   function_scope_scalar = 1; // assignment via overloaded = operator
   function_scope_scalar += 2; // += operator via implicit conversion to int&
+  class_scope_scalar = 3;
   int* ptr = &function_scope_scalar; // conversion to pointer via overloaded & operator
 }
 
 void bar() {
-  /* optional: static */ sylex::work_group_static<int[64]> function_scope_array;
+  /* optional: static */ syclex::work_group_static<int[64]> function_scope_array;
   function_scope_array[0] = 1; // [] operator via implicit conversion to int(&)[64]
   int* ptr = function_scope_array; // conversion to pointer via implicit conversion to int(&)[64]
 }

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

@@ -59,7 +59,8 @@ enum class UnsupportedGraphFeatures {
   sycl_ext_oneapi_device_global = 6,
   sycl_ext_oneapi_bindless_images = 7,
   sycl_ext_oneapi_experimental_cuda_cluster_launch = 8,
-  sycl_ext_codeplay_enqueue_native_command = 9
+  sycl_ext_codeplay_enqueue_native_command = 9,
+  sycl_ext_oneapi_work_group_scratch_memory = 10
 };
 
 inline const char *

sycl/include/sycl/ext/oneapi/work_group_scratch_memory.hpp

@@ -23,7 +23,7 @@ inline void* get_work_group_scratch_memory() {
 #else
   throw sycl::exception(
       sycl::errc::feature_not_supported,
-      "sycl_ext_oneapi_work_group_dynamic extension is not supported on host");
+      "sycl_ext_oneapi_work_scratch_memory extension is not supported on host");
 #endif
 }
 

sycl/include/sycl/ext/oneapi/work_group_static.hpp

@@ -10,7 +10,7 @@
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL_ALWAYS_INLINE
 #include <sycl/exception.hpp>                 // for exception
 
-#include <type_traits> // for enable_if_t
+#include <type_traits> // for enable_if_t, is_trivially_destructible_v ...
 
 namespace sycl {
 inline namespace _V1 {
@@ -34,6 +34,8 @@ template <typename T> class __SYCL_WG_SCOPE work_group_static final {
       std::is_trivially_destructible_v<T> &&
           std::is_trivially_constructible_v<T>,
       "Can only be used with trivially constructible and destructible types");
+  static_assert(!std::is_const_v<T> || !std::is_volatile_v<T>,
+                "Can only be used with non const and non volatile types");
   __SYCL_ALWAYS_INLINE work_group_static() = default;
   work_group_static(const work_group_static &) = delete;
   work_group_static &operator=(const work_group_static &) = delete;

sycl/source/feature_test.hpp.in

@@ -109,7 +109,7 @@ inline namespace _V1 {
 #define SYCL_EXT_ONEAPI_PROFILING_TAG 1
 #define SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND 1
 #define SYCL_EXT_ONEAPI_GET_KERNEL_INFO 1
-#define SYCL_EXT_ONEAPI_WORK_GROUP_DYNAMIC 1
+#define SYCL_EXT_ONEAPI_WORK_GROUP_SCRATCH_MEMORY 1
 #define SYCL_EXT_ONEAPI_WORK_GROUP_STATIC 1
 // In progress yet
 #define SYCL_EXT_ONEAPI_ATOMIC16 0

sycl/source/handler.cpp

@@ -1951,6 +1951,8 @@ void handler::setKernelClusterLaunch(sycl::range<3> ClusterSize, int Dims) {
 }
 
 void handler::setKernelWorkGroupMem(size_t Size) {
+  throwIfGraphAssociated<syclex::detail::UnsupportedGraphFeatures::
+                             sycl_ext_oneapi_work_group_scratch_memory>();
   impl->MKernelWorkGroupMemorySize = Size;
 }
 

sycl/test-e2e/Graph/Inputs/work_group_dynamic_memory.cpp

@@ -2,6 +2,8 @@
 // RUN: %{run} %t.out
 //
 
+// UNSUPPORTED: gpu-intel-gen12
+
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/work_group_scratch_memory.hpp>
 #include <sycl/group_barrier.hpp>

sycl/test-e2e/WorkGroupMemory/Dynamic/copy_dynamic_size.cpp

@@ -2,6 +2,8 @@
 // RUN: %{run} %t.out
 //
 
+// UNSUPPORTED: gpu-intel-gen12
+
 // Test work_group_dynamic extension with allocation size specified at runtime
 // and an additional local accessor.