[SYCL][Docs] Add sycl_ext_named_sub_group_sizes kernel properties

sycl/doc/extensions/proposed/sycl_ext_oneapi_named_sub_group_sizes.asciidoc → sycl/doc/extensions/experimental/sycl_ext_oneapi_named_sub_group_sizes.asciidoc

@@ -51,11 +51,12 @@ This extension also depends on the following other SYCL extensions:
 
 == 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.
+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

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

@@ -0,0 +1,45 @@
+//== named_sub_group_sizes.hpp --- SYCL extension for named sub-group sizes ==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/ext/oneapi/kernel_properties/properties.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+
+struct named_sub_group_size {
+  static constexpr uint32_t primary = -1;
+  static constexpr uint32_t automatic = -2;
+};
+
+inline constexpr sub_group_size_key::value_t<named_sub_group_size::primary>
+    sub_group_size_primary;
+
+inline constexpr sub_group_size_key::value_t<named_sub_group_size::automatic>
+    sub_group_size_automatic;
+
+namespace detail {
+template <>
+struct PropertyMetaInfo<
+    sub_group_size_key::value_t<named_sub_group_size::automatic>> {
+  // sub_group_size_automatic means that the kernel can be compiled with
+  // any sub-group size. That is, if the kernel has the sub_group_size_automatic
+  // property, then no sycl-sub-group-size IR attribute needs to be attached.
+  // Specializing PropertyMetaInfo for sub_group_size_automatic and setting
+  // name to an empty string will result in no sycl-sub-group-size IR being
+  // attached.
+  static constexpr const char *name = "";
+  static constexpr const char *value = 0;
+};
+} // namespace detail
+
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/info/device_traits.def

@@ -208,6 +208,8 @@ __SYCL_PARAM_TRAITS_SPEC(device, ext_oneapi_srgb, bool,
                          UR_DEVICE_INFO_IMAGE_SRGB)
 __SYCL_PARAM_TRAITS_SPEC(device, ext_intel_mem_channel, bool,
                          UR_DEVICE_INFO_MEM_CHANNEL_SUPPORT)
+__SYCL_PARAM_TRAITS_SPEC(device, primary_sub_group_size, uint32_t,
+                         UR_DEVICE_INFO_SUB_GROUP_PRIMARY_SIZE_EXP)
 
 //Deprecated oneapi/intel extension
 //TODO:Remove when possible

sycl/include/sycl/sycl.hpp

@@ -127,6 +127,7 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.")
 #include <sycl/ext/oneapi/experimental/group_helpers_sorters.hpp>
 #include <sycl/ext/oneapi/experimental/group_load_store.hpp>
 #include <sycl/ext/oneapi/experimental/group_sort.hpp>
+#include <sycl/ext/oneapi/experimental/named_sub_group_sizes.hpp>
 #include <sycl/ext/oneapi/experimental/prefetch.hpp>
 #include <sycl/ext/oneapi/experimental/profiling_tag.hpp>
 #include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>

sycl/source/detail/ur_device_info_ret_types.inc

@@ -162,6 +162,7 @@ MAP(UR_DEVICE_INFO_NODE_MASK, uint32_t)
 // These aren't present in the specification, extracted from ur_api.h
 // instead.
 MAP(UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP, ur_exp_device_2d_block_array_capability_flags_t)
+MAP(UR_DEVICE_INFO_SUB_GROUP_PRIMARY_SIZE_EXP, uint32_t)
 MAP(UR_DEVICE_INFO_ASYNC_USM_ALLOCATIONS_SUPPORT_EXP, ur_bool_t)
 MAP(UR_DEVICE_INFO_BINDLESS_IMAGES_1D_USM_SUPPORT_EXP, ur_bool_t)
 MAP(UR_DEVICE_INFO_BINDLESS_IMAGES_2D_USM_SUPPORT_EXP, ur_bool_t)

sycl/source/feature_test.hpp.in

@@ -124,6 +124,7 @@ inline namespace _V1 {
 #define SYCL_KHR_DEFAULT_CONTEXT 1
 #define SYCL_EXT_INTEL_EVENT_MODE 1
 #define SYCL_EXT_ONEAPI_TANGLE 1
+#define SYCL_EXT_ONEAPI_NAMED_SUB_GROUP_SIZES 1
 
 // Unfinished KHR extensions. These extensions are only available if the
 // __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS macro is defined.

sycl/test-e2e/SubGroup/primary_sub_group_size.cpp

@@ -0,0 +1,35 @@
+// REQUIRES: aspect-usm_shared_allocations
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/named_sub_group_sizes.hpp>
+#include <sycl/sub_group.hpp>
+#include <sycl/usm.hpp>
+
+struct SGSizePrimaryKernelFunctor {
+  SGSizePrimaryKernelFunctor(uint32_t *OutPtr) : Out{OutPtr} {}
+
+  void operator()(sycl::nd_item<1> Item) const {
+    *Out = Item.get_sub_group().get_max_local_range()[0];
+  }
+
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
+    return sycl::ext::oneapi::experimental::properties{
+        sycl::ext::oneapi::experimental::sub_group_size_primary};
+  }
+
+  uint32_t *Out;
+};
+
+int main() {
+  sycl::queue Q;
+
+  uint32_t *OutPtr = sycl::malloc_shared<uint32_t>(1, Q);
+  Q.parallel_for(sycl::nd_range<1>{1, 1}, SGSizePrimaryKernelFunctor{OutPtr})
+      .wait();
+
+  assert(*OutPtr ==
+         Q.get_device().get_info<sycl::info::device::primary_sub_group_size>());
+  return 0;
+}

