Skip to content

Commit f18ed8c

Browse files
frasercrmckfrasercrmck
authored
[SYCL] Detect conflicts between kernel properties (#15510)
The `max_work_group_size` and `max_linear_work_group_size` kernel properties conflict with the `work_group_size` property when the required work-group size exceeds either of the maximum sizes.
1 parent b9eb520 commit f18ed8c

File tree

4 files changed

+156
-17
lines changed

4 files changed

+156
-17
lines changed

‎sycl/include/sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp

Lines changed: 3 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -45,23 +45,12 @@ using HasUsmKind = HasProperty<usm_kind_key, PropertyListT>;
4545
template <typename PropertyListT>
4646
using HasBufferLocation = HasProperty<buffer_location_key, PropertyListT>;
4747

48-
// Get the value of a property from a property list
49-
template <typename PropKey, typename ConstType, typename DefaultPropVal,
50-
typename PropertyListT>
51-
struct GetPropertyValueFromPropList {};
52-
5348
template <typename PropKey, typename ConstType, typename DefaultPropVal,
5449
typename... Props>
5550
struct GetPropertyValueFromPropList<PropKey, ConstType, DefaultPropVal,
56-
detail::properties_t<Props...>> {
57-
using prop_val_t = std::conditional_t<
58-
detail::ContainsProperty<PropKey, std::tuple<Props...>>::value,
59-
typename detail::FindCompileTimePropertyValueType<
60-
PropKey, std::tuple<Props...>>::type,
61-
DefaultPropVal>;
62-
static constexpr ConstType value =
63-
detail::PropertyMetaInfo<std::remove_const_t<prop_val_t>>::value;
64-
};
51+
detail::properties_t<Props...>>
52+
: GetPropertyValueFromPropList<PropKey, ConstType, DefaultPropVal,
53+
std::tuple<Props...>> {};
6554

6655
// Get the value of alignment from a property list
6756
// If alignment is not present in the property list, set to default value 0

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

Lines changed: 76 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -9,13 +9,12 @@
99
#pragma once
1010

1111
#include <array> // for array
12+
#include <limits>
1213
#include <stddef.h> // for size_t
1314
#include <stdint.h> // for uint32_T
1415
#include <sycl/aspects.hpp> // for aspect
1516
#include <sycl/ext/oneapi/experimental/forward_progress.hpp> // for forward_progress_guarantee enum
16-
#include <sycl/ext/oneapi/properties/property.hpp> // for PropKind
17-
#include <sycl/ext/oneapi/properties/property_utils.hpp> // for SizeListToStr
18-
#include <sycl/ext/oneapi/properties/property_value.hpp> // for property_value
17+
#include <sycl/ext/oneapi/properties/properties.hpp>
1918
#include <type_traits> // for true_type
2019
#include <utility> // for declval
2120
namespace sycl {
@@ -351,6 +350,80 @@ struct HasKernelPropertiesGetMethod<T,
351350
decltype(std::declval<T>().get(std::declval<properties_tag>()));
352351
};
353352

353+
// Trait for property compile-time meta names and values.
354+
template <typename PropertyT> struct WGSizePropertyMetaInfo {
355+
static constexpr std::array<size_t, 0> WGSize = {};
356+
static constexpr size_t LinearSize = 0;
357+
};
358+
359+
template <size_t Dim0, size_t... Dims>
360+
struct WGSizePropertyMetaInfo<work_group_size_key::value_t<Dim0, Dims...>> {
361+
static constexpr std::array<size_t, sizeof...(Dims) + 1> WGSize = {Dim0,
362+
Dims...};
363+
static constexpr size_t LinearSize = (Dim0 * ... * Dims);
364+
};
365+
366+
template <size_t Dim0, size_t... Dims>
367+
struct WGSizePropertyMetaInfo<max_work_group_size_key::value_t<Dim0, Dims...>> {
368+
static constexpr std::array<size_t, sizeof...(Dims) + 1> WGSize = {Dim0,
369+
Dims...};
370+
static constexpr size_t LinearSize = (Dim0 * ... * Dims);
371+
};
372+
373+
// Get the value of a work-group size related property from a property list
374+
template <typename PropKey, typename PropertiesT>
375+
struct GetWGPropertyFromPropList {};
376+
377+
template <typename PropKey, typename... PropertiesT>
378+
struct GetWGPropertyFromPropList<PropKey, std::tuple<PropertiesT...>> {
379+
using prop_val_t = std::conditional_t<
380+
ContainsProperty<PropKey, std::tuple<PropertiesT...>>::value,
381+
typename FindCompileTimePropertyValueType<
382+
PropKey, std::tuple<PropertiesT...>>::type,
383+
void>;
384+
static constexpr auto WGSize =
385+
WGSizePropertyMetaInfo<std::remove_const_t<prop_val_t>>::WGSize;
386+
static constexpr size_t LinearSize =
387+
WGSizePropertyMetaInfo<std::remove_const_t<prop_val_t>>::LinearSize;
388+
};
389+
390+
// If work_group_size and max_work_group_size coexist, check that the
391+
// dimensionality matches and that the required work-group size doesn't
392+
// trivially exceed the maximum size.
393+
template <typename Properties>
394+
struct ConflictingProperties<max_work_group_size_key, Properties>
395+
: std::false_type {
396+
using WGSizeVal = GetWGPropertyFromPropList<work_group_size_key, Properties>;
397+
using MaxWGSizeVal =
398+
GetWGPropertyFromPropList<max_work_group_size_key, Properties>;
399+
// If work_group_size_key doesn't exist in the list of properties, WGSize is
400+
// an empty array and so Dims == 0.
401+
static constexpr size_t Dims = WGSizeVal::WGSize.size();
402+
static_assert(
403+
Dims == 0 || Dims == MaxWGSizeVal::WGSize.size(),
404+
"work_group_size and max_work_group_size dimensionality must match");
405+
static_assert(Dims < 1 || WGSizeVal::WGSize[0] <= MaxWGSizeVal::WGSize[0],
406+
"work_group_size must not exceed max_work_group_size");
407+
static_assert(Dims < 2 || WGSizeVal::WGSize[1] <= MaxWGSizeVal::WGSize[1],
408+
"work_group_size must not exceed max_work_group_size");
409+
static_assert(Dims < 3 || WGSizeVal::WGSize[2] <= MaxWGSizeVal::WGSize[2],
410+
"work_group_size must not exceed max_work_group_size");
411+
};
412+
413+
// If work_group_size and max_linear_work_group_size coexist, check that the
414+
// required linear work-group size doesn't trivially exceed the maximum size.
415+
template <typename Properties>
416+
struct ConflictingProperties<max_linear_work_group_size_key, Properties>
417+
: std::false_type {
418+
using WGSizeVal = GetWGPropertyFromPropList<work_group_size_key, Properties>;
419+
using MaxLinearWGSizeVal =
420+
GetPropertyValueFromPropList<max_linear_work_group_size_key, size_t, void,
421+
Properties>;
422+
static_assert(WGSizeVal::WGSize.empty() ||
423+
WGSizeVal::LinearSize <= MaxLinearWGSizeVal::value,
424+
"work_group_size must not exceed max_linear_work_group_size");
425+
};
426+
354427
} // namespace detail
355428
} // namespace ext::oneapi::experimental
356429
} // namespace _V1

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

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -150,6 +150,24 @@ struct ExtractProperties<PropertyArgsT,
150150
}
151151
};
152152

