Skip to content

Commit 5eb3536

Browse files
lbushi25lbushi25
authored
[SYCL] Implement forward_progress extension API (#13389)
This PR implements the API for the forward progress guarantees extension defined here: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_forward_progress.asciidoc ~~It is complete except that it does not verify that the forward progress guarantees requested by a kernel submission are actually supported by the device to which the kernel is submitted. That will be done in a later PR.~~
1 parent 34292bb commit 5eb3536

23 files changed

+976
-15
lines changed

‎sycl/include/sycl/detail/pi.h

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -437,7 +437,12 @@ typedef enum {
437437
// The number of max registers per block (device specific)
438438
PI_EXT_CODEPLAY_DEVICE_INFO_MAX_REGISTERS_PER_WORK_GROUP = 0x20009,
439439
PI_EXT_INTEL_DEVICE_INFO_ESIMD_SUPPORT = 0x2000A,
440-
440+
PI_EXT_ONEAPI_DEVICE_INFO_WORK_GROUP_PROGRESS_AT_ROOT_GROUP_LEVEL = 0x2000B,
441+
PI_EXT_ONEAPI_DEVICE_INFO_SUB_GROUP_PROGRESS_AT_ROOT_GROUP_LEVEL = 0x2000C,
442+
PI_EXT_ONEAPI_DEVICE_INFO_SUB_GROUP_PROGRESS_AT_WORK_GROUP_LEVEL = 0x2000D,
443+
PI_EXT_ONEAPI_DEVICE_INFO_WORK_ITEM_PROGRESS_AT_ROOT_GROUP_LEVEL = 0x2000E,
444+
PI_EXT_ONEAPI_DEVICE_INFO_WORK_ITEM_PROGRESS_AT_WORK_GROUP_LEVEL = 0x2000F,
445+
PI_EXT_ONEAPI_DEVICE_INFO_WORK_ITEM_PROGRESS_AT_SUB_GROUP_LEVEL = 0x20010,
441446
// Bindless images, mipmaps, interop
442447
PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_IMAGES_SUPPORT = 0x20100,
443448
PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_IMAGES_SHARED_USM_SUPPORT = 0x20101,

‎sycl/include/sycl/ext/oneapi/experimental/forward_progress.hpp

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,30 @@
1+
//==------- forward_progress.hpp - sycl_ext_oneapi_forward_progress -------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#pragma once
10+
11+
namespace sycl {
12+
inline namespace _V1 {
13+
namespace ext {
14+
namespace oneapi {
15+
namespace experimental {
16+
17+
enum class forward_progress_guarantee { concurrent, parallel, weakly_parallel };
18+
19+
enum class execution_scope {
20+
work_item,
21+
sub_group,
22+
work_group,
23+
root_group,
24+
};
25+
26+
} // namespace experimental
27+
} // namespace oneapi
28+
} // namespace ext
29+
} // namespace _V1
30+
} // namespace sycl

‎sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp

Lines changed: 91 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -8,17 +8,16 @@
88

99
#pragma once
1010

11-
#include <sycl/aspects.hpp> // for aspect
11+
#include <array> // for array
12+
#include <stddef.h> // for size_t
13+
#include <stdint.h> // for uint32_T
14+
#include <sycl/aspects.hpp> // for aspect
15+
#include <sycl/ext/oneapi/experimental/forward_progress.hpp> // for forward_progress_guarantee enum
1216
#include <sycl/ext/oneapi/properties/property.hpp> // for PropKind
1317
#include <sycl/ext/oneapi/properties/property_utils.hpp> // for SizeListToStr
1418
#include <sycl/ext/oneapi/properties/property_value.hpp> // for property_value
15-
16-
#include <array> // for array
17-
#include <stddef.h> // for size_t
18-
#include <stdint.h> // for uint32_t
19-
#include <type_traits> // for true_type
20-
#include <utility> // for declval
21-
19+
#include <type_traits> // for true_type
20+
#include <utility> // for declval
2221
namespace sycl {
2322
inline namespace _V1 {
2423
namespace ext::oneapi::experimental {
@@ -55,7 +54,8 @@ struct sub_group_size_key
5554
std::integral_constant<uint32_t, Size>>;
5655
};
5756

58-
struct device_has_key : detail::compile_time_property_key<detail::PropKind::DeviceHas> {
57+
struct device_has_key
58+
: detail::compile_time_property_key<detail::PropKind::DeviceHas> {
5959
template <aspect... Aspects>
6060
using value_t = property_value<device_has_key,
6161
std::integral_constant<aspect, Aspects>...>;
@@ -126,6 +126,88 @@ inline constexpr sub_group_size_key::value_t<Size> sub_group_size;
126126
template <aspect... Aspects>
127127
inline constexpr device_has_key::value_t<Aspects...> device_has;
128128

129+
struct work_group_progress_key
130+
: detail::compile_time_property_key<detail::PropKind::WorkGroupProgress> {
131+
template <forward_progress_guarantee Guarantee,
132+
execution_scope CoordinationScope>
133+
using value_t = property_value<
134+
work_group_progress_key,
135+
std::integral_constant<forward_progress_guarantee, Guarantee>,
136+
std::integral_constant<execution_scope, CoordinationScope>>;
137+
};
138+
139+
struct sub_group_progress_key
140+
: detail::compile_time_property_key<detail::PropKind::SubGroupProgress> {
141+
template <forward_progress_guarantee Guarantee,
142+
execution_scope CoordinationScope>
143+
using value_t = property_value<
144+
sub_group_progress_key,
145+
std::integral_constant<forward_progress_guarantee, Guarantee>,
146+
std::integral_constant<execution_scope, CoordinationScope>>;
147+
};
148+
149+
struct work_item_progress_key
150+
: detail::compile_time_property_key<detail::PropKind::WorkItemProgress> {
151+
template <forward_progress_guarantee Guarantee,
152+
execution_scope CoordinationScope>
153+
using value_t = property_value<
154+
work_item_progress_key,
155+
std::integral_constant<forward_progress_guarantee, Guarantee>,
156+
std::integral_constant<execution_scope, CoordinationScope>>;
157+
};
158+
159+
template <forward_progress_guarantee Guarantee,
160+
execution_scope CoordinationScope>
161+
struct property_value<
162+
work_group_progress_key,
163+
std::integral_constant<forward_progress_guarantee, Guarantee>,
164+
std::integral_constant<execution_scope, CoordinationScope>> {
165+
using key_t = work_group_progress_key;
166+
static constexpr forward_progress_guarantee guarantee = Guarantee;
167+
static constexpr execution_scope coordinationScope = CoordinationScope;
168+
};
169+
170+
template <forward_progress_guarantee Guarantee,
171+
execution_scope CoordinationScope>
172+
struct property_value<
173+
sub_group_progress_key,
174+
std::integral_constant<forward_progress_guarantee, Guarantee>,
175+
std::integral_constant<execution_scope, CoordinationScope>> {
176+
using key_t = work_group_progress_key;
177+
static constexpr forward_progress_guarantee guarantee = Guarantee;
178+
static constexpr execution_scope coordinationScope = CoordinationScope;
179+
};
180+
181+
template <forward_progress_guarantee Guarantee,
182+
execution_scope CoordinationScope>
183+
struct property_value<
184+
work_item_progress_key,
185+
std::integral_constant<forward_progress_guarantee, Guarantee>,
186+
std::integral_constant<execution_scope, CoordinationScope>> {
187+
using key_t = work_group_progress_key;
188+
static constexpr forward_progress_guarantee guarantee = Guarantee;
189+
static constexpr execution_scope coordinationScope = CoordinationScope;
190+
};
191+
192+
template <forward_progress_guarantee Guarantee,
193+
execution_scope CoordinationScope>
194+
inline constexpr work_group_progress_key::value_t<Guarantee, CoordinationScope>
195+
work_group_progress;
196+
197+
template <forward_progress_guarantee Guarantee,
198+
execution_scope CoordinationScope>
199+
inline constexpr sub_group_progress_key::value_t<Guarantee, CoordinationScope>
200+
sub_group_progress;
201+
202+
template <forward_progress_guarantee Guarantee,
203+
execution_scope CoordinationScope>
204+
inline constexpr work_item_progress_key::value_t<Guarantee, CoordinationScope>
205+
work_item_progress;
206+
207+
template <> struct is_property_key<work_group_progress_key> : std::true_type {};
208+
template <> struct is_property_key<sub_group_progress_key> : std::true_type {};
209+
template <> struct is_property_key<work_item_progress_key> : std::true_type {};
210+
129211
namespace detail {
130212
template <size_t Dim0, size_t... Dims>
131213
struct PropertyMetaInfo<work_group_size_key::value_t<Dim0, Dims...>> {

‎sycl/include/sycl/ext/oneapi/properties/property.hpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -200,8 +200,11 @@ enum PropKind : uint32_t {
200200
ContiguousMemory = 59,
201201
FullGroup = 60,
202202
Naive = 61,
203+
WorkGroupProgress = 62,
204+
SubGroupProgress = 63,
205+
WorkItemProgress = 64,
203206
// PropKindSize must always be the last value.
204-
PropKindSize = 62,
207+
PropKindSize = 65,
205208
};
206209

207210
struct property_key_base_tag {};

‎sycl/include/sycl/handler.hpp

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -938,6 +938,11 @@ class __SYCL_EXPORT handler {
938938
}
939939
}
940940

941+
void verifyDeviceHasProgressGuarantee(
942+
sycl::ext::oneapi::experimental::forward_progress_guarantee guarantee,
943+
sycl::ext::oneapi::experimental::execution_scope threadScope,
944+
sycl::ext::oneapi::experimental::execution_scope coordinationScope);
945+
941946
/// Process kernel properties.
942947
///
943948
/// Stores information about kernel properties into the handler.
@@ -972,6 +977,36 @@ class __SYCL_EXPORT handler {
972977
constexpr bool UsesRootSync = PropertiesT::template has_property<
973978
sycl::ext::oneapi::experimental::use_root_sync_key>();
974979
setKernelIsCooperative(UsesRootSync);
980+
if constexpr (PropertiesT::template has_property<
981+
sycl::ext::oneapi::experimental::
982+
work_group_progress_key>()) {
983+
auto prop = Props.template get_property<
984+
sycl::ext::oneapi::experimental::work_group_progress_key>();
985+
verifyDeviceHasProgressGuarantee(
986+
prop.guarantee,
987+
sycl::ext::oneapi::experimental::execution_scope::work_group,
988+
prop.coordinationScope);
989+
}
990+
if constexpr (PropertiesT::template has_property<
991+
sycl::ext::oneapi::experimental::
992+
sub_group_progress_key>()) {
993+
auto prop = Props.template get_property<
994+
sycl::ext::oneapi::experimental::sub_group_progress_key>();
995+
verifyDeviceHasProgressGuarantee(
996+
prop.guarantee,
997+
sycl::ext::oneapi::experimental::execution_scope::sub_group,
998+
prop.coordinationScope);
999+
}
1000+
if constexpr (PropertiesT::template has_property<
1001+
sycl::ext::oneapi::experimental::
1002+
work_item_progress_key>()) {
1003+
auto prop = Props.template get_property<
1004+
sycl::ext::oneapi::experimental::work_item_progress_key>();
1005+
verifyDeviceHasProgressGuarantee(
1006+
prop.guarantee,
1007+
sycl::ext::oneapi::experimental::execution_scope::work_item,
1008+
prop.coordinationScope);
1009+
}
9751010
}
9761011

9771012
/// Checks whether it is possible to copy the source shape to the destination

‎sycl/include/sycl/info/ext_oneapi_device_traits.def

Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,45 @@ __SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental,device, max_global_work_group
66
__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(ext::oneapi::experimental,device, max_work_groups<1>, id<1>, PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_1D)
77
__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(ext::oneapi::experimental,device, max_work_groups<2>, id<2>, PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_2D)
88
__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(ext::oneapi::experimental,device, max_work_groups<3>, id<3>, PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_3D)
9+
10+
// Forward progress guarantees
11+
__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(
12+
ext::oneapi::experimental, device,
13+
work_group_progress_capabilities<
14+
ext::oneapi::experimental::execution_scope::root_group>,
15+
std::vector<ext::oneapi::experimental::forward_progress_guarantee>,
16+
PI_EXT_ONEAPI_DEVICE_INFO_WORK_GROUP_PROGRESS_AT_ROOT_GROUP_LEVEL)
17+
__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(
18+
ext::oneapi::experimental, device,
19+
sub_group_progress_capabilities<
20+
ext::oneapi::experimental::execution_scope::root_group>,
21+
std::vector<ext::oneapi::experimental::forward_progress_guarantee>,
22+
PI_EXT_ONEAPI_DEVICE_INFO_SUB_GROUP_PROGRESS_AT_ROOT_GROUP_LEVEL)
23+
__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(
24+
ext::oneapi::experimental, device,
25+
sub_group_progress_capabilities<
26+
ext::oneapi::experimental::execution_scope::work_group>,
27+
std::vector<ext::oneapi::experimental::forward_progress_guarantee>,
28+
PI_EXT_ONEAPI_DEVICE_INFO_SUB_GROUP_PROGRESS_AT_WORK_GROUP_LEVEL)
29+
__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(
30+
ext::oneapi::experimental, device,
31+
work_item_progress_capabilities<
32+
ext::oneapi::experimental::execution_scope::root_group>,
33+
std::vector<ext::oneapi::experimental::forward_progress_guarantee>,
34+
PI_EXT_ONEAPI_DEVICE_INFO_WORK_ITEM_PROGRESS_AT_ROOT_GROUP_LEVEL)
35+
__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(
36+
ext::oneapi::experimental, device,
37+
work_item_progress_capabilities<
38+
ext::oneapi::experimental::execution_scope::work_group>,
39+
std::vector<ext::oneapi::experimental::forward_progress_guarantee>,
40+
PI_EXT_ONEAPI_DEVICE_INFO_WORK_ITEM_PROGRESS_AT_WORK_GROUP_LEVEL)
41+
__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(
42+
ext::oneapi::experimental, device,
43+
work_item_progress_capabilities<
44+
ext::oneapi::experimental::execution_scope::sub_group>,
45+
std::vector<ext::oneapi::experimental::forward_progress_guarantee>,
46+
PI_EXT_ONEAPI_DEVICE_INFO_WORK_ITEM_PROGRESS_AT_SUB_GROUP_LEVEL)
47+
948
__SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, device, architecture,
1049
ext::oneapi::experimental::architecture,
1150
PI_EXT_ONEAPI_DEVICE_INFO_IP_VERSION)

‎sycl/include/sycl/info/info_desc.hpp

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -17,6 +17,7 @@
1717
#include <sycl/aspects.hpp>
1818
#include <sycl/detail/type_traits.hpp>
1919
#include <sycl/ext/oneapi/experimental/device_architecture.hpp>
20+
#include <sycl/ext/oneapi/experimental/forward_progress.hpp>
2021
#include <sycl/ext/oneapi/matrix/query-types.hpp>
2122

2223
#include <sycl/range.hpp>
@@ -195,6 +196,13 @@ template <typename T, T param> struct compatibility_param_traits {};
195196

196197
namespace ext::oneapi::experimental::info::device {
197198
template <int Dimensions> struct max_work_groups;
199+
template <ext::oneapi::experimental::execution_scope CoordinationScope>
200+
struct work_group_progress_capabilities;
201+
template <ext::oneapi::experimental::execution_scope CoordinationScope>
202+
struct sub_group_progress_capabilities;
203+
template <ext::oneapi::experimental::execution_scope CoordinationScope>
204+
struct work_item_progress_capabilities;
205+
198206
} // namespace ext::oneapi::experimental::info::device
199207
#include <sycl/info/ext_codeplay_device_traits.def>
200208
#include <sycl/info/ext_intel_device_traits.def>

‎sycl/include/sycl/kernel.hpp

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,8 @@
88

99
#pragma once
1010

11+
#include <cstddef> // for size_t
12+
#include <memory> // for shared_ptr, hash, opera...
1113
#include <sycl/backend_types.hpp> // for backend, backend_return_t
1214
#include <sycl/context.hpp> // for context
1315
#include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEPRECATED
@@ -21,10 +23,7 @@
2123
#include <sycl/device.hpp> // for device
2224
#include <sycl/kernel_bundle_enums.hpp> // for bundle_state
2325
#include <sycl/range.hpp> // for range
24-
25-
#include <cstddef> // for size_t
26-
#include <memory> // for shared_ptr, hash, opera...
27-
#include <variant> // for hash
26+
#include <variant> // for hash
2827

2928
namespace sycl {
3029
inline namespace _V1 {

‎sycl/include/sycl/sycl.hpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -93,6 +93,7 @@
9393
#include <sycl/ext/oneapi/experimental/cuda/barrier.hpp>
9494
#include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
9595
#include <sycl/ext/oneapi/experimental/fixed_size_group.hpp>
96+
#include <sycl/ext/oneapi/experimental/forward_progress.hpp>
9697
#include <sycl/ext/oneapi/experimental/group_load_store.hpp>
9798
#include <sycl/ext/oneapi/experimental/opportunistic_group.hpp>
9899
#include <sycl/ext/oneapi/experimental/prefetch.hpp>

0 commit comments

Comments
 (0)