sycl/test/abi/sycl_symbols_linux.dump

@@ -3970,6 +3970,7 @@ _ZNK4sycl3_V16device13get_info_implINS0_4info6device21image_max_buffer_sizeEEENS
 _ZNK4sycl3_V16device13get_info_implINS0_4info6device21is_compiler_availableEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_4info6device22execution_capabilitiesEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_4info6device22ext_intel_gpu_eu_countEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv
+_ZNK4sycl3_V16device13get_info_implINS0_4info6device22primary_sub_group_sizeEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_4info6device22usm_device_allocationsEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_4info6device22usm_shared_allocationsEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_4info6device22usm_system_allocationsEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv

sycl/test/extensions/properties/properties_kernel_named_sub_group_size.cpp

@@ -0,0 +1,42 @@
+// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
+// expected-no-diagnostics
+#include <sycl/sycl.hpp>
+
+struct SGSizePrimaryKernelFunctor {
+  SGSizePrimaryKernelFunctor() {}
+
+  void operator()(sycl::nd_item<1>) const {}
+
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
+    return sycl::ext::oneapi::experimental::properties{
+        sycl::ext::oneapi::experimental::sub_group_size_primary};
+  }
+};
+
+struct SGSizeAutoKernelFunctor {
+  SGSizeAutoKernelFunctor() {}
+
+  void operator()(sycl::nd_item<1>) const {}
+
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
+    return sycl::ext::oneapi::experimental::properties{
+        sycl::ext::oneapi::experimental::sub_group_size_automatic};
+  }
+};
+
+int main() {
+  sycl::queue Q;
+  sycl::nd_range<1> NDRange{6, 2};
+
+  // CHECK: spir_kernel void @{{.*}}SGSizePrimaryKernelFunctor()
+  // CHECK-SAME: !intel_reqd_sub_group_size ![[SGSizeAttr:[0-9]+]]
+  Q.parallel_for(NDRange, SGSizePrimaryKernelFunctor{});
+
+  // CHECK: spir_kernel void @{{.*}}SGSizeAutoKernelFunctor()
+  // CHECK-NOT: intel_reqd_sub_group_size
+  // CHECK-SAME: {
+  Q.parallel_for(NDRange, SGSizeAutoKernelFunctor{});
+}
+
+// CHECK: ![[SGSizeAttr]] = !{i32 -1}

unified-runtime/scripts/core/EXP-SUB-GROUP-PRIMARY-SIZE.rst

@@ -0,0 +1,60 @@
+<%
+    OneApi=tags['$OneApi']
+    x=tags['$x']
+    X=x.upper()
+%>
+
+.. _experimental-sub-group-primary-size:
+
+================================================================================
+Sub-group primary size
+================================================================================
+
+.. warning::
+
+    Experimental features:
+
+    *   May be replaced, updated, or removed at any time.
+    *   Do not require maintaining API/ABI stability of their own additions over
+        time.
+    *   Do not require conformance testing of their own additions.
+
+
+Motivation
+--------------------------------------------------------------------------------
+Some devices expose a "primary" sub-group size, which is a device-specific named
+size that is independent of the kernels run on it. Usually, this sub-group size
+can be specified by name in kernel code, but in order for the host code to know
+this size, the corresponding device info query is introduced.
+
+API
+--------------------------------------------------------------------------------
+
+Enums
+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
+
+* ${x}_device_info_t
+    * ${X}_DEVICE_INFO_SUB_GROUP_PRIMARY_SIZE_EXP
+
+Changelog
+--------------------------------------------------------------------------------
+
++-----------+------------------------+
+| Revision  | Changes                |
++===========+========================+
+| 1.0       | Initial Draft          |
++-----------+------------------------+
+
+
+Support
+--------------------------------------------------------------------------------
+
+Adapters which support this experimental feature *must* return ${X}_RESULT_SUCCESS from
+the ${x}DeviceGetInfo call with the new ${X}_DEVICE_INFO_SUB_GROUP_PRIMARY_SIZE_EXP
+device descriptor.
+
+
+Contributors
+--------------------------------------------------------------------------------
+
+* Steffen Larsen `[email protected] <[email protected]>`_

unified-runtime/scripts/core/exp-sub-group-primary-size.yml

