intel · dm-vodopyanov · May 12, 2026 · Apr 10, 2026 · Apr 14, 2026 · Apr 14, 2026
@@ -112,6 +112,27 @@ supports.
  feature-test macro always has this value.
 |===
 
+=== Headers
+
+The APIs defined in this extension are provided by either of the following header files:
+
+* `<sycl/sycl.hpp>`
+* `<sycl/ext/oneapi/kernel_properties.hpp>`
+
+In addition, the following lightweight header provides a subset of the APIs from this extension:
+
+* `<sycl/ext/oneapi/free_function_kernel_properties.hpp>`
+
+This lightweight header provides `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY`, and the
+properties `nd_range_kernel<Dims>` and `single_task_kernel`.  See
+link:../experimental/sycl_ext_oneapi_kernel_properties.asciidoc[
+sycl_ext_oneapi_kernel_properties] for other APIs that are provided by this
+header.
+
+[_Note:_ The lightweight header is intended for cases where fast compilation 
+time is a priority.
+_{endnote}_]
+
 === Defining a free function kernel
 
 A free function kernel is a normal C++ function definition, where the function
@@ -804,6 +825,8 @@ sycl_ext_oneapi_kernel_properties] by applying the properties to the function
 declaration as illustrated below.
 
 ```
+#include <sycl/ext/oneapi/kernel_properties.hpp>
+
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size<32>))
 void iota(float start, float *ptr) {

@@ -9,6 +9,7 @@
 :encoding: utf-8
 :lang: en
 :dpcpp: pass:[DPC++]
+:endnote: &#8212;{nbsp}end{nbsp}note
 
 :blank: pass:[ +]
 
@@ -115,6 +116,28 @@ supports.
  feature-test macro always has this value.
 |===
 
+=== Headers
+
+The APIs defined in this extension are provided by either of the following header files:
+
+* `<sycl/sycl.hpp>`
+* `<sycl/ext/oneapi/kernel_properties.hpp>`
+
+In addition, the following lightweight header provides a subset of the APIs that
+is suitable for applications that define kernels using the
+link:../experimental/sycl_ext_oneapi_free_function_kernels.asciidoc[
+sycl_ext_oneapi_free_function_kernels] extension:
+
+* `<sycl/ext/oneapi/free_function_kernel_properties.hpp>`
+
+This lightweight header provides `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY`, 
+and the properties `work_group_size`, `work_group_size_hint`, 
+`sub_group_size`, `max_work_group_size`, and `max_linear_work_group_size`.
+
+[_Note:_ The lightweight header is intended for cases where 
+kernels are defined using the free-function kernel syntax and fast 
+compilation time is a priority._{endnote}_]
+
 === Kernel Properties
 
 The kernel properties below correspond to kernel attributes defined in

@@ -19,7 +19,7 @@
 #include <sycl/ext/oneapi/experimental/graph.hpp>
 #include <sycl/ext/oneapi/experimental/use_root_sync_prop.hpp>
 #include <sycl/ext/oneapi/experimental/virtual_functions.hpp>
-#include <sycl/ext/oneapi/kernel_properties/properties.hpp>
+#include <sycl/ext/oneapi/kernel_properties.hpp>
 #include <sycl/ext/oneapi/work_group_scratch_memory.hpp>
 #include <sycl/kernel_handler.hpp>
 
@@ -361,6 +361,48 @@ using KernelPropertyHolderStructTy =
                 sycl::ext::oneapi::experimental::cuda::cluster_size_key<2>,
                 sycl::ext::oneapi::experimental::cuda::cluster_size_key<3>>;
 
+template <typename PropertiesT> constexpr void validateKernelProperties() {
+  using namespace sycl::ext::oneapi::experimental;
+
+  if constexpr (PropertiesT::template has_property<work_group_size_key>()) {
+    constexpr auto WGSize =
+        PropertiesT::template get_property<work_group_size_key>();
+
+    if constexpr (PropertiesT::template has_property<
+                      max_work_group_size_key>()) {
+      constexpr auto MaxWGSize =
+          PropertiesT::template get_property<max_work_group_size_key>();
+      constexpr auto WGDimensions = decltype(WGSize)::dimensions;
+      constexpr auto MaxWGDimensions = decltype(MaxWGSize)::dimensions;
+
+      static_assert(
+          WGDimensions == MaxWGDimensions,
+          "work_group_size and max_work_group_size dimensionality must match");
+      if constexpr (WGDimensions == MaxWGDimensions) {
+        static_assert(WGDimensions < 1 || WGSize[0] <= MaxWGSize[0],
+                      "work_group_size must not exceed max_work_group_size");
+        static_assert(WGDimensions < 2 || WGSize[1] <= MaxWGSize[1],
+                      "work_group_size must not exceed max_work_group_size");
+        static_assert(WGDimensions < 3 || WGSize[2] <= MaxWGSize[2],
+                      "work_group_size must not exceed max_work_group_size");
+      }
+    }
+
+    if constexpr (PropertiesT::template has_property<
+                      max_linear_work_group_size_key>()) {
+      constexpr auto Dimensions = decltype(WGSize)::dimensions;
+      constexpr auto LinearSize = WGSize[0] * (Dimensions > 1 ? WGSize[1] : 1) *
+                                  (Dimensions > 2 ? WGSize[2] : 1);
+      constexpr auto MaxLinearWGSize =
+          PropertiesT::template get_property<max_linear_work_group_size_key>();
+
+      static_assert(
+          LinearSize < MaxLinearWGSize.value,
+          "work_group_size must not exceed max_linear_work_group_size");
+    }
+  }
+}
+
 /// Note: it is important that this function *does not* depend on kernel
 /// name or kernel type, because then it will be instantiated for every
 /// kernel, even though body of those instantiated functions could be almost
@@ -370,6 +412,7 @@ template <bool IsESIMDKernel = false, typename PropertiesT,
               ext::oneapi::experimental::is_property_list_v<PropertiesT>>>
 constexpr KernelPropertyHolderStructTy
 extractKernelProperties(PropertiesT Props) {
+  validateKernelProperties<PropertiesT>();
   static_assert(
       !PropertiesT::template has_property<
           sycl::ext::intel::experimental::fp_control_key>() ||

@@ -13,7 +13,7 @@
 #include <sycl/detail/helpers.hpp>
 #include <sycl/detail/iostream_proxy.hpp>
 #include <sycl/device.hpp>
-#include <sycl/ext/oneapi/kernel_properties/properties.hpp>
+#include <sycl/ext/oneapi/kernel_properties.hpp>
 #include <sycl/id.hpp>
 #include <sycl/item.hpp>
 #include <sycl/kernel_handler.hpp>

@@ -10,7 +10,7 @@
 
 #include <sycl/ext/intel/esimd/detail/defines_elementary.hpp>
 #include <sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp>
-#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/ext/oneapi/properties/property_value.hpp>
 #include <utility>
 

@@ -10,7 +10,7 @@
 
 #include <sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp>
 #include <sycl/ext/oneapi/experimental/prefetch.hpp>
-#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/ext/oneapi/properties/property.hpp>
 #include <sycl/ext/oneapi/properties/property_value.hpp>
 

@@ -8,7 +8,7 @@
 
 #pragma once
 
-#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/ext/oneapi/properties/property.hpp>
 #include <sycl/ext/oneapi/properties/property_value.hpp>
 

@@ -11,7 +11,7 @@
 #include <sycl/access/access.hpp>                       // for address_space
 #include <sycl/exception.hpp>                           // for make_error_code
 #include <sycl/ext/oneapi/device_global/properties.hpp> // for device_image...
-#include <sycl/ext/oneapi/properties/properties.hpp>    // for properties_t
+#include <sycl/ext/oneapi/properties.hpp>               // for properties_t
 #include <sycl/multi_ptr.hpp>                           // for multi_ptr
 #include <sycl/pointers.hpp>                            // for decorated_gl...
 

@@ -13,7 +13,7 @@
 #include <sycl/ext/intel/experimental/cache_control_properties.hpp>
 #include <sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr_properties.hpp>
 #include <sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp>
-#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/ext/oneapi/properties/property.hpp>
 #include <sycl/ext/oneapi/properties/property_value.hpp>
 #include <sycl/pointers.hpp>

@@ -10,7 +10,7 @@
 #pragma once
 
 #include <sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp>
-#include <sycl/ext/oneapi/properties/properties.hpp> // for properties_t
+#include <sycl/ext/oneapi/properties.hpp> // for properties_t
 #include <sycl/usm/usm_enums.hpp>
 
 #include <type_traits> // for false_type, con...

@@ -8,7 +8,7 @@
 
 #pragma once
 
-#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/range.hpp>
 
 namespace sycl {

@@ -8,7 +8,7 @@
 
 #pragma once
 
-#include <sycl/ext/oneapi/properties/properties.hpp> // for properties_t
+#include <sycl/ext/oneapi/properties.hpp> // for properties_t
 
 #include <type_traits> // for false_type, con...
 #include <utility>     // for declval

@@ -15,7 +15,7 @@
 #include <sycl/ext/oneapi/experimental/enqueue_types.hpp>
 #include <sycl/ext/oneapi/experimental/free_function_traits.hpp>
 #include <sycl/ext/oneapi/experimental/graph.hpp>
-#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/handler.hpp>
 #include <sycl/nd_range.hpp>
 #include <sycl/queue.hpp>

@@ -14,7 +14,7 @@
 #include <sycl/detail/export.hpp>      // for __SYCL_EXPORT
 #include <sycl/detail/kernel_desc.hpp> // for kernel_param_kind_t
 #include <sycl/ext/oneapi/experimental/work_group_memory.hpp> // for work_group_memory
-#include <sycl/ext/oneapi/properties/properties.hpp> // for empty_properties_t
+#include <sycl/ext/oneapi/properties.hpp> // for empty_properties_t
 
 #include <functional> // for function
 #include <memory>     // for shared_ptr

@@ -14,7 +14,7 @@
 #include <sycl/builtins.hpp>            // for min
 #include <sycl/exception.hpp>           // for sycl_category, exception
 #include <sycl/ext/oneapi/bfloat16.hpp> // for bfloat16
-#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/memory_enums.hpp> // for memory_scope
 #include <sycl/range.hpp>        // for range
 #include <sycl/sycl_span.hpp>    // for span

@@ -13,7 +13,7 @@
 #include <sycl/detail/address_space_cast.hpp>
 #include <sycl/detail/helpers.hpp>
 #include <sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp>
-#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/group_barrier.hpp>
 #include <sycl/sycl_span.hpp>
 

@@ -10,7 +10,7 @@
 
 #include <sycl/__spirv/spirv_ops.hpp>
 #include <sycl/detail/address_space_cast.hpp>
-#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/id.hpp>
 #include <sycl/vector.hpp>
 

@@ -9,7 +9,7 @@
 #pragma once
 
 #include <sycl/ext/oneapi/experimental/syclbin_properties.hpp>
-#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/kernel_bundle.hpp>
 
 #include <fstream>

@@ -8,7 +8,7 @@
 
 #pragma once
 
-#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/kernel_bundle.hpp>
 
 namespace sycl {

@@ -11,7 +11,7 @@
 
 #pragma once
 
-#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties.hpp>
 
 namespace sycl {
 inline namespace _V1 {

@@ -1,5 +1,6 @@
 #pragma once
 
+#include <sycl/ext/oneapi/free_function_kernel_properties.hpp>
 #include <sycl/ext/oneapi/properties/property.hpp>
 #include <sycl/ext/oneapi/properties/property_utils.hpp>
 #include <sycl/ext/oneapi/properties/property_value.hpp>
@@ -49,6 +50,17 @@ struct PropertyMetaInfo<indirectly_callable_key::value_t<Set>> {
 #endif
 };
 
+template <typename Set>
+struct FunctionPropertyMetaInfo<indirectly_callable_key::value_t<Set>> {
+  static constexpr const char *name = "indirectly-callable";
+  static constexpr const char *value =
+#ifdef __SYCL_DEVICE_ONLY__
+      __builtin_sycl_unique_stable_name(Set);
+#else
+      "";
+#endif
+};
+
 #ifdef __SYCL_DEVICE_ONLY__
 // Helper to concatenate several lists of characters into a single string.
 // Lists are separated from each other with comma within the resulting string.
@@ -105,6 +117,17 @@ struct PropertyMetaInfo<calls_indirectly_key::value_t<SetIds...>> {
 #endif
 };
 
+template <typename... SetIds>
+struct FunctionPropertyMetaInfo<calls_indirectly_key::value_t<SetIds...>> {
+  static constexpr const char *name = "calls-indirectly";
+  static constexpr const char *value =
+#ifdef __SYCL_DEVICE_ONLY__
+      UniqueStableNameListStr<SetIds...>::value;
+#else
+      "";
+#endif
+};
+
 } // namespace detail
 
 } // namespace ext::oneapi::experimental

@@ -10,7 +10,7 @@
 #include <sycl/access/access.hpp>
 #include <sycl/detail/defines.hpp>
 #include <sycl/ext/oneapi/experimental/graph.hpp>
-#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/multi_ptr.hpp>
 
 #include <cstddef>