153+
// Get the value of a property from a property list
154+
template <typename PropKey, typename ConstType, typename DefaultPropVal,
155+
typename PropertiesT>
156+
struct GetPropertyValueFromPropList {};
157+
158+
template <typename PropKey, typename ConstType, typename DefaultPropVal,
159+
typename... PropertiesT>
160+
struct GetPropertyValueFromPropList<PropKey, ConstType, DefaultPropVal,
161+
std::tuple<PropertiesT...>> {
162+
using prop_val_t = std::conditional_t<
163+
ContainsProperty<PropKey, std::tuple<PropertiesT...>>::value,
164+
typename FindCompileTimePropertyValueType<
165+
PropKey, std::tuple<PropertiesT...>>::type,
166+
DefaultPropVal>;
167+
static constexpr ConstType value =
168+
PropertyMetaInfo<std::remove_const_t<prop_val_t>>::value;
169+
};
170+
153171
} // namespace detail
154172

155173
template <typename PropertiesT> class properties {

‎sycl/test/extensions/properties/properties_kernel_negative.cpp

Lines changed: 59 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -317,9 +317,68 @@ void check_sub_group_size() {
317317
KernelFunctorWithSGSize<2>{});
318318
}
319319

320+
void check_max_work_group_size() {
321+
sycl::queue Q;
322+
323+
// expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size and max_work_group_size dimensionality must match}}
324+
Q.single_task(
325+
sycl::ext::oneapi::experimental::properties{
326+
sycl::ext::oneapi::experimental::work_group_size<2, 2>,
327+
sycl::ext::oneapi::experimental::max_work_group_size<1>},
328+
[]() {});
329+
330+
// expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size must not exceed max_work_group_size}}
331+
Q.single_task(
332+
sycl::ext::oneapi::experimental::properties{
333+
sycl::ext::oneapi::experimental::work_group_size<2>,
334+
sycl::ext::oneapi::experimental::max_work_group_size<1>},
335+
[]() {});
336+
337+
// expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size must not exceed max_work_group_size}}
338+
Q.single_task(
339+
sycl::ext::oneapi::experimental::properties{
340+
sycl::ext::oneapi::experimental::work_group_size<2, 2>,
341+
sycl::ext::oneapi::experimental::max_work_group_size<2, 1>},
342+
[]() {});
343+
344+
// expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size must not exceed max_work_group_size}}
345+
Q.single_task(
346+
sycl::ext::oneapi::experimental::properties{
347+
sycl::ext::oneapi::experimental::work_group_size<2, 2, 2>,
348+
sycl::ext::oneapi::experimental::max_work_group_size<2, 2, 1>},
349+
[]() {});
350+
}
351+
352+
void check_max_linear_work_group_size() {
353+
sycl::queue Q;
354+
355+
// expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size must not exceed max_linear_work_group_size}}
356+
Q.single_task(
357+
sycl::ext::oneapi::experimental::properties{
358+
sycl::ext::oneapi::experimental::work_group_size<2>,
359+
sycl::ext::oneapi::experimental::max_linear_work_group_size<1>},
360+
[]() {});
361+
362+
// expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size must not exceed max_linear_work_group_size}}
363+
Q.single_task(
364+
sycl::ext::oneapi::experimental::properties{
365+
sycl::ext::oneapi::experimental::work_group_size<2, 4>,
366+
sycl::ext::oneapi::experimental::max_linear_work_group_size<7>},
367+
[]() {});
368+
369+
// expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size must not exceed max_linear_work_group_size}}
370+
Q.single_task(
371+
sycl::ext::oneapi::experimental::properties{
372+
sycl::ext::oneapi::experimental::work_group_size<2, 4, 2>,
373+
sycl::ext::oneapi::experimental::max_linear_work_group_size<15>},
374+
[]() {});
375+
}
376+
320377
int main() {
321378
check_work_group_size();
322379
check_work_group_size_hint();
323380
check_sub_group_size();
381+
check_max_work_group_size();
382+
check_max_linear_work_group_size();
324383
return 0;
325384
}

0 commit comments

Comments
 (0)