@@ -0,0 +1,24 @@
+#
+# Copyright (C) 2025 Intel Corporation
+#
+# Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions.
+# See LICENSE.TXT
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#
+# See YaML.md for syntax definition
+#
+--- #--------------------------------------------------------------------------
+type: header
+desc: "Intel $OneApi Unified Runtime Experimental device descriptor for querying the primary sub-group size"
+ordinal: "99"
+--- #--------------------------------------------------------------------------
+type: enum
+extend: true
+typed_etors: true
+desc: "Extension enum to $x_device_info_t to query primary sub-group size."
+name: $x_device_info_t
+etors:
+    - name: SUB_GROUP_PRIMARY_SIZE_EXP
+      value: "0x2023"
+      desc: "[uint32_t][optional-query] return a 32-bit unsigned integer representing the primary sub-group size of the device."
+

unified-runtime/source/adapters/cuda/device.cpp

@@ -1146,6 +1146,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
   case UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP:
     return ReturnValue(
         static_cast<ur_exp_device_2d_block_array_capability_flags_t>(0));
+  case UR_DEVICE_INFO_SUB_GROUP_PRIMARY_SIZE_EXP: {
+    // NVIDIA devices only support one sub-group size (the warp size)
+    int WarpSize = 0;
+    UR_CHECK_ERROR(cuDeviceGetAttribute(
+        &WarpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, hDevice->get()));
+    return ReturnValue(static_cast<uint32_t>(WarpSize));
+  }
   case UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP:
   case UR_DEVICE_INFO_COMMAND_BUFFER_EVENT_SUPPORT_EXP:
     return ReturnValue(true);

unified-runtime/source/adapters/hip/device.cpp

@@ -999,6 +999,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
   case UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP:
     return ReturnValue(
         static_cast<ur_exp_device_2d_block_array_capability_flags_t>(0));
+  case UR_DEVICE_INFO_SUB_GROUP_PRIMARY_SIZE_EXP: {
+    int WarpSize = 0;
+    UR_CHECK_ERROR(hipDeviceGetAttribute(&WarpSize, hipDeviceAttributeWarpSize,
+                                         hDevice->get()));
+    return ReturnValue(static_cast<uint32_t>(WarpSize));
+  }
   case UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP: {
     int RuntimeVersion = 0;
     UR_CHECK_ERROR(hipRuntimeGetVersion(&RuntimeVersion));

unified-runtime/source/adapters/level_zero/device.cpp

@@ -1271,6 +1271,9 @@ ur_result_t urDeviceGetInfo(
     return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION;
 #endif
   }
+  case UR_DEVICE_INFO_SUB_GROUP_PRIMARY_SIZE_EXP:
+    // Currently not supported.
+    return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION;
   case UR_DEVICE_INFO_ASYNC_BARRIER:
     return ReturnValue(false);
   case UR_DEVICE_INFO_HOST_PIPE_READ_WRITE_SUPPORT:

unified-runtime/source/adapters/native_cpu/device.cpp

@@ -376,6 +376,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
   case UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP:
     return ReturnValue(
         static_cast<ur_exp_device_2d_block_array_capability_flags_t>(0));
+  case UR_DEVICE_INFO_SUB_GROUP_PRIMARY_SIZE_EXP:
+    return ReturnValue(static_cast<uint32_t>(1));
   case UR_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES: {
     // Currently for Native CPU fences are implemented using OCK
     // builtins, so we have different capabilities than atomic operations

unified-runtime/source/adapters/opencl/device.cpp

@@ -1359,6 +1359,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
     return ReturnValue(UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_LOAD |
                        UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_STORE);
   }
+  case UR_DEVICE_INFO_SUB_GROUP_PRIMARY_SIZE_EXP:
+    // Currently not supported. Depends on
+    // https://github.com/intel/llvm/pull/11301.
+    return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION;
   case UR_DEVICE_INFO_BFLOAT16_CONVERSIONS_NATIVE: {
     bool Supported = false;
     UR_RETURN_ON_FAILURE(hDevice->checkDeviceExtensions(

unified-runtime/test/conformance/device/urDeviceGetInfo.cpp

@@ -2527,6 +2527,21 @@ TEST_P(urDeviceGetInfoTest, Success2DBlockArrayCapabilities) {
             0);
 }
 
+TEST_P(urDeviceGetInfoTest, SuccessSubGroupPrimarySize) {
+  size_t property_size = 0;
+  const ur_device_info_t property_name = SUB_GROUP_PRIMARY_SIZE_EXP;
+
+  UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(
+      urDeviceGetInfo(device, property_name, 0, nullptr, &property_size));
+  ASSERT_EQ(property_size, sizeof(uint32_t));
+
+  uint32_t property_value = 0;
+  ASSERT_QUERY_RETURNS_VALUE(urDeviceGetInfo(device, property_name,
+                                             property_size, &property_value,
+                                             nullptr),
+                             property_value);
+}
+
 TEST_P(urDeviceGetInfoTest, SuccessUseNativeAssert) {
   size_t property_size = 0;
   const ur_device_info_t property_name = UR_DEVICE_INFO_USE_NATIVE_ASSERT;