|
12 | 12 | #include <sycl/detail/compile_time_kernel_info.hpp> |
13 | 13 | #include <sycl/detail/helpers.hpp> |
14 | 14 | #include <sycl/detail/is_device_copyable.hpp> |
| 15 | +#include <sycl/detail/type_traits.hpp> |
15 | 16 | #include <sycl/ext/intel/experimental/fp_control_kernel_properties.hpp> |
16 | 17 | #include <sycl/ext/intel/experimental/kernel_execution_properties.hpp> |
| 18 | +#include <sycl/ext/oneapi/experimental/cluster_group_prop.hpp> |
| 19 | +#include <sycl/ext/oneapi/experimental/graph.hpp> |
| 20 | +#include <sycl/ext/oneapi/experimental/use_root_sync_prop.hpp> |
17 | 21 | #include <sycl/ext/oneapi/experimental/virtual_functions.hpp> |
18 | 22 | #include <sycl/ext/oneapi/kernel_properties/properties.hpp> |
19 | 23 | #include <sycl/ext/oneapi/work_group_scratch_memory.hpp> |
@@ -253,23 +257,130 @@ struct KernelWrapper< |
253 | 257 | } |
254 | 258 | }; // KernelWrapper struct |
255 | 259 |
|
256 | | -struct KernelLaunchPropertyWrapper { |
257 | | - template <typename KernelName, typename PropertyProcessor, |
258 | | - typename KernelType> |
259 | | - static void parseProperties([[maybe_unused]] PropertyProcessor h, |
260 | | - [[maybe_unused]] const KernelType &KernelFunc) { |
261 | | -#ifndef __SYCL_DEVICE_ONLY__ |
262 | | - // If there are properties provided by get method then process them. |
263 | | - if constexpr (ext::oneapi::experimental::detail:: |
264 | | - HasKernelPropertiesGetMethod<const KernelType &>::value) { |
265 | | - |
266 | | - h->template processProperties< |
267 | | - detail::CompileTimeKernelInfo<KernelName>.IsESIMD>( |
268 | | - KernelFunc.get(ext::oneapi::experimental::properties_tag{})); |
| 260 | +// This namespace encapsulates everything related to parsing kernel launch |
| 261 | +// properties. |
| 262 | +inline namespace kernel_launch_properties_v1 { |
| 263 | + |
| 264 | +template <typename key, typename = void> struct MarshalledProperty; |
| 265 | + |
| 266 | +// Generic implementation for runtime properties. |
| 267 | +template <typename PropertyTy> |
| 268 | +struct MarshalledProperty< |
| 269 | + PropertyTy, |
| 270 | + std::enable_if_t<!std::is_empty_v<PropertyTy> && |
| 271 | + std::is_same_v<PropertyTy, typename PropertyTy::key_t>>> { |
| 272 | + std::optional<PropertyTy> MProperty; |
| 273 | + |
| 274 | + template <typename InputPropertyTy> |
| 275 | + MarshalledProperty(const InputPropertyTy &Props) { |
| 276 | + (void)Props; |
| 277 | + if constexpr (InputPropertyTy::template has_property<PropertyTy>()) |
| 278 | + MProperty = Props.template get_property<PropertyTy>(); |
| 279 | + } |
| 280 | + |
| 281 | + MarshalledProperty() = default; |
| 282 | +}; |
| 283 | + |
| 284 | +// Generic implementation for properties with non-template value_t. |
| 285 | +template <typename PropertyTy> |
| 286 | +struct MarshalledProperty<PropertyTy, |
| 287 | + std::void_t<typename PropertyTy::value_t>> { |
| 288 | + bool MPresent = false; |
| 289 | + |
| 290 | + template <typename InputPropertyTy> |
| 291 | + MarshalledProperty(const InputPropertyTy &) { |
| 292 | + using namespace sycl::ext::oneapi::experimental; |
| 293 | + MPresent = InputPropertyTy::template has_property< |
| 294 | + sycl::ext::oneapi::experimental::use_root_sync_key>(); |
| 295 | + } |
| 296 | + |
| 297 | + MarshalledProperty() = default; |
| 298 | +}; |
| 299 | + |
| 300 | +// Specialization for work group progress property. |
| 301 | +template <typename PropertyTy> |
| 302 | +struct MarshalledProperty< |
| 303 | + PropertyTy, |
| 304 | + std::enable_if_t<check_type_in_v< |
| 305 | + PropertyTy, sycl::ext::oneapi::experimental::work_group_progress_key, |
| 306 | + sycl::ext::oneapi::experimental::sub_group_progress_key, |
| 307 | + sycl::ext::oneapi::experimental::work_item_progress_key>>> { |
| 308 | + |
| 309 | + using forward_progress_guarantee = |
| 310 | + sycl::ext::oneapi::experimental::forward_progress_guarantee; |
| 311 | + using execution_scope = sycl::ext::oneapi::experimental::execution_scope; |
| 312 | + |
| 313 | + std::optional<forward_progress_guarantee> MFPGuarantee; |
| 314 | + std::optional<execution_scope> MFPCoordinationScope; |
| 315 | + |
| 316 | + template <typename InputPropertyTy> |
| 317 | + MarshalledProperty(const InputPropertyTy &Props) { |
| 318 | + (void)Props; |
| 319 | + |
| 320 | + if constexpr (InputPropertyTy::template has_property<PropertyTy>()) { |
| 321 | + MFPGuarantee = Props.template get_property<PropertyTy>().guarantee; |
| 322 | + MFPCoordinationScope = |
| 323 | + Props.template get_property<PropertyTy>().coordinationScope; |
269 | 324 | } |
270 | | -#endif |
271 | 325 | } |
272 | | -}; // KernelLaunchPropertyWrapper struct |
| 326 | + |
| 327 | + MarshalledProperty() = default; |
| 328 | +}; |
| 329 | + |
| 330 | +template <typename... keys> struct PropsHolder : MarshalledProperty<keys>... { |
| 331 | + bool MEmpty = true; |
| 332 | + |
| 333 | + template <typename PropertiesT, |
| 334 | + class = typename std::enable_if_t< |
| 335 | + ext::oneapi::experimental::is_property_list_v<PropertiesT>>> |
| 336 | + PropsHolder(PropertiesT Props) |
| 337 | + : MarshalledProperty<keys>(Props)..., |
| 338 | + MEmpty(((!PropertiesT::template has_property<keys>() && ...))) {} |
| 339 | + |
| 340 | + PropsHolder() = default; |
| 341 | + |
| 342 | + constexpr bool isEmpty() const { return MEmpty; } |
| 343 | + |
| 344 | + template <typename PropertyCastKey> constexpr auto get() const { |
| 345 | + return static_cast<const MarshalledProperty<PropertyCastKey> *>(this); |
| 346 | + } |
| 347 | +}; |
| 348 | + |
| 349 | +using KernelPropertyHolderStructTy = |
| 350 | + PropsHolder<sycl::ext::oneapi::experimental::work_group_scratch_size, |
| 351 | + sycl::ext::intel::experimental::cache_config_key, |
| 352 | + sycl::ext::oneapi::experimental::use_root_sync_key, |
| 353 | + sycl::ext::oneapi::experimental::work_group_progress_key, |
| 354 | + sycl::ext::oneapi::experimental::sub_group_progress_key, |
| 355 | + sycl::ext::oneapi::experimental::work_item_progress_key, |
| 356 | + sycl::ext::oneapi::experimental::cuda::cluster_size_key<1>, |
| 357 | + sycl::ext::oneapi::experimental::cuda::cluster_size_key<2>, |
| 358 | + sycl::ext::oneapi::experimental::cuda::cluster_size_key<3>>; |
| 359 | + |
| 360 | +/// Note: it is important that this function *does not* depend on kernel |
| 361 | +/// name or kernel type, because then it will be instantiated for every |
| 362 | +/// kernel, even though body of those instantiated functions could be almost |
| 363 | +/// the same, thus unnecessary increasing compilation time. |
| 364 | +template <bool IsESIMDKernel = false, typename PropertiesT, |
| 365 | + class = typename std::enable_if_t< |
| 366 | + ext::oneapi::experimental::is_property_list_v<PropertiesT>>> |
| 367 | +constexpr KernelPropertyHolderStructTy |
| 368 | +extractKernelProperties(PropertiesT Props) { |
| 369 | + static_assert( |
| 370 | + !PropertiesT::template has_property< |
| 371 | + sycl::ext::intel::experimental::fp_control_key>() || |
| 372 | + (PropertiesT::template has_property< |
| 373 | + sycl::ext::intel::experimental::fp_control_key>() && |
| 374 | + IsESIMDKernel), |
| 375 | + "Floating point control property is supported for ESIMD kernels only."); |
| 376 | + static_assert( |
| 377 | + !PropertiesT::template has_property< |
| 378 | + sycl::ext::oneapi::experimental::indirectly_callable_key>(), |
| 379 | + "indirectly_callable property cannot be applied to SYCL kernels"); |
| 380 | + |
| 381 | + return KernelPropertyHolderStructTy(Props); |
| 382 | +} |
| 383 | +} // namespace kernel_launch_properties_v1 |
273 | 384 |
|
274 | 385 | } // namespace detail |
275 | 386 | } // namespace _V1 |
|
0 commit comments