From d83b1fd6e0c23158e43b9fcdaa93ea2378fe0867 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Fri, 18 Jul 2025 05:48:28 +0200 Subject: [PATCH 01/11] [SYCL][NFC] Move Kernel Wrapper outside the handler. --- sycl/include/sycl/handler.hpp | 277 +---------------- sycl/include/sycl/kernel_helper.hpp | 290 ++++++++++++++++++ .../handler/unnamed-lambda-negative.cpp | 14 +- .../check_device_code/device_has_kernel.cpp | 4 +- .../esimd/slm_init_specconst_size.cpp | 2 +- .../extensions/annotated_arg/unaliased.cpp | 2 +- .../properties_kernel_sub_group_size.cpp | 25 +- .../properties_kernel_work_group_size.cpp | 30 +- ...properties_kernel_work_group_size_hint.cpp | 27 +- sycl/test/check_device_code/no_offset.cpp | 4 +- sycl/test/esimd/global_var.cpp | 2 +- .../non_esimd_kernel_fp_control.cpp | 4 +- .../properties/properties_kernel_negative.cpp | 2 +- .../include_deps/sycl_detail_core.hpp.cpp | 1 + .../virtual-functions/properties-negative.cpp | 6 +- 15 files changed, 367 insertions(+), 323 deletions(-) create mode 100644 sycl/include/sycl/kernel_helper.hpp diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 70429152f0ea1..2ee8f5ab5724b 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -47,6 +47,7 @@ #include #include #include +#include #include #include #include @@ -264,28 +265,6 @@ __SYCL_EXPORT void *getValueFromDynamicParameter( ext::oneapi::experimental::detail::dynamic_parameter_base &DynamicParamBase); -// Helper for merging properties with ones defined in an optional kernel functor -// getter. -template -struct GetMergedKernelProperties { - using type = PropertiesT; -}; -template -struct GetMergedKernelProperties< - KernelType, PropertiesT, - std::enable_if_t::value>> { - using get_method_properties = - typename ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod< - KernelType>::properties_t; - static_assert( - ext::oneapi::experimental::is_property_list::value, - "get(sycl::ext::oneapi::experimental::properties_tag) member in kernel " - "functor class must return a valid property list."); - using type = ext::oneapi::experimental::detail::merged_properties_t< - PropertiesT, get_method_properties>; -}; - template class RoundedRangeIDGenerator { id Id; id InitId; @@ -923,6 +902,7 @@ class __SYCL_EXPORT handler { } } +public: /// Process runtime kernel properties. /// /// Stores information about kernel properties into the handler. @@ -988,36 +968,7 @@ class __SYCL_EXPORT handler { checkAndSetClusterRange(Props); } - /// Process kernel properties. - /// - /// Stores information about kernel properties into the handler. - /// - /// 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 - /// the same, thus unnecessary increasing compilation time. - template < - bool IsESIMDKernel, - typename PropertiesT = ext::oneapi::experimental::empty_properties_t> - void processProperties(PropertiesT Props) { - static_assert( - ext::oneapi::experimental::is_property_list::value, - "Template type is not a property list."); - static_assert( - !PropertiesT::template has_property< - sycl::ext::intel::experimental::fp_control_key>() || - (PropertiesT::template has_property< - sycl::ext::intel::experimental::fp_control_key>() && - IsESIMDKernel), - "Floating point control property is supported for ESIMD kernels only."); - static_assert( - !PropertiesT::template has_property< - sycl::ext::oneapi::experimental::indirectly_callable_key>(), - "indirectly_callable property cannot be applied to SYCL kernels"); - - processLaunchProperties(Props); - } - +private: /// Checks whether it is possible to copy the source shape to the destination /// shape(the shapes are described by the accessor ranges) by using /// copying by regions of memory and not copying element by element @@ -1320,7 +1271,8 @@ class __SYCL_EXPORT handler { decltype(Wrapper), NameWT>; KernelWrapper::wrap(this, Wrapper); + TransformedArgType, decltype(this), + PropertiesT>::wrap(this, Wrapper); #ifndef __SYCL_DEVICE_ONLY__ constexpr detail::string_view Name{detail::getKernelName()}; verifyUsedKernelBundleInternal(Name); @@ -1345,7 +1297,7 @@ class __SYCL_EXPORT handler { // If parallel_for range rounding is forced then only range rounded // kernel is generated KernelWrapper::wrap(this, KernelFunc); + decltype(this), PropertiesT>::wrap(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ constexpr detail::string_view Name{detail::getKernelName()}; @@ -1412,215 +1364,6 @@ class __SYCL_EXPORT handler { #endif } -#ifdef SYCL_LANGUAGE_VERSION -#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS -#define __SYCL_KERNEL_ATTR__ [[clang::sycl_kernel_entry_point(KernelName)]] -#else -#define __SYCL_KERNEL_ATTR__ [[clang::sycl_kernel]] -#endif // __INTEL_SYCL_USE_INTEGRATION_HEADERS -#else -#define __SYCL_KERNEL_ATTR__ -#endif // SYCL_LANGUAGE_VERSION - - // NOTE: the name of this function - "kernel_single_task" - is used by the - // Front End to determine kernel invocation kind. - template -#ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_attributes_function( - "sycl-single-task", - ext::oneapi::experimental::detail::PropertyMetaInfo::name..., - nullptr, - ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] -#endif - - __SYCL_KERNEL_ATTR__ static void - kernel_single_task(const KernelType &KernelFunc) { -#ifdef __SYCL_DEVICE_ONLY__ - KernelFunc(); -#else - (void)KernelFunc; -#endif - } - - // NOTE: the name of this function - "kernel_single_task" - is used by the - // Front End to determine kernel invocation kind. - template -#ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_attributes_function( - "sycl-single-task", - ext::oneapi::experimental::detail::PropertyMetaInfo::name..., - nullptr, - ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] -#endif - __SYCL_KERNEL_ATTR__ static void - kernel_single_task(const KernelType &KernelFunc, kernel_handler KH) { -#ifdef __SYCL_DEVICE_ONLY__ - KernelFunc(KH); -#else - (void)KernelFunc; - (void)KH; -#endif - } - - // NOTE: the name of these functions - "kernel_parallel_for" - are used by the - // Front End to determine kernel invocation kind. - template -#ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_attributes_function( - ext::oneapi::experimental::detail::PropertyMetaInfo::name..., - ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] -#endif - __SYCL_KERNEL_ATTR__ static void - kernel_parallel_for(const KernelType &KernelFunc) { -#ifdef __SYCL_DEVICE_ONLY__ - KernelFunc(detail::Builder::getElement(detail::declptr())); -#else - (void)KernelFunc; -#endif - } - - // NOTE: the name of these functions - "kernel_parallel_for" - are used by the - // Front End to determine kernel invocation kind. - template -#ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_attributes_function( - ext::oneapi::experimental::detail::PropertyMetaInfo::name..., - ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] -#endif - __SYCL_KERNEL_ATTR__ static void - kernel_parallel_for(const KernelType &KernelFunc, kernel_handler KH) { -#ifdef __SYCL_DEVICE_ONLY__ - KernelFunc(detail::Builder::getElement(detail::declptr()), KH); -#else - (void)KernelFunc; - (void)KH; -#endif - } - - // NOTE: the name of this function - "kernel_parallel_for_work_group" - is - // used by the Front End to determine kernel invocation kind. - template -#ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_attributes_function( - ext::oneapi::experimental::detail::PropertyMetaInfo::name..., - ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] -#endif - __SYCL_KERNEL_ATTR__ static void - kernel_parallel_for_work_group(const KernelType &KernelFunc) { -#ifdef __SYCL_DEVICE_ONLY__ - KernelFunc(detail::Builder::getElement(detail::declptr())); -#else - (void)KernelFunc; -#endif - } - - // NOTE: the name of this function - "kernel_parallel_for_work_group" - is - // used by the Front End to determine kernel invocation kind. - template -#ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_attributes_function( - ext::oneapi::experimental::detail::PropertyMetaInfo::name..., - ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] -#endif - __SYCL_KERNEL_ATTR__ static void - kernel_parallel_for_work_group(const KernelType &KernelFunc, - kernel_handler KH) { -#ifdef __SYCL_DEVICE_ONLY__ - KernelFunc(detail::Builder::getElement(detail::declptr()), KH); -#else - (void)KernelFunc; - (void)KH; -#endif - } - - // The KernelWrapper below has two purposes. - // - // First, from SYCL 2020, Table 129 (Member functions of the `handler ` class) - // > The callable ... can optionally take a `kernel_handler` ... in - // > which case the SYCL runtime will construct an instance of - // > `kernel_handler` and pass it to the callable. - // - // Note: "..." due to slight wording variability between - // single_task/parallel_for (e.g. only parameter vs last). This helper class - // calls `kernel_*` entry points (both hardcoded names known to FE and special - // device-specific entry point attributes) with proper arguments (with/without - // `kernel_handler` argument, depending on the signature of the SYCL kernel - // function). - // - // Second, it performs a few checks and some properties processing (including - // the one provided via `sycl_ext_oneapi_kernel_properties` extension by - // embedding them into the kernel's type). - - enum class WrapAs { single_task, parallel_for, parallel_for_work_group }; - - template < - WrapAs WrapAsVal, typename KernelName, typename KernelType, - typename ElementType, - typename PropertiesT = ext::oneapi::experimental::empty_properties_t, - typename MergedPropertiesT = typename detail::GetMergedKernelProperties< - KernelType, PropertiesT>::type> - struct KernelWrapper; - template - struct KernelWrapper< - WrapAsVal, KernelName, KernelType, ElementType, PropertiesT, - ext::oneapi::experimental::detail::properties_t> { - static void wrap(handler *h, const KernelType &KernelFunc) { -#ifdef __SYCL_DEVICE_ONLY__ - detail::CheckDeviceCopyable(); -#else - // If there are properties provided by get method then process them. - if constexpr (ext::oneapi::experimental::detail:: - HasKernelPropertiesGetMethod< - const KernelType &>::value) { - h->processProperties()>( - KernelFunc.get(ext::oneapi::experimental::properties_tag{})); - } -#endif - // Note: the static_assert below need to be run on both the host and the - // device ends to avoid test issues, so don't put it into the #ifdef - // __SYCL_DEVICE_ONLY__ directive above print out diagnostic message if - // the kernel functor has a get(properties_tag) member, but it's not const - static_assert( - (ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod< - const KernelType &>::value) || - !(ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod< - KernelType>::value), - "get(sycl::ext::oneapi::experimental::properties_tag) member in " - "kernel functor class must be declared as a const member function"); - auto L = [&](auto &&...args) { - if constexpr (WrapAsVal == WrapAs::single_task) { - h->kernel_single_task( - std::forward(args)...); - } else if constexpr (WrapAsVal == WrapAs::parallel_for) { - h->kernel_parallel_for( - std::forward(args)...); - } else if constexpr (WrapAsVal == WrapAs::parallel_for_work_group) { - h->kernel_parallel_for_work_group( - std::forward(args)...); - } else { - // Always false, but template-dependent. Can't compare `WrapAsVal` - // with itself because of `-Wtautological-compare` warning. - static_assert(!std::is_same_v, - "Unexpected WrapAsVal"); - } - }; - if constexpr (detail::KernelLambdaHasKernelHandlerArgT< - KernelType, ElementType>::value) { - kernel_handler KH; - L(KernelFunc, KH); - } else { - L(KernelFunc); - } - } - }; - template < WrapAs WrapAsVal, typename KernelName, typename ElementType = void, int Dims = 1, bool SetNumWorkGroups = false, @@ -1633,8 +1376,8 @@ class __SYCL_EXPORT handler { using NameT = typename detail::get_kernel_name_t::name; (void)Props; - KernelWrapper::wrap( - this, KernelFunc); + KernelWrapper::wrap(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ if constexpr (WrapAsVal == WrapAs::single_task) { throwOnKernelParameterMisuse(); @@ -1674,8 +1417,8 @@ class __SYCL_EXPORT handler { typename detail::get_kernel_name_t::name; (void)Props; (void)Kernel; - KernelWrapper::wrap( - this, KernelFunc); + KernelWrapper::wrap(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ if constexpr (WrapAsVal == WrapAs::single_task) { throwOnKernelParameterMisuse(); diff --git a/sycl/include/sycl/kernel_helper.hpp b/sycl/include/sycl/kernel_helper.hpp new file mode 100644 index 0000000000000..7c4ffe38e8959 --- /dev/null +++ b/sycl/include/sycl/kernel_helper.hpp @@ -0,0 +1,290 @@ +//==-------- kernel_helper.hpp --- SYCL kernel helper utilities ---------==// +// +// 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 +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +namespace sycl { +inline namespace _V1 { + +namespace detail { +// Helper for merging properties with ones defined in an optional kernel functor +// getter. +template +struct GetMergedKernelProperties { + using type = PropertiesT; +}; +template +struct GetMergedKernelProperties< + KernelType, PropertiesT, + std::enable_if_t::value>> { + using get_method_properties = + typename ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod< + KernelType>::properties_t; + static_assert( + ext::oneapi::experimental::is_property_list::value, + "get(sycl::ext::oneapi::experimental::properties_tag) member in kernel " + "functor class must return a valid property list."); + using type = ext::oneapi::experimental::detail::merged_properties_t< + PropertiesT, get_method_properties>; +}; +} // namespace detail + +#ifdef SYCL_LANGUAGE_VERSION +#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS +#define __SYCL_KERNEL_ATTR__ [[clang::sycl_kernel_entry_point(KernelName)]] +#else +#define __SYCL_KERNEL_ATTR__ [[clang::sycl_kernel]] +#endif // __INTEL_SYCL_USE_INTEGRATION_HEADERS +#else +#define __SYCL_KERNEL_ATTR__ +#endif // SYCL_LANGUAGE_VERSION + +// NOTE: the name of this function - "kernel_single_task" - is used by the +// Front End to determine kernel invocation kind. +template +#ifdef __SYCL_DEVICE_ONLY__ +[[__sycl_detail__::add_ir_attributes_function( + "sycl-single-task", + ext::oneapi::experimental::detail::PropertyMetaInfo::name..., + nullptr, + ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] +#endif + +__SYCL_KERNEL_ATTR__ static void +kernel_single_task(const KernelType &KernelFunc) { +#ifdef __SYCL_DEVICE_ONLY__ + KernelFunc(); +#else + (void)KernelFunc; +#endif +} + +// NOTE: the name of this function - "kernel_single_task" - is used by the +// Front End to determine kernel invocation kind. +template +#ifdef __SYCL_DEVICE_ONLY__ +[[__sycl_detail__::add_ir_attributes_function( + "sycl-single-task", + ext::oneapi::experimental::detail::PropertyMetaInfo::name..., + nullptr, + ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] +#endif +__SYCL_KERNEL_ATTR__ static void +kernel_single_task(const KernelType &KernelFunc, kernel_handler KH) { +#ifdef __SYCL_DEVICE_ONLY__ + KernelFunc(KH); +#else + (void)KernelFunc; + (void)KH; +#endif +} + +// NOTE: the name of these functions - "kernel_parallel_for" - are used by the +// Front End to determine kernel invocation kind. +template +#ifdef __SYCL_DEVICE_ONLY__ +[[__sycl_detail__::add_ir_attributes_function( + ext::oneapi::experimental::detail::PropertyMetaInfo::name..., + ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] +#endif +__SYCL_KERNEL_ATTR__ static void +kernel_parallel_for(const KernelType &KernelFunc) { +#ifdef __SYCL_DEVICE_ONLY__ + KernelFunc(detail::Builder::getElement(detail::declptr())); +#else + (void)KernelFunc; +#endif +} + +// NOTE: the name of these functions - "kernel_parallel_for" - are used by the +// Front End to determine kernel invocation kind. +template +#ifdef __SYCL_DEVICE_ONLY__ +[[__sycl_detail__::add_ir_attributes_function( + ext::oneapi::experimental::detail::PropertyMetaInfo::name..., + ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] +#endif +__SYCL_KERNEL_ATTR__ static void +kernel_parallel_for(const KernelType &KernelFunc, kernel_handler KH) { +#ifdef __SYCL_DEVICE_ONLY__ + KernelFunc(detail::Builder::getElement(detail::declptr()), KH); +#else + (void)KernelFunc; + (void)KH; +#endif +} + +// NOTE: the name of this function - "kernel_parallel_for_work_group" - is +// used by the Front End to determine kernel invocation kind. +template +#ifdef __SYCL_DEVICE_ONLY__ +[[__sycl_detail__::add_ir_attributes_function( + ext::oneapi::experimental::detail::PropertyMetaInfo::name..., + ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] +#endif +__SYCL_KERNEL_ATTR__ static void +kernel_parallel_for_work_group(const KernelType &KernelFunc) { +#ifdef __SYCL_DEVICE_ONLY__ + KernelFunc(detail::Builder::getElement(detail::declptr())); +#else + (void)KernelFunc; +#endif +} + +// NOTE: the name of this function - "kernel_parallel_for_work_group" - is +// used by the Front End to determine kernel invocation kind. +template +#ifdef __SYCL_DEVICE_ONLY__ +[[__sycl_detail__::add_ir_attributes_function( + ext::oneapi::experimental::detail::PropertyMetaInfo::name..., + ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] +#endif +__SYCL_KERNEL_ATTR__ static void +kernel_parallel_for_work_group(const KernelType &KernelFunc, + kernel_handler KH) { +#ifdef __SYCL_DEVICE_ONLY__ + KernelFunc(detail::Builder::getElement(detail::declptr()), KH); +#else + (void)KernelFunc; + (void)KH; +#endif +} + +/// Process kernel properties. +/// +/// Stores information about kernel properties into the handler. +/// +/// 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 +/// the same, thus unnecessary increasing compilation time. +template +void processProperties([[maybe_unused]] PropertiesT Props) { + static_assert(ext::oneapi::experimental::is_property_list::value, + "Template type is not a property list."); + static_assert( + !PropertiesT::template has_property< + sycl::ext::intel::experimental::fp_control_key>() || + (PropertiesT::template has_property< + sycl::ext::intel::experimental::fp_control_key>() && + IsESIMDKernel), + "Floating point control property is supported for ESIMD kernels only."); + static_assert( + !PropertiesT::template has_property< + sycl::ext::oneapi::experimental::indirectly_callable_key>(), + "indirectly_callable property cannot be applied to SYCL kernels"); +} + +// The KernelWrapper below has two purposes. +// +// First, from SYCL 2020, Table 129 (Member functions of the `handler ` class) +// > The callable ... can optionally take a `kernel_handler` ... in +// > which case the SYCL runtime will construct an instance of +// > `kernel_handler` and pass it to the callable. +// +// Note: "..." due to slight wording variability between +// single_task/parallel_for (e.g. only parameter vs last). This helper class +// calls `kernel_*` entry points (both hardcoded names known to FE and special +// device-specific entry point attributes) with proper arguments (with/without +// `kernel_handler` argument, depending on the signature of the SYCL kernel +// function). +// +// Second, it performs a few checks and some properties processing (including +// the one provided via `sycl_ext_oneapi_kernel_properties` extension by +// embedding them into the kernel's type). + +enum class WrapAs { single_task, parallel_for, parallel_for_work_group }; + +template ::type> +struct KernelWrapper; +template +struct KernelWrapper< + WrapAsVal, KernelName, KernelType, ElementType, PropertyProcessor, + PropertiesT, + ext::oneapi::experimental::detail::properties_t> { + static void wrap([[maybe_unused]] PropertyProcessor h, + [[maybe_unused]] const KernelType &KernelFunc) { +#ifdef __SYCL_DEVICE_ONLY__ + detail::CheckDeviceCopyable(); +#else + // If there are properties provided by get method then process them. + if constexpr (ext::oneapi::experimental::detail:: + HasKernelPropertiesGetMethod::value) { + + processProperties()>( + KernelFunc.get(ext::oneapi::experimental::properties_tag{})); + h->processLaunchProperties( + KernelFunc.get(ext::oneapi::experimental::properties_tag{})); + } +#endif + // Note: the static_assert below need to be run on both the host and the + // device ends to avoid test issues, so don't put it into the #ifdef + // __SYCL_DEVICE_ONLY__ directive above print out diagnostic message if + // the kernel functor has a get(properties_tag) member, but it's not const + static_assert( + (ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod< + const KernelType &>::value) || + !(ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod< + KernelType>::value), + "get(sycl::ext::oneapi::experimental::properties_tag) member in " + "kernel functor class must be declared as a const member function"); + auto L = [&](auto &&...args) { + if constexpr (WrapAsVal == WrapAs::single_task) { + kernel_single_task( + std::forward(args)...); + } else if constexpr (WrapAsVal == WrapAs::parallel_for) { + kernel_parallel_for( + std::forward(args)...); + } else if constexpr (WrapAsVal == WrapAs::parallel_for_work_group) { + kernel_parallel_for_work_group( + std::forward(args)...); + } else { + // Always false, but template-dependent. Can't compare `WrapAsVal` + // with itself because of `-Wtautological-compare` warning. + static_assert(!std::is_same_v, + "Unexpected WrapAsVal"); + } + }; + if constexpr (detail::KernelLambdaHasKernelHandlerArgT< + KernelType, ElementType>::value) { + kernel_handler KH; + L(KernelFunc, KH); + } else { + L(KernelFunc); + } + } +}; // KernelWrapper struct + +} // namespace _V1 +} // namespace sycl \ No newline at end of file diff --git a/sycl/test/basic_tests/handler/unnamed-lambda-negative.cpp b/sycl/test/basic_tests/handler/unnamed-lambda-negative.cpp index 30144d2366c57..c17a60db250a0 100644 --- a/sycl/test/basic_tests/handler/unnamed-lambda-negative.cpp +++ b/sycl/test/basic_tests/handler/unnamed-lambda-negative.cpp @@ -8,26 +8,26 @@ int main() { sycl::queue q; - // expected-error-re@sycl/handler.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/kernel_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} q.single_task([=](){}); - // expected-error-re@sycl/handler.hpp:* {{unnamed type 'sycl::detail::RoundedRangeKernel<{{.*}}>' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} - // expected-error-re@sycl/handler.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/kernel_helper.hpp:* {{unnamed type 'sycl::detail::RoundedRangeKernel<{{.*}}>' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/kernel_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} q.parallel_for(sycl::range<1>{1}, [=](sycl::item<1>) {}); q.submit([&](sycl::handler &cgh) { - // expected-error-re@sycl/handler.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/kernel_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} cgh.single_task([=](){}); - // expected-error-re@sycl/handler.hpp:* {{unnamed type 'sycl::detail::RoundedRangeKernel<{{.*}}>' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} - // expected-error-re@sycl/handler.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/kernel_helper.hpp:* {{unnamed type 'sycl::detail::RoundedRangeKernel<{{.*}}>' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/kernel_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} cgh.parallel_for(sycl::range<1>{1}, [=](sycl::item<1>) {}); - // expected-error-re@sycl/handler.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/kernel_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} cgh.parallel_for_work_group(sycl::range<1>{1}, [=](sycl::group<1>) {}); }); diff --git a/sycl/test/check_device_code/device_has_kernel.cpp b/sycl/test/check_device_code/device_has_kernel.cpp index 191dfe49d2030..f25bc27eccce8 100644 --- a/sycl/test/check_device_code/device_has_kernel.cpp +++ b/sycl/test/check_device_code/device_has_kernel.cpp @@ -10,7 +10,7 @@ using namespace sycl; queue q; -// CHECK-ASPECTS: define weak_odr dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] {{.*}} +// CHECK-ASPECTS: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] {{.*}} // CHECK-ASPECTS: define {{.*}}spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] // CHECK-ASPECTS-SAME: !sycl_used_aspects ![[ASPECTS1]] @@ -53,7 +53,7 @@ void foo() { q.submit([&](handler &h) { KernelFunctor f1; h.single_task(f1); - // CHECK-ASPECTS: define weak_odr dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]] + // CHECK-ASPECTS: define dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]] h.single_task( []() [[sycl::device_has(sycl::aspect::gpu)]] {}); }); diff --git a/sycl/test/check_device_code/esimd/slm_init_specconst_size.cpp b/sycl/test/check_device_code/esimd/slm_init_specconst_size.cpp index a2f76ff0981ab..1991616c3f698 100644 --- a/sycl/test/check_device_code/esimd/slm_init_specconst_size.cpp +++ b/sycl/test/check_device_code/esimd/slm_init_specconst_size.cpp @@ -21,7 +21,7 @@ int main() { [=](sycl::kernel_handler kh) SYCL_ESIMD_KERNEL { slm_init(kh.get_specialization_constant()); }); - // CHECK: define weak_odr dso_local spir_kernel void @{{.*}}() local_unnamed_addr #1 + // CHECK: define dso_local spir_kernel void @{{.*}}() local_unnamed_addr #1 }); } diff --git a/sycl/test/check_device_code/extensions/annotated_arg/unaliased.cpp b/sycl/test/check_device_code/extensions/annotated_arg/unaliased.cpp index 028799fbe5214..e603b79d4e5e2 100644 --- a/sycl/test/check_device_code/extensions/annotated_arg/unaliased.cpp +++ b/sycl/test/check_device_code/extensions/annotated_arg/unaliased.cpp @@ -19,4 +19,4 @@ int main() { return 0; } -// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef align 4 "sycl-unaliased" +// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef writeonly align 4 captures(none) "sycl-unaliased" diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp index ad81d1db1fe0b..1a49166819e42 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp @@ -79,28 +79,28 @@ int main() { Q.parallel_for(R3, {Ev}, Props, Redu1, [](sycl::id<3>, auto &) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel27(){{.*}} #[[SGSizeAttr2]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel27(){{.*}} #[[SGSizeAttr4:[0-9]+]] Q.parallel_for(NDR1, Props, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel28(){{.*}} #[[SGSizeAttr2]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel28(){{.*}} #[[SGSizeAttr4]] Q.parallel_for(NDR1, Ev, Props, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel29(){{.*}} #[[SGSizeAttr2]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel29(){{.*}} #[[SGSizeAttr4]] Q.parallel_for(NDR1, {Ev}, Props, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel30(){{.*}} #[[SGSizeAttr2]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel30(){{.*}} #[[SGSizeAttr4]] Q.parallel_for(NDR2, Props, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel31(){{.*}} #[[SGSizeAttr2]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel31(){{.*}} #[[SGSizeAttr4]] Q.parallel_for(NDR2, Ev, Props, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel32(){{.*}} #[[SGSizeAttr2]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel32(){{.*}} #[[SGSizeAttr4]] Q.parallel_for(NDR2, {Ev}, Props, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel33(){{.*}} #[[SGSizeAttr2]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel33(){{.*}} #[[SGSizeAttr4]] Q.parallel_for(NDR3, Props, [](sycl::nd_item<3>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel34(){{.*}} #[[SGSizeAttr2]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel34(){{.*}} #[[SGSizeAttr4]] Q.parallel_for(NDR3, Ev, Props, [](sycl::nd_item<3>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel35(){{.*}} #[[SGSizeAttr2]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel35(){{.*}} #[[SGSizeAttr4]] Q.parallel_for(NDR3, {Ev}, Props, [](sycl::nd_item<3>) {}); @@ -202,17 +202,17 @@ int main() { [](sycl::id<3>, auto &) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel63(){{.*}} #[[SGSizeAttr2]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel63(){{.*}} #[[SGSizeAttr4]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR1, Props, [](sycl::nd_item<1>) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel64(){{.*}} #[[SGSizeAttr2]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel64(){{.*}} #[[SGSizeAttr4]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR2, Props, [](sycl::nd_item<2>) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel65(){{.*}} #[[SGSizeAttr2]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel65(){{.*}} #[[SGSizeAttr4]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR3, Props, [](sycl::nd_item<3>) {}); @@ -278,3 +278,4 @@ int main() { // CHECK-IR: attributes #[[SGSizeAttr1]] = { {{.*}}"sycl-sub-group-size"="1" // CHECK-IR: attributes #[[SGSizeAttr2]] = { {{.*}}"sycl-sub-group-size"="1" // CHECK-IR: attributes #[[SGSizeAttr3]] = { {{.*}}"sycl-sub-group-size"="1" +// CHECK-IR: attributes #[[SGSizeAttr4]] = { {{.*}}"sycl-sub-group-size"="1" diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp index 63280fcc638f3..18546b5250be9 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp @@ -95,28 +95,28 @@ int main() { Q.parallel_for(R3, {Ev}, Props3, Redu1, [](sycl::id<3>, auto &) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel27(){{.*}} #[[WGSizeAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel27(){{.*}} #[[WGSizeAttr10:[0-9]+]] Q.parallel_for(NDR1, Props1, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel28(){{.*}} #[[WGSizeAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel28(){{.*}} #[[WGSizeAttr10]] Q.parallel_for(NDR1, Ev, Props1, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel29(){{.*}} #[[WGSizeAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel29(){{.*}} #[[WGSizeAttr10]] Q.parallel_for(NDR1, {Ev}, Props1, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel30(){{.*}} #[[WGSizeAttr5]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel30(){{.*}} #[[WGSizeAttr11:[0-9]+]] Q.parallel_for(NDR2, Props2, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel31(){{.*}} #[[WGSizeAttr5]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel31(){{.*}} #[[WGSizeAttr11]] Q.parallel_for(NDR2, Ev, Props2, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel32(){{.*}} #[[WGSizeAttr5]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel32(){{.*}} #[[WGSizeAttr11]] Q.parallel_for(NDR2, {Ev}, Props2, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel33(){{.*}} #[[WGSizeAttr6]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel33(){{.*}} #[[WGSizeAttr12:[0-9]+]] Q.parallel_for(NDR3, Props3, [](sycl::nd_item<3>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel34(){{.*}} #[[WGSizeAttr6]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel34(){{.*}} #[[WGSizeAttr12]] Q.parallel_for(NDR3, Ev, Props3, [](sycl::nd_item<3>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel35(){{.*}} #[[WGSizeAttr6]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel35(){{.*}} #[[WGSizeAttr12]] Q.parallel_for(NDR3, {Ev}, Props3, [](sycl::nd_item<3>) {}); @@ -218,17 +218,20 @@ int main() { [](sycl::id<3>, auto &) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel63(){{.*}} #[[WGSizeAttr4]] + // 15 + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel63(){{.*}} #[[WGSizeAttr10]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR1, Props1, [](sycl::nd_item<1>) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel64(){{.*}} #[[WGSizeAttr5]] + // 16 + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel64(){{.*}} #[[WGSizeAttr11]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR2, Props2, [](sycl::nd_item<2>) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel65(){{.*}} #[[WGSizeAttr6]] + // 17 + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel65(){{.*}} #[[WGSizeAttr12]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR3, Props3, [](sycl::nd_item<3>) {}); @@ -300,3 +303,6 @@ int main() { // CHECK-IR: attributes #[[WGSizeAttr7]] = { {{.*}}"sycl-work-group-size"="1" // CHECK-IR: attributes #[[WGSizeAttr8]] = { {{.*}}"sycl-work-group-size"="1,2" // CHECK-IR: attributes #[[WGSizeAttr9]] = { {{.*}}"sycl-work-group-size"="1,2,3" +// CHECK-IR: attributes #[[WGSizeAttr10]] = { {{.*}}"sycl-work-group-size"="1" +// CHECK-IR: attributes #[[WGSizeAttr11]] = { {{.*}}"sycl-work-group-size"="1,2" +// CHECK-IR: attributes #[[WGSizeAttr12]] = { {{.*}}"sycl-work-group-size"="1,2,3" \ No newline at end of file diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp index a0bae31ad8004..502b28244ba20 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp @@ -98,31 +98,31 @@ int main() { Q.parallel_for(R3, {Ev}, Props3, Redu1, [](sycl::id<3>, auto &) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel27(){{.*}} #[[WGSizeHintAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel27(){{.*}} #[[WGSizeHintAttr10:[0-9]+]] Q.parallel_for(NDR1, Props1, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel28(){{.*}} #[[WGSizeHintAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel28(){{.*}} #[[WGSizeHintAttr10]] Q.parallel_for(NDR1, Ev, Props1, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel29(){{.*}} #[[WGSizeHintAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel29(){{.*}} #[[WGSizeHintAttr10]] Q.parallel_for(NDR1, {Ev}, Props1, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel30(){{.*}} #[[WGSizeHintAttr5]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel30(){{.*}} #[[WGSizeHintAttr11:[0-9]+]] Q.parallel_for(NDR2, Props2, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel31(){{.*}} #[[WGSizeHintAttr5]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel31(){{.*}} #[[WGSizeHintAttr11]] Q.parallel_for(NDR2, Ev, Props2, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel32(){{.*}} #[[WGSizeHintAttr5]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel32(){{.*}} #[[WGSizeHintAttr11]] Q.parallel_for(NDR2, {Ev}, Props2, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel33(){{.*}} #[[WGSizeHintAttr6]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel33(){{.*}} #[[WGSizeHintAttr12:[0-9]+]] Q.parallel_for(NDR3, Props3, [](sycl::nd_item<3>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel34(){{.*}} #[[WGSizeHintAttr6]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel34(){{.*}} #[[WGSizeHintAttr12]] Q.parallel_for(NDR3, Ev, Props3, [](sycl::nd_item<3>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel35(){{.*}} #[[WGSizeHintAttr6]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel35(){{.*}} #[[WGSizeHintAttr12]] Q.parallel_for(NDR3, {Ev}, Props3, [](sycl::nd_item<3>) {}); @@ -227,17 +227,17 @@ int main() { [](sycl::id<3>, auto &) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel63(){{.*}} #[[WGSizeHintAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel63(){{.*}} #[[WGSizeHintAttr10]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR1, Props1, [](sycl::nd_item<1>) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel64(){{.*}} #[[WGSizeHintAttr5]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel64(){{.*}} #[[WGSizeHintAttr11]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR2, Props2, [](sycl::nd_item<2>) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel65(){{.*}} #[[WGSizeHintAttr6]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel65(){{.*}} #[[WGSizeHintAttr12]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR3, Props3, [](sycl::nd_item<3>) {}); @@ -309,3 +309,6 @@ int main() { // CHECK-IR: attributes #[[WGSizeHintAttr7]] = { {{.*}}"sycl-work-group-size-hint"="1" // CHECK-IR: attributes #[[WGSizeHintAttr8]] = { {{.*}}"sycl-work-group-size-hint"="1,2" // CHECK-IR: attributes #[[WGSizeHintAttr9]] = { {{.*}}"sycl-work-group-size-hint"="1,2,3" +// CHECK-IR: attributes #[[WGSizeHintAttr10]] = { {{.*}}"sycl-work-group-size-hint"="1" +// CHECK-IR: attributes #[[WGSizeHintAttr11]] = { {{.*}}"sycl-work-group-size-hint"="1,2" +// CHECK-IR: attributes #[[WGSizeHintAttr12]] = { {{.*}}"sycl-work-group-size-hint"="1,2,3" diff --git a/sycl/test/check_device_code/no_offset.cpp b/sycl/test/check_device_code/no_offset.cpp index f311c7dba40cf..d2c1a2ce43df8 100644 --- a/sycl/test/check_device_code/no_offset.cpp +++ b/sycl/test/check_device_code/no_offset.cpp @@ -15,7 +15,7 @@ int main() { sycl::ext::oneapi::accessor_property_list PL{sycl::ext::oneapi::no_offset, sycl::no_init}; sycl::accessor acc_a(a, cgh, sycl::write_only, PL); sycl::accessor acc_b{b, cgh, sycl::read_only}; - // CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlT_E_(ptr addrspace(1) {{.*}}, ptr addrspace(1) noundef readonly {{.*}}, ptr noundef byval(%"class.sycl::_V1::id") align 8 {{.*}}) + // CHECK: define dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlT_E_(ptr addrspace(1) {{.*}}, ptr addrspace(1) noundef readonly {{.*}}, ptr noundef readonly byval(%"class.sycl::_V1::id") align 8 {{.*}}) cgh.parallel_for(size, [=](auto i) { acc_a[i] = acc_b[i]; }); @@ -33,7 +33,7 @@ int main() { q.submit([&](sycl::handler &cgh) { sycl::accessor acc_a(a, cgh, sycl::write_only); sycl::accessor acc_b{b, cgh, sycl::read_only}; - // CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_EUlT_E_(ptr addrspace(1) {{.*}}, ptr noundef byval(%"class.sycl::_V1::id") align 8 {{.*}}, ptr addrspace(1) noundef readonly {{.*}}, ptr noundef byval(%"class.sycl::_V1::id") align 8 {{.*}}) + // CHECK: define dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_EUlT_E_(ptr addrspace(1) {{.*}}, ptr noundef readonly byval(%"class.sycl::_V1::id") align 8 {{.*}}, ptr noundef readonly byval(%"class.sycl::_V1::id") align 8 {{.*}}) cgh.parallel_for(size, [=](auto i) { acc_a[i] = acc_b[i]; }); diff --git a/sycl/test/esimd/global_var.cpp b/sycl/test/esimd/global_var.cpp index 7c10f04c8a618..61ff4c2f39d37 100644 --- a/sycl/test/esimd/global_var.cpp +++ b/sycl/test/esimd/global_var.cpp @@ -42,7 +42,7 @@ void kernel_call() { q.submit([&](sycl::handler &cgh) { cgh.parallel_for(nd_range<1>(1, 1), [=](nd_item<1> ndi) { - //expected-note@sycl/handler.hpp:* 2{{called by 'kernel_parallel_for}} + //expected-note@sycl/kernel_helper.hpp:* 2{{called by 'kernel_parallel_for}} //expected-error@+1{{ESIMD globals cannot be used in a SYCL context}} vc = 0; //expected-note@+1{{called by}} diff --git a/sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp b/sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp index 46d11eccdfe54..a4152802a1b17 100644 --- a/sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp +++ b/sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp @@ -20,7 +20,7 @@ struct ESIMDKernel { int main(void) { queue q; - // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.+}}: Floating point control property is supported for ESIMD kernels only.}} + // expected-error-re@sycl/kernel_helper.hpp:* {{static assertion failed due to requirement {{.+}}: Floating point control property is supported for ESIMD kernels only.}} syclex::properties properties7{ intelex::fp_control}; @@ -28,7 +28,7 @@ int main(void) { cgh.single_task(properties7, [=]() {}); }); - // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.+}}: Floating point control property is supported for ESIMD kernels only.}} + // expected-error-re@sycl/kernel_helper.hpp:* {{static assertion failed due to requirement {{.+}}: Floating point control property is supported for ESIMD kernels only.}} ESIMDKernel Kern; q.submit([&](handler &cgh) { cgh.parallel_for(range<1>(1), Kern); }); diff --git a/sycl/test/extensions/properties/properties_kernel_negative.cpp b/sycl/test/extensions/properties/properties_kernel_negative.cpp index 03f57711e0842..c88a649e024d4 100644 --- a/sycl/test/extensions/properties/properties_kernel_negative.cpp +++ b/sycl/test/extensions/properties/properties_kernel_negative.cpp @@ -399,7 +399,7 @@ struct TestKernelNoGetter { void check_non_const_getter_warning() { sycl::queue Q; - // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.+}}: get(sycl::ext::oneapi::experimental::properties_tag) member in kernel functor class must be declared as a const member function}} + // expected-error-re@sycl/kernel_helper.hpp:* {{static assertion failed due to requirement {{.+}}: get(sycl::ext::oneapi::experimental::properties_tag) member in kernel functor class must be declared as a const member function}} Q.single_task(TestKernelNonConstGetter()); // No error expected for kernel functor with a const get(properties_tag) diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp index 5c91a32bbdb04..5342539cce89f 100644 --- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp +++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp @@ -146,6 +146,7 @@ // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp // CHECK-NEXT: detail/sycl_local_mem_builtins.hpp +// CHECK-NEXT: kernel_helper.hpp // CHECK-NEXT: kernel.hpp // CHECK-NEXT: sampler.hpp // CHECK-NEXT: feature_test.hpp diff --git a/sycl/test/virtual-functions/properties-negative.cpp b/sycl/test/virtual-functions/properties-negative.cpp index b8e1b75f1d9a9..1a28969f86a8d 100644 --- a/sycl/test/virtual-functions/properties-negative.cpp +++ b/sycl/test/virtual-functions/properties-negative.cpp @@ -17,15 +17,15 @@ int main() { oneapi::properties props_int{oneapi::indirectly_callable_in}; oneapi::properties props_user{oneapi::indirectly_callable_in}; - // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}} + // expected-error-re@sycl/kernel_helper.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}} q.single_task(props_empty, [=]() {}); // When both "props_empty" and "props_void" are in use, we won't see the // static assert firing for the second one, because there will be only one // instantiation of handler::processProperties. q.single_task(props_void, [=]() {}); - // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}} + // expected-error-re@sycl/kernel_helper.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}} q.single_task(props_int, [=]() {}); - // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}} + // expected-error-re@sycl/kernel_helper.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}} q.single_task(props_user, [=]() {}); return 0; From cbb32ec308c7cde8dc6e8f5aee510cb97daf787d Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Fri, 18 Jul 2025 17:02:16 +0200 Subject: [PATCH 02/11] Don't pollute the global namespace --- sycl/include/sycl/handler.hpp | 2 +- sycl/include/sycl/kernel_helper.hpp | 163 ++++++++++++++-------------- 2 files changed, 85 insertions(+), 80 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 2ee8f5ab5724b..1c54196066747 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1892,7 +1892,7 @@ class __SYCL_EXPORT handler { using NameT = typename detail::get_kernel_name_t::name; (void)Kernel; - kernel_single_task(KernelFunc); + KernelWrapperSingletonFunc::kernel_single_task(KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ throwIfActionIsCreated(); constexpr detail::string_view Name{detail::getKernelName()}; diff --git a/sycl/include/sycl/kernel_helper.hpp b/sycl/include/sycl/kernel_helper.hpp index 7c4ffe38e8959..e99f618ab5894 100644 --- a/sycl/include/sycl/kernel_helper.hpp +++ b/sycl/include/sycl/kernel_helper.hpp @@ -47,6 +47,8 @@ struct GetMergedKernelProperties< }; } // namespace detail +struct KernelWrapperSingletonFunc { + #ifdef SYCL_LANGUAGE_VERSION #ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS #define __SYCL_KERNEL_ATTR__ [[clang::sycl_kernel_entry_point(KernelName)]] @@ -57,120 +59,121 @@ struct GetMergedKernelProperties< #define __SYCL_KERNEL_ATTR__ #endif // SYCL_LANGUAGE_VERSION -// NOTE: the name of this function - "kernel_single_task" - is used by the -// Front End to determine kernel invocation kind. -template + // NOTE: the name of this function - "kernel_single_task" - is used by the + // Front End to determine kernel invocation kind. + template #ifdef __SYCL_DEVICE_ONLY__ -[[__sycl_detail__::add_ir_attributes_function( - "sycl-single-task", - ext::oneapi::experimental::detail::PropertyMetaInfo::name..., - nullptr, - ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] + [[__sycl_detail__::add_ir_attributes_function( + "sycl-single-task", + ext::oneapi::experimental::detail::PropertyMetaInfo::name..., + nullptr, + ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif -__SYCL_KERNEL_ATTR__ static void -kernel_single_task(const KernelType &KernelFunc) { + __SYCL_KERNEL_ATTR__ static void + kernel_single_task(const KernelType &KernelFunc) { #ifdef __SYCL_DEVICE_ONLY__ - KernelFunc(); + KernelFunc(); #else - (void)KernelFunc; + (void)KernelFunc; #endif -} + } -// NOTE: the name of this function - "kernel_single_task" - is used by the -// Front End to determine kernel invocation kind. -template + // NOTE: the name of this function - "kernel_single_task" - is used by the + // Front End to determine kernel invocation kind. + template #ifdef __SYCL_DEVICE_ONLY__ -[[__sycl_detail__::add_ir_attributes_function( - "sycl-single-task", - ext::oneapi::experimental::detail::PropertyMetaInfo::name..., - nullptr, - ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] + [[__sycl_detail__::add_ir_attributes_function( + "sycl-single-task", + ext::oneapi::experimental::detail::PropertyMetaInfo::name..., + nullptr, + ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif -__SYCL_KERNEL_ATTR__ static void -kernel_single_task(const KernelType &KernelFunc, kernel_handler KH) { + __SYCL_KERNEL_ATTR__ static void + kernel_single_task(const KernelType &KernelFunc, kernel_handler KH) { #ifdef __SYCL_DEVICE_ONLY__ - KernelFunc(KH); + KernelFunc(KH); #else - (void)KernelFunc; - (void)KH; + (void)KernelFunc; + (void)KH; #endif -} + } -// NOTE: the name of these functions - "kernel_parallel_for" - are used by the -// Front End to determine kernel invocation kind. -template + // NOTE: the name of these functions - "kernel_parallel_for" - are used by the + // Front End to determine kernel invocation kind. + template #ifdef __SYCL_DEVICE_ONLY__ -[[__sycl_detail__::add_ir_attributes_function( - ext::oneapi::experimental::detail::PropertyMetaInfo::name..., - ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] + [[__sycl_detail__::add_ir_attributes_function( + ext::oneapi::experimental::detail::PropertyMetaInfo::name..., + ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif -__SYCL_KERNEL_ATTR__ static void -kernel_parallel_for(const KernelType &KernelFunc) { + __SYCL_KERNEL_ATTR__ static void + kernel_parallel_for(const KernelType &KernelFunc) { #ifdef __SYCL_DEVICE_ONLY__ - KernelFunc(detail::Builder::getElement(detail::declptr())); + KernelFunc(detail::Builder::getElement(detail::declptr())); #else - (void)KernelFunc; + (void)KernelFunc; #endif -} + } -// NOTE: the name of these functions - "kernel_parallel_for" - are used by the -// Front End to determine kernel invocation kind. -template + // NOTE: the name of these functions - "kernel_parallel_for" - are used by the + // Front End to determine kernel invocation kind. + template #ifdef __SYCL_DEVICE_ONLY__ -[[__sycl_detail__::add_ir_attributes_function( - ext::oneapi::experimental::detail::PropertyMetaInfo::name..., - ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] + [[__sycl_detail__::add_ir_attributes_function( + ext::oneapi::experimental::detail::PropertyMetaInfo::name..., + ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif -__SYCL_KERNEL_ATTR__ static void -kernel_parallel_for(const KernelType &KernelFunc, kernel_handler KH) { + __SYCL_KERNEL_ATTR__ static void + kernel_parallel_for(const KernelType &KernelFunc, kernel_handler KH) { #ifdef __SYCL_DEVICE_ONLY__ - KernelFunc(detail::Builder::getElement(detail::declptr()), KH); + KernelFunc(detail::Builder::getElement(detail::declptr()), KH); #else - (void)KernelFunc; - (void)KH; + (void)KernelFunc; + (void)KH; #endif -} + } -// NOTE: the name of this function - "kernel_parallel_for_work_group" - is -// used by the Front End to determine kernel invocation kind. -template + // NOTE: the name of this function - "kernel_parallel_for_work_group" - is + // used by the Front End to determine kernel invocation kind. + template #ifdef __SYCL_DEVICE_ONLY__ -[[__sycl_detail__::add_ir_attributes_function( - ext::oneapi::experimental::detail::PropertyMetaInfo::name..., - ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] + [[__sycl_detail__::add_ir_attributes_function( + ext::oneapi::experimental::detail::PropertyMetaInfo::name..., + ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif -__SYCL_KERNEL_ATTR__ static void -kernel_parallel_for_work_group(const KernelType &KernelFunc) { + __SYCL_KERNEL_ATTR__ static void + kernel_parallel_for_work_group(const KernelType &KernelFunc) { #ifdef __SYCL_DEVICE_ONLY__ - KernelFunc(detail::Builder::getElement(detail::declptr())); + KernelFunc(detail::Builder::getElement(detail::declptr())); #else - (void)KernelFunc; + (void)KernelFunc; #endif -} + } -// NOTE: the name of this function - "kernel_parallel_for_work_group" - is -// used by the Front End to determine kernel invocation kind. -template + // NOTE: the name of this function - "kernel_parallel_for_work_group" - is + // used by the Front End to determine kernel invocation kind. + template #ifdef __SYCL_DEVICE_ONLY__ -[[__sycl_detail__::add_ir_attributes_function( - ext::oneapi::experimental::detail::PropertyMetaInfo::name..., - ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] + [[__sycl_detail__::add_ir_attributes_function( + ext::oneapi::experimental::detail::PropertyMetaInfo::name..., + ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif -__SYCL_KERNEL_ATTR__ static void -kernel_parallel_for_work_group(const KernelType &KernelFunc, - kernel_handler KH) { + __SYCL_KERNEL_ATTR__ static void + kernel_parallel_for_work_group(const KernelType &KernelFunc, + kernel_handler KH) { #ifdef __SYCL_DEVICE_ONLY__ - KernelFunc(detail::Builder::getElement(detail::declptr()), KH); + KernelFunc(detail::Builder::getElement(detail::declptr()), KH); #else - (void)KernelFunc; - (void)KH; + (void)KernelFunc; + (void)KH; #endif -} + } +}; // KernelWrapperSingletonFunc /// Process kernel properties. /// @@ -230,7 +233,9 @@ template > { + ext::oneapi::experimental::detail::properties_t> + : public KernelWrapperSingletonFunc { + static void wrap([[maybe_unused]] PropertyProcessor h, [[maybe_unused]] const KernelType &KernelFunc) { #ifdef __SYCL_DEVICE_ONLY__ From 4c11bd259af6e58928865f872f296669b787b249 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Fri, 18 Jul 2025 20:30:10 +0200 Subject: [PATCH 03/11] Process properties back in handler --- sycl/include/sycl/handler.hpp | 34 +++++++++++++++++-- sycl/include/sycl/kernel_helper.hpp | 34 ++----------------- .../check_device_code/device_has_kernel.cpp | 4 +-- .../esimd/slm_init_specconst_size.cpp | 2 +- .../extensions/annotated_arg/unaliased.cpp | 2 +- .../properties_kernel_sub_group_size.cpp | 25 +++++++------- .../properties_kernel_work_group_size.cpp | 30 +++++++--------- ...properties_kernel_work_group_size_hint.cpp | 27 +++++++-------- sycl/test/check_device_code/no_offset.cpp | 4 +-- .../non_esimd_kernel_fp_control.cpp | 4 +-- .../include_deps/sycl_detail_core.hpp.cpp | 2 +- .../virtual-functions/properties-negative.cpp | 6 ++-- 12 files changed, 83 insertions(+), 91 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 1c54196066747..27650d6ab3163 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -902,7 +902,6 @@ class __SYCL_EXPORT handler { } } -public: /// Process runtime kernel properties. /// /// Stores information about kernel properties into the handler. @@ -968,6 +967,37 @@ class __SYCL_EXPORT handler { checkAndSetClusterRange(Props); } +public: + /// Process kernel properties. + /// + /// Stores information about kernel properties into the handler. + /// + /// 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 + /// the same, thus unnecessary increasing compilation time. + template < + bool IsESIMDKernel, + typename PropertiesT = ext::oneapi::experimental::empty_properties_t> + void processProperties(PropertiesT Props) { + static_assert( + ext::oneapi::experimental::is_property_list::value, + "Template type is not a property list."); + static_assert( + !PropertiesT::template has_property< + sycl::ext::intel::experimental::fp_control_key>() || + (PropertiesT::template has_property< + sycl::ext::intel::experimental::fp_control_key>() && + IsESIMDKernel), + "Floating point control property is supported for ESIMD kernels only."); + static_assert( + !PropertiesT::template has_property< + sycl::ext::oneapi::experimental::indirectly_callable_key>(), + "indirectly_callable property cannot be applied to SYCL kernels"); + + processLaunchProperties(Props); + } + private: /// Checks whether it is possible to copy the source shape to the destination /// shape(the shapes are described by the accessor ranges) by using @@ -1892,7 +1922,7 @@ class __SYCL_EXPORT handler { using NameT = typename detail::get_kernel_name_t::name; (void)Kernel; - KernelWrapperSingletonFunc::kernel_single_task(KernelFunc); + KernelWrapperSingletonFuncs::kernel_single_task(KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ throwIfActionIsCreated(); constexpr detail::string_view Name{detail::getKernelName()}; diff --git a/sycl/include/sycl/kernel_helper.hpp b/sycl/include/sycl/kernel_helper.hpp index e99f618ab5894..6cf7c8f565b9e 100644 --- a/sycl/include/sycl/kernel_helper.hpp +++ b/sycl/include/sycl/kernel_helper.hpp @@ -47,7 +47,7 @@ struct GetMergedKernelProperties< }; } // namespace detail -struct KernelWrapperSingletonFunc { +struct KernelWrapperSingletonFuncs { #ifdef SYCL_LANGUAGE_VERSION #ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS @@ -175,32 +175,6 @@ struct KernelWrapperSingletonFunc { } }; // KernelWrapperSingletonFunc -/// Process kernel properties. -/// -/// Stores information about kernel properties into the handler. -/// -/// 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 -/// the same, thus unnecessary increasing compilation time. -template -void processProperties([[maybe_unused]] PropertiesT Props) { - static_assert(ext::oneapi::experimental::is_property_list::value, - "Template type is not a property list."); - static_assert( - !PropertiesT::template has_property< - sycl::ext::intel::experimental::fp_control_key>() || - (PropertiesT::template has_property< - sycl::ext::intel::experimental::fp_control_key>() && - IsESIMDKernel), - "Floating point control property is supported for ESIMD kernels only."); - static_assert( - !PropertiesT::template has_property< - sycl::ext::oneapi::experimental::indirectly_callable_key>(), - "indirectly_callable property cannot be applied to SYCL kernels"); -} - // The KernelWrapper below has two purposes. // // First, from SYCL 2020, Table 129 (Member functions of the `handler ` class) @@ -234,7 +208,7 @@ struct KernelWrapper< WrapAsVal, KernelName, KernelType, ElementType, PropertyProcessor, PropertiesT, ext::oneapi::experimental::detail::properties_t> - : public KernelWrapperSingletonFunc { + : public KernelWrapperSingletonFuncs { static void wrap([[maybe_unused]] PropertyProcessor h, [[maybe_unused]] const KernelType &KernelFunc) { @@ -245,9 +219,7 @@ struct KernelWrapper< if constexpr (ext::oneapi::experimental::detail:: HasKernelPropertiesGetMethod::value) { - processProperties()>( - KernelFunc.get(ext::oneapi::experimental::properties_tag{})); - h->processLaunchProperties( + h->template processProperties()>( KernelFunc.get(ext::oneapi::experimental::properties_tag{})); } #endif diff --git a/sycl/test/check_device_code/device_has_kernel.cpp b/sycl/test/check_device_code/device_has_kernel.cpp index f25bc27eccce8..191dfe49d2030 100644 --- a/sycl/test/check_device_code/device_has_kernel.cpp +++ b/sycl/test/check_device_code/device_has_kernel.cpp @@ -10,7 +10,7 @@ using namespace sycl; queue q; -// CHECK-ASPECTS: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] {{.*}} +// CHECK-ASPECTS: define weak_odr dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] {{.*}} // CHECK-ASPECTS: define {{.*}}spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] // CHECK-ASPECTS-SAME: !sycl_used_aspects ![[ASPECTS1]] @@ -53,7 +53,7 @@ void foo() { q.submit([&](handler &h) { KernelFunctor f1; h.single_task(f1); - // CHECK-ASPECTS: define dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]] + // CHECK-ASPECTS: define weak_odr dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]] h.single_task( []() [[sycl::device_has(sycl::aspect::gpu)]] {}); }); diff --git a/sycl/test/check_device_code/esimd/slm_init_specconst_size.cpp b/sycl/test/check_device_code/esimd/slm_init_specconst_size.cpp index 1991616c3f698..a2f76ff0981ab 100644 --- a/sycl/test/check_device_code/esimd/slm_init_specconst_size.cpp +++ b/sycl/test/check_device_code/esimd/slm_init_specconst_size.cpp @@ -21,7 +21,7 @@ int main() { [=](sycl::kernel_handler kh) SYCL_ESIMD_KERNEL { slm_init(kh.get_specialization_constant()); }); - // CHECK: define dso_local spir_kernel void @{{.*}}() local_unnamed_addr #1 + // CHECK: define weak_odr dso_local spir_kernel void @{{.*}}() local_unnamed_addr #1 }); } diff --git a/sycl/test/check_device_code/extensions/annotated_arg/unaliased.cpp b/sycl/test/check_device_code/extensions/annotated_arg/unaliased.cpp index e603b79d4e5e2..028799fbe5214 100644 --- a/sycl/test/check_device_code/extensions/annotated_arg/unaliased.cpp +++ b/sycl/test/check_device_code/extensions/annotated_arg/unaliased.cpp @@ -19,4 +19,4 @@ int main() { return 0; } -// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef writeonly align 4 captures(none) "sycl-unaliased" +// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef align 4 "sycl-unaliased" diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp index 1a49166819e42..ad81d1db1fe0b 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp @@ -79,28 +79,28 @@ int main() { Q.parallel_for(R3, {Ev}, Props, Redu1, [](sycl::id<3>, auto &) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel27(){{.*}} #[[SGSizeAttr4:[0-9]+]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel27(){{.*}} #[[SGSizeAttr2]] Q.parallel_for(NDR1, Props, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel28(){{.*}} #[[SGSizeAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel28(){{.*}} #[[SGSizeAttr2]] Q.parallel_for(NDR1, Ev, Props, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel29(){{.*}} #[[SGSizeAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel29(){{.*}} #[[SGSizeAttr2]] Q.parallel_for(NDR1, {Ev}, Props, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel30(){{.*}} #[[SGSizeAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel30(){{.*}} #[[SGSizeAttr2]] Q.parallel_for(NDR2, Props, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel31(){{.*}} #[[SGSizeAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel31(){{.*}} #[[SGSizeAttr2]] Q.parallel_for(NDR2, Ev, Props, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel32(){{.*}} #[[SGSizeAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel32(){{.*}} #[[SGSizeAttr2]] Q.parallel_for(NDR2, {Ev}, Props, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel33(){{.*}} #[[SGSizeAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel33(){{.*}} #[[SGSizeAttr2]] Q.parallel_for(NDR3, Props, [](sycl::nd_item<3>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel34(){{.*}} #[[SGSizeAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel34(){{.*}} #[[SGSizeAttr2]] Q.parallel_for(NDR3, Ev, Props, [](sycl::nd_item<3>) {}); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel35(){{.*}} #[[SGSizeAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel35(){{.*}} #[[SGSizeAttr2]] Q.parallel_for(NDR3, {Ev}, Props, [](sycl::nd_item<3>) {}); @@ -202,17 +202,17 @@ int main() { [](sycl::id<3>, auto &) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel63(){{.*}} #[[SGSizeAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel63(){{.*}} #[[SGSizeAttr2]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR1, Props, [](sycl::nd_item<1>) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel64(){{.*}} #[[SGSizeAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel64(){{.*}} #[[SGSizeAttr2]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR2, Props, [](sycl::nd_item<2>) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel65(){{.*}} #[[SGSizeAttr4]] + // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel65(){{.*}} #[[SGSizeAttr2]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR3, Props, [](sycl::nd_item<3>) {}); @@ -278,4 +278,3 @@ int main() { // CHECK-IR: attributes #[[SGSizeAttr1]] = { {{.*}}"sycl-sub-group-size"="1" // CHECK-IR: attributes #[[SGSizeAttr2]] = { {{.*}}"sycl-sub-group-size"="1" // CHECK-IR: attributes #[[SGSizeAttr3]] = { {{.*}}"sycl-sub-group-size"="1" -// CHECK-IR: attributes #[[SGSizeAttr4]] = { {{.*}}"sycl-sub-group-size"="1" diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp index 18546b5250be9..63280fcc638f3 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp @@ -95,28 +95,28 @@ int main() { Q.parallel_for(R3, {Ev}, Props3, Redu1, [](sycl::id<3>, auto &) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel27(){{.*}} #[[WGSizeAttr10:[0-9]+]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel27(){{.*}} #[[WGSizeAttr4]] Q.parallel_for(NDR1, Props1, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel28(){{.*}} #[[WGSizeAttr10]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel28(){{.*}} #[[WGSizeAttr4]] Q.parallel_for(NDR1, Ev, Props1, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel29(){{.*}} #[[WGSizeAttr10]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel29(){{.*}} #[[WGSizeAttr4]] Q.parallel_for(NDR1, {Ev}, Props1, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel30(){{.*}} #[[WGSizeAttr11:[0-9]+]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel30(){{.*}} #[[WGSizeAttr5]] Q.parallel_for(NDR2, Props2, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel31(){{.*}} #[[WGSizeAttr11]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel31(){{.*}} #[[WGSizeAttr5]] Q.parallel_for(NDR2, Ev, Props2, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel32(){{.*}} #[[WGSizeAttr11]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel32(){{.*}} #[[WGSizeAttr5]] Q.parallel_for(NDR2, {Ev}, Props2, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel33(){{.*}} #[[WGSizeAttr12:[0-9]+]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel33(){{.*}} #[[WGSizeAttr6]] Q.parallel_for(NDR3, Props3, [](sycl::nd_item<3>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel34(){{.*}} #[[WGSizeAttr12]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel34(){{.*}} #[[WGSizeAttr6]] Q.parallel_for(NDR3, Ev, Props3, [](sycl::nd_item<3>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel35(){{.*}} #[[WGSizeAttr12]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel35(){{.*}} #[[WGSizeAttr6]] Q.parallel_for(NDR3, {Ev}, Props3, [](sycl::nd_item<3>) {}); @@ -218,20 +218,17 @@ int main() { [](sycl::id<3>, auto &) {}); }); - // 15 - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel63(){{.*}} #[[WGSizeAttr10]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel63(){{.*}} #[[WGSizeAttr4]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR1, Props1, [](sycl::nd_item<1>) {}); }); - // 16 - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel64(){{.*}} #[[WGSizeAttr11]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel64(){{.*}} #[[WGSizeAttr5]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR2, Props2, [](sycl::nd_item<2>) {}); }); - // 17 - // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel65(){{.*}} #[[WGSizeAttr12]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel65(){{.*}} #[[WGSizeAttr6]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR3, Props3, [](sycl::nd_item<3>) {}); @@ -303,6 +300,3 @@ int main() { // CHECK-IR: attributes #[[WGSizeAttr7]] = { {{.*}}"sycl-work-group-size"="1" // CHECK-IR: attributes #[[WGSizeAttr8]] = { {{.*}}"sycl-work-group-size"="1,2" // CHECK-IR: attributes #[[WGSizeAttr9]] = { {{.*}}"sycl-work-group-size"="1,2,3" -// CHECK-IR: attributes #[[WGSizeAttr10]] = { {{.*}}"sycl-work-group-size"="1" -// CHECK-IR: attributes #[[WGSizeAttr11]] = { {{.*}}"sycl-work-group-size"="1,2" -// CHECK-IR: attributes #[[WGSizeAttr12]] = { {{.*}}"sycl-work-group-size"="1,2,3" \ No newline at end of file diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp index 502b28244ba20..a0bae31ad8004 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp @@ -98,31 +98,31 @@ int main() { Q.parallel_for(R3, {Ev}, Props3, Redu1, [](sycl::id<3>, auto &) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel27(){{.*}} #[[WGSizeHintAttr10:[0-9]+]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel27(){{.*}} #[[WGSizeHintAttr4]] Q.parallel_for(NDR1, Props1, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel28(){{.*}} #[[WGSizeHintAttr10]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel28(){{.*}} #[[WGSizeHintAttr4]] Q.parallel_for(NDR1, Ev, Props1, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel29(){{.*}} #[[WGSizeHintAttr10]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel29(){{.*}} #[[WGSizeHintAttr4]] Q.parallel_for(NDR1, {Ev}, Props1, [](sycl::nd_item<1>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel30(){{.*}} #[[WGSizeHintAttr11:[0-9]+]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel30(){{.*}} #[[WGSizeHintAttr5]] Q.parallel_for(NDR2, Props2, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel31(){{.*}} #[[WGSizeHintAttr11]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel31(){{.*}} #[[WGSizeHintAttr5]] Q.parallel_for(NDR2, Ev, Props2, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel32(){{.*}} #[[WGSizeHintAttr11]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel32(){{.*}} #[[WGSizeHintAttr5]] Q.parallel_for(NDR2, {Ev}, Props2, [](sycl::nd_item<2>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel33(){{.*}} #[[WGSizeHintAttr12:[0-9]+]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel33(){{.*}} #[[WGSizeHintAttr6]] Q.parallel_for(NDR3, Props3, [](sycl::nd_item<3>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel34(){{.*}} #[[WGSizeHintAttr12]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel34(){{.*}} #[[WGSizeHintAttr6]] Q.parallel_for(NDR3, Ev, Props3, [](sycl::nd_item<3>) {}); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel35(){{.*}} #[[WGSizeHintAttr12]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel35(){{.*}} #[[WGSizeHintAttr6]] Q.parallel_for(NDR3, {Ev}, Props3, [](sycl::nd_item<3>) {}); @@ -227,17 +227,17 @@ int main() { [](sycl::id<3>, auto &) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel63(){{.*}} #[[WGSizeHintAttr10]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel63(){{.*}} #[[WGSizeHintAttr4]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR1, Props1, [](sycl::nd_item<1>) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel64(){{.*}} #[[WGSizeHintAttr11]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel64(){{.*}} #[[WGSizeHintAttr5]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR2, Props2, [](sycl::nd_item<2>) {}); }); - // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel65(){{.*}} #[[WGSizeHintAttr12]] + // CHECK-IR: spir_kernel void @{{.*}}WGSizeHintKernel65(){{.*}} #[[WGSizeHintAttr6]] Q.submit([&](sycl::handler &CGH) { CGH.parallel_for(NDR3, Props3, [](sycl::nd_item<3>) {}); @@ -309,6 +309,3 @@ int main() { // CHECK-IR: attributes #[[WGSizeHintAttr7]] = { {{.*}}"sycl-work-group-size-hint"="1" // CHECK-IR: attributes #[[WGSizeHintAttr8]] = { {{.*}}"sycl-work-group-size-hint"="1,2" // CHECK-IR: attributes #[[WGSizeHintAttr9]] = { {{.*}}"sycl-work-group-size-hint"="1,2,3" -// CHECK-IR: attributes #[[WGSizeHintAttr10]] = { {{.*}}"sycl-work-group-size-hint"="1" -// CHECK-IR: attributes #[[WGSizeHintAttr11]] = { {{.*}}"sycl-work-group-size-hint"="1,2" -// CHECK-IR: attributes #[[WGSizeHintAttr12]] = { {{.*}}"sycl-work-group-size-hint"="1,2,3" diff --git a/sycl/test/check_device_code/no_offset.cpp b/sycl/test/check_device_code/no_offset.cpp index d2c1a2ce43df8..f311c7dba40cf 100644 --- a/sycl/test/check_device_code/no_offset.cpp +++ b/sycl/test/check_device_code/no_offset.cpp @@ -15,7 +15,7 @@ int main() { sycl::ext::oneapi::accessor_property_list PL{sycl::ext::oneapi::no_offset, sycl::no_init}; sycl::accessor acc_a(a, cgh, sycl::write_only, PL); sycl::accessor acc_b{b, cgh, sycl::read_only}; - // CHECK: define dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlT_E_(ptr addrspace(1) {{.*}}, ptr addrspace(1) noundef readonly {{.*}}, ptr noundef readonly byval(%"class.sycl::_V1::id") align 8 {{.*}}) + // CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlT_E_(ptr addrspace(1) {{.*}}, ptr addrspace(1) noundef readonly {{.*}}, ptr noundef byval(%"class.sycl::_V1::id") align 8 {{.*}}) cgh.parallel_for(size, [=](auto i) { acc_a[i] = acc_b[i]; }); @@ -33,7 +33,7 @@ int main() { q.submit([&](sycl::handler &cgh) { sycl::accessor acc_a(a, cgh, sycl::write_only); sycl::accessor acc_b{b, cgh, sycl::read_only}; - // CHECK: define dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_EUlT_E_(ptr addrspace(1) {{.*}}, ptr noundef readonly byval(%"class.sycl::_V1::id") align 8 {{.*}}, ptr noundef readonly byval(%"class.sycl::_V1::id") align 8 {{.*}}) + // CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_EUlT_E_(ptr addrspace(1) {{.*}}, ptr noundef byval(%"class.sycl::_V1::id") align 8 {{.*}}, ptr addrspace(1) noundef readonly {{.*}}, ptr noundef byval(%"class.sycl::_V1::id") align 8 {{.*}}) cgh.parallel_for(size, [=](auto i) { acc_a[i] = acc_b[i]; }); diff --git a/sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp b/sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp index a4152802a1b17..46d11eccdfe54 100644 --- a/sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp +++ b/sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp @@ -20,7 +20,7 @@ struct ESIMDKernel { int main(void) { queue q; - // expected-error-re@sycl/kernel_helper.hpp:* {{static assertion failed due to requirement {{.+}}: Floating point control property is supported for ESIMD kernels only.}} + // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.+}}: Floating point control property is supported for ESIMD kernels only.}} syclex::properties properties7{ intelex::fp_control}; @@ -28,7 +28,7 @@ int main(void) { cgh.single_task(properties7, [=]() {}); }); - // expected-error-re@sycl/kernel_helper.hpp:* {{static assertion failed due to requirement {{.+}}: Floating point control property is supported for ESIMD kernels only.}} + // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.+}}: Floating point control property is supported for ESIMD kernels only.}} ESIMDKernel Kern; q.submit([&](handler &cgh) { cgh.parallel_for(range<1>(1), Kern); }); diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp index 5342539cce89f..68b65f5f4ff3e 100644 --- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp +++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp @@ -146,8 +146,8 @@ // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp // CHECK-NEXT: detail/sycl_local_mem_builtins.hpp -// CHECK-NEXT: kernel_helper.hpp // CHECK-NEXT: kernel.hpp +// CHECK-NEXT: kernel_helper.hpp // CHECK-NEXT: sampler.hpp // CHECK-NEXT: feature_test.hpp // CHECK-EMPTY: diff --git a/sycl/test/virtual-functions/properties-negative.cpp b/sycl/test/virtual-functions/properties-negative.cpp index 1a28969f86a8d..b8e1b75f1d9a9 100644 --- a/sycl/test/virtual-functions/properties-negative.cpp +++ b/sycl/test/virtual-functions/properties-negative.cpp @@ -17,15 +17,15 @@ int main() { oneapi::properties props_int{oneapi::indirectly_callable_in}; oneapi::properties props_user{oneapi::indirectly_callable_in}; - // expected-error-re@sycl/kernel_helper.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}} + // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}} q.single_task(props_empty, [=]() {}); // When both "props_empty" and "props_void" are in use, we won't see the // static assert firing for the second one, because there will be only one // instantiation of handler::processProperties. q.single_task(props_void, [=]() {}); - // expected-error-re@sycl/kernel_helper.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}} + // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}} q.single_task(props_int, [=]() {}); - // expected-error-re@sycl/kernel_helper.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}} + // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}} q.single_task(props_user, [=]() {}); return 0; From b539f28d5e613fb71c0894ce8c765aabc2680682 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Mon, 21 Jul 2025 17:18:49 +0200 Subject: [PATCH 04/11] Make KernelWrapper a friend of sycl::handler --- sycl/include/sycl/handler.hpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 27650d6ab3163..95059a9176646 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -967,7 +967,6 @@ class __SYCL_EXPORT handler { checkAndSetClusterRange(Props); } -public: /// Process kernel properties. /// /// Stores information about kernel properties into the handler. @@ -998,7 +997,6 @@ class __SYCL_EXPORT handler { processLaunchProperties(Props); } -private: /// Checks whether it is possible to copy the source shape to the destination /// shape(the shapes are described by the accessor ranges) by using /// copying by regions of memory and not copying element by element @@ -3649,6 +3647,9 @@ class __SYCL_EXPORT handler { void instantiateKernelOnHost(void *InstantiateKernelOnHostPtr); friend class detail::HandlerAccess; + template + friend struct KernelWrapper; #ifdef __INTEL_PREVIEW_BREAKING_CHANGES __SYCL_DLL_LOCAL detail::handler_impl *get_impl() { return impl; } From 06b8e1958c525ba20ffe8ff0ad8f8353d5b31275 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Mon, 21 Jul 2025 18:04:02 +0200 Subject: [PATCH 05/11] Move kernelWrapper in the detail namespace --- sycl/include/sycl/handler.hpp | 14 +++++++------- sycl/include/sycl/kernel_helper.hpp | 6 +++--- 2 files changed, 10 insertions(+), 10 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 95059a9176646..32ad5182c86dc 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1298,7 +1298,7 @@ class __SYCL_EXPORT handler { using KName = std::conditional_t::value, decltype(Wrapper), NameWT>; - KernelWrapper::wrap(this, Wrapper); #ifndef __SYCL_DEVICE_ONLY__ @@ -1324,7 +1324,7 @@ class __SYCL_EXPORT handler { #ifndef __SYCL_FORCE_PARALLEL_FOR_RANGE_ROUNDING__ // If parallel_for range rounding is forced then only range rounded // kernel is generated - KernelWrapper::wrap(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ constexpr detail::string_view Name{detail::getKernelName()}; @@ -1404,7 +1404,7 @@ class __SYCL_EXPORT handler { using NameT = typename detail::get_kernel_name_t::name; (void)Props; - KernelWrapper::wrap(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ if constexpr (WrapAsVal == WrapAs::single_task) { @@ -1445,7 +1445,7 @@ class __SYCL_EXPORT handler { typename detail::get_kernel_name_t::name; (void)Props; (void)Kernel; - KernelWrapper::wrap(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ if constexpr (WrapAsVal == WrapAs::single_task) { @@ -1480,7 +1480,7 @@ class __SYCL_EXPORT handler { #endif // __INTEL_PREVIEW_BREAKING_CHANGES // NOTE: to support kernel_handler argument in kernel lambdas, only - // KernelWrapper<...>::wrap() must be called in this code. + // detail::KernelWrapper<...>::wrap() must be called in this code. void setStateExplicitKernelBundle(); void setStateSpecConstSet(); @@ -1920,7 +1920,7 @@ class __SYCL_EXPORT handler { using NameT = typename detail::get_kernel_name_t::name; (void)Kernel; - KernelWrapperSingletonFuncs::kernel_single_task(KernelFunc); + detail::KernelWrapperSingletonFuncs::kernel_single_task(KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ throwIfActionIsCreated(); constexpr detail::string_view Name{detail::getKernelName()}; @@ -3649,7 +3649,7 @@ class __SYCL_EXPORT handler { friend class detail::HandlerAccess; template - friend struct KernelWrapper; + friend struct detail::KernelWrapper; #ifdef __INTEL_PREVIEW_BREAKING_CHANGES __SYCL_DLL_LOCAL detail::handler_impl *get_impl() { return impl; } diff --git a/sycl/include/sycl/kernel_helper.hpp b/sycl/include/sycl/kernel_helper.hpp index 6cf7c8f565b9e..7f9a396b54d43 100644 --- a/sycl/include/sycl/kernel_helper.hpp +++ b/sycl/include/sycl/kernel_helper.hpp @@ -23,6 +23,8 @@ namespace sycl { inline namespace _V1 { +enum class WrapAs { single_task, parallel_for, parallel_for_work_group }; + namespace detail { // Helper for merging properties with ones defined in an optional kernel functor // getter. @@ -45,7 +47,6 @@ struct GetMergedKernelProperties< using type = ext::oneapi::experimental::detail::merged_properties_t< PropertiesT, get_method_properties>; }; -} // namespace detail struct KernelWrapperSingletonFuncs { @@ -193,8 +194,6 @@ struct KernelWrapperSingletonFuncs { // the one provided via `sycl_ext_oneapi_kernel_properties` extension by // embedding them into the kernel's type). -enum class WrapAs { single_task, parallel_for, parallel_for_work_group }; - template Date: Mon, 21 Jul 2025 18:34:32 +0200 Subject: [PATCH 06/11] Rename helper header to kernel_launch_helper. --- sycl/include/sycl/handler.hpp | 8 +------- ...{kernel_helper.hpp => kernel_launch_helper.hpp} | 7 +++++-- .../handler/unnamed-lambda-negative.cpp | 14 +++++++------- sycl/test/esimd/global_var.cpp | 2 +- .../properties/properties_kernel_negative.cpp | 2 +- sycl/test/include_deps/sycl_detail_core.hpp.cpp | 10 +++++----- 6 files changed, 20 insertions(+), 23 deletions(-) rename sycl/include/sycl/{kernel_helper.hpp => kernel_launch_helper.hpp} (97%) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 32ad5182c86dc..fff62c0f2c659 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -10,7 +10,6 @@ #include #include -#include #include #include #include @@ -27,8 +26,6 @@ #include #include #include -#include -#include #include #include #include @@ -37,17 +34,14 @@ #include #include #include -#include #include #include -#include #include #include #include #include #include -#include -#include +#include #include #include #include diff --git a/sycl/include/sycl/kernel_helper.hpp b/sycl/include/sycl/kernel_launch_helper.hpp similarity index 97% rename from sycl/include/sycl/kernel_helper.hpp rename to sycl/include/sycl/kernel_launch_helper.hpp index 7f9a396b54d43..7054a5e137ef6 100644 --- a/sycl/include/sycl/kernel_helper.hpp +++ b/sycl/include/sycl/kernel_launch_helper.hpp @@ -1,4 +1,4 @@ -//==-------- kernel_helper.hpp --- SYCL kernel helper utilities ---------==// +//==-------- kernel_launch_helper.hpp --- SYCL kernel launch utilities ----==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -23,6 +23,8 @@ namespace sycl { inline namespace _V1 { +// TODO: Move WrapAs to detail:: namespace as well and move this header to +// include/sycl/detail directory. enum class WrapAs { single_task, parallel_for, parallel_for_work_group }; namespace detail { @@ -218,6 +220,7 @@ struct KernelWrapper< if constexpr (ext::oneapi::experimental::detail:: HasKernelPropertiesGetMethod::value) { + // TODO: decouple property processing from KernelWrapper. h->template processProperties()>( KernelFunc.get(ext::oneapi::experimental::properties_tag{})); } @@ -264,4 +267,4 @@ struct KernelWrapper< } // namespace detail } // namespace _V1 -} // namespace sycl \ No newline at end of file +} // namespace sycl diff --git a/sycl/test/basic_tests/handler/unnamed-lambda-negative.cpp b/sycl/test/basic_tests/handler/unnamed-lambda-negative.cpp index c17a60db250a0..f250854f64e9c 100644 --- a/sycl/test/basic_tests/handler/unnamed-lambda-negative.cpp +++ b/sycl/test/basic_tests/handler/unnamed-lambda-negative.cpp @@ -8,26 +8,26 @@ int main() { sycl::queue q; - // expected-error-re@sycl/kernel_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} q.single_task([=](){}); - // expected-error-re@sycl/kernel_helper.hpp:* {{unnamed type 'sycl::detail::RoundedRangeKernel<{{.*}}>' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} - // expected-error-re@sycl/kernel_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/kernel_launch_helper.hpp:* {{unnamed type 'sycl::detail::RoundedRangeKernel<{{.*}}>' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} q.parallel_for(sycl::range<1>{1}, [=](sycl::item<1>) {}); q.submit([&](sycl::handler &cgh) { - // expected-error-re@sycl/kernel_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} cgh.single_task([=](){}); - // expected-error-re@sycl/kernel_helper.hpp:* {{unnamed type 'sycl::detail::RoundedRangeKernel<{{.*}}>' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} - // expected-error-re@sycl/kernel_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/kernel_launch_helper.hpp:* {{unnamed type 'sycl::detail::RoundedRangeKernel<{{.*}}>' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} cgh.parallel_for(sycl::range<1>{1}, [=](sycl::item<1>) {}); - // expected-error-re@sycl/kernel_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} cgh.parallel_for_work_group(sycl::range<1>{1}, [=](sycl::group<1>) {}); }); diff --git a/sycl/test/esimd/global_var.cpp b/sycl/test/esimd/global_var.cpp index 61ff4c2f39d37..2002cc82f4af4 100644 --- a/sycl/test/esimd/global_var.cpp +++ b/sycl/test/esimd/global_var.cpp @@ -42,7 +42,7 @@ void kernel_call() { q.submit([&](sycl::handler &cgh) { cgh.parallel_for(nd_range<1>(1, 1), [=](nd_item<1> ndi) { - //expected-note@sycl/kernel_helper.hpp:* 2{{called by 'kernel_parallel_for}} + //expected-note@sycl/kernel_launch_helper.hpp:* 2{{called by 'kernel_parallel_for}} //expected-error@+1{{ESIMD globals cannot be used in a SYCL context}} vc = 0; //expected-note@+1{{called by}} diff --git a/sycl/test/extensions/properties/properties_kernel_negative.cpp b/sycl/test/extensions/properties/properties_kernel_negative.cpp index c88a649e024d4..33b20ed9961ac 100644 --- a/sycl/test/extensions/properties/properties_kernel_negative.cpp +++ b/sycl/test/extensions/properties/properties_kernel_negative.cpp @@ -399,7 +399,7 @@ struct TestKernelNoGetter { void check_non_const_getter_warning() { sycl::queue Q; - // expected-error-re@sycl/kernel_helper.hpp:* {{static assertion failed due to requirement {{.+}}: get(sycl::ext::oneapi::experimental::properties_tag) member in kernel functor class must be declared as a const member function}} + // expected-error-re@sycl/kernel_launch_helper.hpp:* {{static assertion failed due to requirement {{.+}}: get(sycl::ext::oneapi::experimental::properties_tag) member in kernel functor class must be declared as a const member function}} Q.single_task(TestKernelNonConstGetter()); // No error expected for kernel functor with a const get(properties_tag) diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp index 68b65f5f4ff3e..2ef598ae49659 100644 --- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp +++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp @@ -135,19 +135,19 @@ // CHECK-NEXT: detail/reduction_forward.hpp // CHECK-NEXT: detail/ur.hpp // CHECK-NEXT: ur_api_funcs.def -// CHECK-NEXT: ext/intel/experimental/fp_control_kernel_properties.hpp -// CHECK-NEXT: ext/intel/experimental/kernel_execution_properties.hpp // CHECK-NEXT: ext/oneapi/bindless_images_interop.hpp // CHECK-NEXT: ext/oneapi/bindless_images_mem_handle.hpp // CHECK-NEXT: ext/oneapi/experimental/cluster_group_prop.hpp // CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp // CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp -// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp +// CHECK-NEXT: kernel.hpp +// CHECK-NEXT: kernel_launch_helper.hpp +// CHECK-NEXT: ext/intel/experimental/fp_control_kernel_properties.hpp +// CHECK-NEXT: ext/intel/experimental/kernel_execution_properties.hpp +// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp // CHECK-NEXT: detail/sycl_local_mem_builtins.hpp -// CHECK-NEXT: kernel.hpp -// CHECK-NEXT: kernel_helper.hpp // CHECK-NEXT: sampler.hpp // CHECK-NEXT: feature_test.hpp // CHECK-EMPTY: From 4e83f9aec0e8e6919a9cffb3b80a71dd120583e0 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Mon, 21 Jul 2025 18:35:10 +0200 Subject: [PATCH 07/11] clang formatting --- sycl/include/sycl/handler.hpp | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index fff62c0f2c659..f627fe1775204 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1293,8 +1293,8 @@ class __SYCL_EXPORT handler { decltype(Wrapper), NameWT>; detail::KernelWrapper::wrap(this, Wrapper); + TransformedArgType, decltype(this), + PropertiesT>::wrap(this, Wrapper); #ifndef __SYCL_DEVICE_ONLY__ constexpr detail::string_view Name{detail::getKernelName()}; verifyUsedKernelBundleInternal(Name); @@ -1318,8 +1318,9 @@ class __SYCL_EXPORT handler { #ifndef __SYCL_FORCE_PARALLEL_FOR_RANGE_ROUNDING__ // If parallel_for range rounding is forced then only range rounded // kernel is generated - detail::KernelWrapper::wrap(this, KernelFunc); + detail::KernelWrapper::wrap(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ constexpr detail::string_view Name{detail::getKernelName()}; @@ -1398,8 +1399,8 @@ class __SYCL_EXPORT handler { using NameT = typename detail::get_kernel_name_t::name; (void)Props; - detail::KernelWrapper::wrap(this, KernelFunc); + detail::KernelWrapper::wrap(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ if constexpr (WrapAsVal == WrapAs::single_task) { throwOnKernelParameterMisuse(); @@ -1439,8 +1440,8 @@ class __SYCL_EXPORT handler { typename detail::get_kernel_name_t::name; (void)Props; (void)Kernel; - detail::KernelWrapper::wrap(this, KernelFunc); + detail::KernelWrapper::wrap(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ if constexpr (WrapAsVal == WrapAs::single_task) { throwOnKernelParameterMisuse(); @@ -3641,8 +3642,7 @@ class __SYCL_EXPORT handler { void instantiateKernelOnHost(void *InstantiateKernelOnHostPtr); friend class detail::HandlerAccess; - template + template friend struct detail::KernelWrapper; #ifdef __INTEL_PREVIEW_BREAKING_CHANGES From 91e479099c95c8e49b1a01ed3a509ccc0ba6d6a3 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Mon, 21 Jul 2025 19:10:17 +0200 Subject: [PATCH 08/11] Move WrapAs to detail:: namespace --- sycl/include/sycl/handler.hpp | 40 +++++++++++----------- sycl/include/sycl/kernel_launch_helper.hpp | 4 +-- 2 files changed, 21 insertions(+), 23 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index f627fe1775204..9a6cd2a192f69 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1292,7 +1292,7 @@ class __SYCL_EXPORT handler { using KName = std::conditional_t::value, decltype(Wrapper), NameWT>; - detail::KernelWrapper::wrap(this, Wrapper); #ifndef __SYCL_DEVICE_ONLY__ @@ -1318,7 +1318,7 @@ class __SYCL_EXPORT handler { #ifndef __SYCL_FORCE_PARALLEL_FOR_RANGE_ROUNDING__ // If parallel_for range rounding is forced then only range rounded // kernel is generated - detail::KernelWrapper::wrap(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ @@ -1388,7 +1388,7 @@ class __SYCL_EXPORT handler { } template < - WrapAs WrapAsVal, typename KernelName, typename ElementType = void, + detail::WrapAs WrapAsVal, typename KernelName, typename ElementType = void, int Dims = 1, bool SetNumWorkGroups = false, typename PropertiesT = ext::oneapi::experimental::empty_properties_t, typename KernelType, typename... RangeParams> @@ -1402,7 +1402,7 @@ class __SYCL_EXPORT handler { detail::KernelWrapper::wrap(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ - if constexpr (WrapAsVal == WrapAs::single_task) { + if constexpr (WrapAsVal == detail::WrapAs::single_task) { throwOnKernelParameterMisuse(); } throwIfActionIsCreated(); @@ -1427,7 +1427,7 @@ class __SYCL_EXPORT handler { // Implementation for something that had to be removed long ago but now stuck // until next major release... template < - WrapAs WrapAsVal, typename KernelName, typename ElementType = void, + detail::WrapAs WrapAsVal, typename KernelName, typename ElementType = void, int Dims = 1, bool SetNumWorkGroups = false, typename PropertiesT = ext::oneapi::experimental::empty_properties_t, typename KernelType, typename... RangeParams> @@ -1443,7 +1443,7 @@ class __SYCL_EXPORT handler { detail::KernelWrapper::wrap(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ - if constexpr (WrapAsVal == WrapAs::single_task) { + if constexpr (WrapAsVal == detail::WrapAs::single_task) { throwOnKernelParameterMisuse(); } throwIfActionIsCreated(); @@ -1717,7 +1717,7 @@ class __SYCL_EXPORT handler { /// \param KernelFunc is a SYCL kernel function. template void single_task(const KernelType &KernelFunc) { - wrap_kernel(KernelFunc, {} /*Props*/, + wrap_kernel(KernelFunc, {} /*Props*/, range<1>{1}); } @@ -1784,7 +1784,7 @@ class __SYCL_EXPORT handler { using TransformedArgType = std::conditional_t< std::is_integral::value && Dims == 1, item, typename TransformUserItemType::type>; - wrap_kernel( + wrap_kernel( KernelFunc, {} /*Props*/, NumWorkItems, WorkItemOffset); } @@ -1802,7 +1802,7 @@ class __SYCL_EXPORT handler { int Dims> void parallel_for_work_group(range NumWorkGroups, const KernelType &KernelFunc) { - wrap_kernel>, Dims, /*SetNumWorkGroups=*/true>(KernelFunc, {} /*Props*/, NumWorkGroups); @@ -1825,7 +1825,7 @@ class __SYCL_EXPORT handler { void parallel_for_work_group(range NumWorkGroups, range WorkGroupSize, const KernelType &KernelFunc) { - wrap_kernel>, Dims>( KernelFunc, {} /*Props*/, nd_range{NumWorkGroups * WorkGroupSize, WorkGroupSize}); @@ -1951,7 +1951,7 @@ class __SYCL_EXPORT handler { // Ignore any set kernel bundles and use the one associated with the kernel setHandlerKernelBundle(Kernel); using LambdaArgType = sycl::detail::lambda_arg_type>; - wrap_kernel_legacy( + wrap_kernel_legacy( KernelFunc, Kernel, {} /*Props*/, NumWorkItems); } @@ -1970,7 +1970,7 @@ class __SYCL_EXPORT handler { void parallel_for(kernel Kernel, range NumWorkItems, id WorkItemOffset, const KernelType &KernelFunc) { using LambdaArgType = sycl::detail::lambda_arg_type>; - wrap_kernel_legacy( + wrap_kernel_legacy( KernelFunc, Kernel, {} /*Props*/, NumWorkItems, WorkItemOffset); } @@ -1990,7 +1990,7 @@ class __SYCL_EXPORT handler { const KernelType &KernelFunc) { using LambdaArgType = sycl::detail::lambda_arg_type>; - wrap_kernel_legacy( + wrap_kernel_legacy( KernelFunc, Kernel, {} /*Props*/, NDRange); } @@ -2014,7 +2014,7 @@ class __SYCL_EXPORT handler { const KernelType &KernelFunc) { using LambdaArgType = sycl::detail::lambda_arg_type>; - wrap_kernel_legacy(KernelFunc, Kernel, {} /*Props*/, NumWorkGroups); @@ -2045,7 +2045,7 @@ class __SYCL_EXPORT handler { sycl::detail::lambda_arg_type>; nd_range ExecRange = nd_range(NumWorkGroups * WorkGroupSize, WorkGroupSize); - wrap_kernel_legacy(KernelFunc, Kernel, {} /*Props*/, ExecRange); } @@ -2060,7 +2060,7 @@ class __SYCL_EXPORT handler { std::enable_if_t::value> single_task(PropertiesT Props, const KernelType &KernelFunc) { - wrap_kernel(KernelFunc, Props, + wrap_kernel(KernelFunc, Props, range<1>{1}); } @@ -2124,7 +2124,7 @@ class __SYCL_EXPORT handler { "must be either sycl::nd_item or be convertible from sycl::nd_item"); using TransformedArgType = sycl::nd_item; - wrap_kernel( + wrap_kernel( KernelFunc, Properties, Range); } @@ -2256,7 +2256,7 @@ class __SYCL_EXPORT handler { "member function instead.") void parallel_for_work_group(range NumWorkGroups, PropertiesT Props, const KernelType &KernelFunc) { - wrap_kernel>, Dims, /*SetNumWorkGroups=*/true>(KernelFunc, Props, NumWorkGroups); } @@ -2270,7 +2270,7 @@ class __SYCL_EXPORT handler { void parallel_for_work_group(range NumWorkGroups, range WorkGroupSize, PropertiesT Props, const KernelType &KernelFunc) { - wrap_kernel>, Dims>( KernelFunc, Props, nd_range{NumWorkGroups * WorkGroupSize, WorkGroupSize}); @@ -3642,7 +3642,7 @@ class __SYCL_EXPORT handler { void instantiateKernelOnHost(void *InstantiateKernelOnHostPtr); friend class detail::HandlerAccess; - template + template friend struct detail::KernelWrapper; #ifdef __INTEL_PREVIEW_BREAKING_CHANGES diff --git a/sycl/include/sycl/kernel_launch_helper.hpp b/sycl/include/sycl/kernel_launch_helper.hpp index 7054a5e137ef6..101f27affcd76 100644 --- a/sycl/include/sycl/kernel_launch_helper.hpp +++ b/sycl/include/sycl/kernel_launch_helper.hpp @@ -22,12 +22,10 @@ namespace sycl { inline namespace _V1 { +namespace detail { -// TODO: Move WrapAs to detail:: namespace as well and move this header to -// include/sycl/detail directory. enum class WrapAs { single_task, parallel_for, parallel_for_work_group }; -namespace detail { // Helper for merging properties with ones defined in an optional kernel functor // getter. template From cdfd679dd5dd9e63a8211981fbe41189c0f87561 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Mon, 21 Jul 2025 19:39:27 +0200 Subject: [PATCH 09/11] Move kernel_launch_wrapper to detail --- .../sycl/{ => detail}/kernel_launch_helper.hpp | 0 sycl/include/sycl/handler.hpp | 2 +- .../handler/unnamed-lambda-negative.cpp | 14 +++++++------- sycl/test/esimd/global_var.cpp | 2 +- .../properties/properties_kernel_negative.cpp | 2 +- sycl/test/include_deps/sycl_detail_core.hpp.cpp | 2 +- 6 files changed, 11 insertions(+), 11 deletions(-) rename sycl/include/sycl/{ => detail}/kernel_launch_helper.hpp (100%) diff --git a/sycl/include/sycl/kernel_launch_helper.hpp b/sycl/include/sycl/detail/kernel_launch_helper.hpp similarity index 100% rename from sycl/include/sycl/kernel_launch_helper.hpp rename to sycl/include/sycl/detail/kernel_launch_helper.hpp diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 9a6cd2a192f69..daf8ce672fd92 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -41,7 +41,7 @@ #include #include #include -#include +#include #include #include #include diff --git a/sycl/test/basic_tests/handler/unnamed-lambda-negative.cpp b/sycl/test/basic_tests/handler/unnamed-lambda-negative.cpp index f250854f64e9c..bff98b8a4679e 100644 --- a/sycl/test/basic_tests/handler/unnamed-lambda-negative.cpp +++ b/sycl/test/basic_tests/handler/unnamed-lambda-negative.cpp @@ -8,26 +8,26 @@ int main() { sycl::queue q; - // expected-error-re@sycl/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} q.single_task([=](){}); - // expected-error-re@sycl/kernel_launch_helper.hpp:* {{unnamed type 'sycl::detail::RoundedRangeKernel<{{.*}}>' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} - // expected-error-re@sycl/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{unnamed type 'sycl::detail::RoundedRangeKernel<{{.*}}>' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} q.parallel_for(sycl::range<1>{1}, [=](sycl::item<1>) {}); q.submit([&](sycl::handler &cgh) { - // expected-error-re@sycl/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} cgh.single_task([=](){}); - // expected-error-re@sycl/kernel_launch_helper.hpp:* {{unnamed type 'sycl::detail::RoundedRangeKernel<{{.*}}>' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} - // expected-error-re@sycl/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{unnamed type 'sycl::detail::RoundedRangeKernel<{{.*}}>' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} cgh.parallel_for(sycl::range<1>{1}, [=](sycl::item<1>) {}); - // expected-error-re@sycl/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} + // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{unnamed type '{{.*}}' is invalid; provide a kernel name, or use '-fsycl-unnamed-lambda' to enable unnamed kernel lambdas}} // expected-note@+1 {{in instantiation of function template}} cgh.parallel_for_work_group(sycl::range<1>{1}, [=](sycl::group<1>) {}); }); diff --git a/sycl/test/esimd/global_var.cpp b/sycl/test/esimd/global_var.cpp index 2002cc82f4af4..d725f6dd1013f 100644 --- a/sycl/test/esimd/global_var.cpp +++ b/sycl/test/esimd/global_var.cpp @@ -42,7 +42,7 @@ void kernel_call() { q.submit([&](sycl::handler &cgh) { cgh.parallel_for(nd_range<1>(1, 1), [=](nd_item<1> ndi) { - //expected-note@sycl/kernel_launch_helper.hpp:* 2{{called by 'kernel_parallel_for}} + //expected-note@sycl/detail/kernel_launch_helper.hpp:* 2{{called by 'kernel_parallel_for}} //expected-error@+1{{ESIMD globals cannot be used in a SYCL context}} vc = 0; //expected-note@+1{{called by}} diff --git a/sycl/test/extensions/properties/properties_kernel_negative.cpp b/sycl/test/extensions/properties/properties_kernel_negative.cpp index 33b20ed9961ac..e3c7314182208 100644 --- a/sycl/test/extensions/properties/properties_kernel_negative.cpp +++ b/sycl/test/extensions/properties/properties_kernel_negative.cpp @@ -399,7 +399,7 @@ struct TestKernelNoGetter { void check_non_const_getter_warning() { sycl::queue Q; - // expected-error-re@sycl/kernel_launch_helper.hpp:* {{static assertion failed due to requirement {{.+}}: get(sycl::ext::oneapi::experimental::properties_tag) member in kernel functor class must be declared as a const member function}} + // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{static assertion failed due to requirement {{.+}}: get(sycl::ext::oneapi::experimental::properties_tag) member in kernel functor class must be declared as a const member function}} Q.single_task(TestKernelNonConstGetter()); // No error expected for kernel functor with a const get(properties_tag) diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp index 2ef598ae49659..cc80138dd235a 100644 --- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp +++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp @@ -142,7 +142,7 @@ // CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp // CHECK-NEXT: kernel.hpp -// CHECK-NEXT: kernel_launch_helper.hpp +// CHECK-NEXT: detail/kernel_launch_helper.hpp // CHECK-NEXT: ext/intel/experimental/fp_control_kernel_properties.hpp // CHECK-NEXT: ext/intel/experimental/kernel_execution_properties.hpp // CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp From e9d4eae213f47f079459fdee7fdafbc4d31cfe1d Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Mon, 21 Jul 2025 19:45:39 +0200 Subject: [PATCH 10/11] Update doc + rename KernelWrapperSingletonFuncs --- sycl/doc/design/CompileTimeProperties.md | 4 +- .../sycl/detail/kernel_launch_helper.hpp | 4 +- sycl/include/sycl/handler.hpp | 48 ++++++++++--------- 3 files changed, 29 insertions(+), 27 deletions(-) diff --git a/sycl/doc/design/CompileTimeProperties.md b/sycl/doc/design/CompileTimeProperties.md index 9c65356329ac9..02f7580940717 100644 --- a/sycl/doc/design/CompileTimeProperties.md +++ b/sycl/doc/design/CompileTimeProperties.md @@ -362,11 +362,11 @@ string if it is not already a string. [9]: **NOTE**: The intention is to replace the existing member functions like -`handler::kernel_single_task()` with wrapper classes like +`detail::KernelWrapperHelperFuncs::kernel_single_task()` with wrapper classes like `KernelSingleTaskWrapper`. We believe this will not cause problems for the device compiler front-end because it recognizes kernel functions via the `[[clang::sycl_kernel]]` attribute, not by the name -`handler::kernel_single_task()`. +`detail::KernelWrapperHelperFuncs::kernel_single_task()`. ## Properties on a non-global variable type diff --git a/sycl/include/sycl/detail/kernel_launch_helper.hpp b/sycl/include/sycl/detail/kernel_launch_helper.hpp index 101f27affcd76..f90d0c4efd497 100644 --- a/sycl/include/sycl/detail/kernel_launch_helper.hpp +++ b/sycl/include/sycl/detail/kernel_launch_helper.hpp @@ -48,7 +48,7 @@ struct GetMergedKernelProperties< PropertiesT, get_method_properties>; }; -struct KernelWrapperSingletonFuncs { +struct KernelWrapperHelperFuncs { #ifdef SYCL_LANGUAGE_VERSION #ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS @@ -207,7 +207,7 @@ struct KernelWrapper< WrapAsVal, KernelName, KernelType, ElementType, PropertyProcessor, PropertiesT, ext::oneapi::experimental::detail::properties_t> - : public KernelWrapperSingletonFuncs { + : public KernelWrapperHelperFuncs { static void wrap([[maybe_unused]] PropertyProcessor h, [[maybe_unused]] const KernelType &KernelFunc) { diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index daf8ce672fd92..778f178bc537b 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -17,6 +17,7 @@ #include #include #include +#include #include #include #include @@ -41,7 +42,6 @@ #include #include #include -#include #include #include #include @@ -1292,9 +1292,9 @@ class __SYCL_EXPORT handler { using KName = std::conditional_t::value, decltype(Wrapper), NameWT>; - detail::KernelWrapper::wrap(this, Wrapper); + detail::KernelWrapper::wrap(this, Wrapper); #ifndef __SYCL_DEVICE_ONLY__ constexpr detail::string_view Name{detail::getKernelName()}; verifyUsedKernelBundleInternal(Name); @@ -1388,8 +1388,8 @@ class __SYCL_EXPORT handler { } template < - detail::WrapAs WrapAsVal, typename KernelName, typename ElementType = void, - int Dims = 1, bool SetNumWorkGroups = false, + detail::WrapAs WrapAsVal, typename KernelName, + typename ElementType = void, int Dims = 1, bool SetNumWorkGroups = false, typename PropertiesT = ext::oneapi::experimental::empty_properties_t, typename KernelType, typename... RangeParams> void wrap_kernel(const KernelType &KernelFunc, const PropertiesT &Props, @@ -1427,8 +1427,8 @@ class __SYCL_EXPORT handler { // Implementation for something that had to be removed long ago but now stuck // until next major release... template < - detail::WrapAs WrapAsVal, typename KernelName, typename ElementType = void, - int Dims = 1, bool SetNumWorkGroups = false, + detail::WrapAs WrapAsVal, typename KernelName, + typename ElementType = void, int Dims = 1, bool SetNumWorkGroups = false, typename PropertiesT = ext::oneapi::experimental::empty_properties_t, typename KernelType, typename... RangeParams> void wrap_kernel_legacy(const KernelType &KernelFunc, kernel &Kernel, @@ -1717,8 +1717,8 @@ class __SYCL_EXPORT handler { /// \param KernelFunc is a SYCL kernel function. template void single_task(const KernelType &KernelFunc) { - wrap_kernel(KernelFunc, {} /*Props*/, - range<1>{1}); + wrap_kernel( + KernelFunc, {} /*Props*/, range<1>{1}); } template @@ -1784,8 +1784,8 @@ class __SYCL_EXPORT handler { using TransformedArgType = std::conditional_t< std::is_integral::value && Dims == 1, item, typename TransformUserItemType::type>; - wrap_kernel( - KernelFunc, {} /*Props*/, NumWorkItems, WorkItemOffset); + wrap_kernel(KernelFunc, {} /*Props*/, NumWorkItems, WorkItemOffset); } /// Hierarchical kernel invocation method of a kernel defined as a lambda @@ -1915,7 +1915,7 @@ class __SYCL_EXPORT handler { using NameT = typename detail::get_kernel_name_t::name; (void)Kernel; - detail::KernelWrapperSingletonFuncs::kernel_single_task(KernelFunc); + detail::KernelWrapperHelperFuncs::kernel_single_task(KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ throwIfActionIsCreated(); constexpr detail::string_view Name{detail::getKernelName()}; @@ -1951,8 +1951,8 @@ class __SYCL_EXPORT handler { // Ignore any set kernel bundles and use the one associated with the kernel setHandlerKernelBundle(Kernel); using LambdaArgType = sycl::detail::lambda_arg_type>; - wrap_kernel_legacy( - KernelFunc, Kernel, {} /*Props*/, NumWorkItems); + wrap_kernel_legacy(KernelFunc, Kernel, {} /*Props*/, NumWorkItems); } /// Defines and invokes a SYCL kernel function for the specified range and @@ -1970,8 +1970,9 @@ class __SYCL_EXPORT handler { void parallel_for(kernel Kernel, range NumWorkItems, id WorkItemOffset, const KernelType &KernelFunc) { using LambdaArgType = sycl::detail::lambda_arg_type>; - wrap_kernel_legacy( - KernelFunc, Kernel, {} /*Props*/, NumWorkItems, WorkItemOffset); + wrap_kernel_legacy(KernelFunc, Kernel, {} /*Props*/, NumWorkItems, + WorkItemOffset); } /// Defines and invokes a SYCL kernel function for the specified range and @@ -1990,8 +1991,8 @@ class __SYCL_EXPORT handler { const KernelType &KernelFunc) { using LambdaArgType = sycl::detail::lambda_arg_type>; - wrap_kernel_legacy( - KernelFunc, Kernel, {} /*Props*/, NDRange); + wrap_kernel_legacy(KernelFunc, Kernel, {} /*Props*/, NDRange); } /// Hierarchical kernel invocation method of a kernel. @@ -2061,7 +2062,7 @@ class __SYCL_EXPORT handler { PropertiesT>::value> single_task(PropertiesT Props, const KernelType &KernelFunc) { wrap_kernel(KernelFunc, Props, - range<1>{1}); + range<1>{1}); } template ; - wrap_kernel( - KernelFunc, Properties, Range); + wrap_kernel(KernelFunc, Properties, Range); } /// Reductions @{ @@ -3642,7 +3643,8 @@ class __SYCL_EXPORT handler { void instantiateKernelOnHost(void *InstantiateKernelOnHostPtr); friend class detail::HandlerAccess; - template + template friend struct detail::KernelWrapper; #ifdef __INTEL_PREVIEW_BREAKING_CHANGES From 085ad44118c1eac7b34444b04613646c9db52ed9 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Tue, 22 Jul 2025 16:30:06 +0200 Subject: [PATCH 11/11] Update test/include_deps/sycl_detail_core.hpp.cpp --- sycl/test/include_deps/sycl_detail_core.hpp.cpp | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp index cc80138dd235a..d050d3255a90e 100644 --- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp +++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp @@ -130,6 +130,13 @@ // CHECK-NEXT: CL/cl_platform.h // CHECK-NEXT: CL/cl_ext.h // CHECK-NEXT: detail/id_queries_fit_in_int.hpp +// CHECK-NEXT: detail/kernel_launch_helper.hpp +// CHECK-NEXT: ext/intel/experimental/fp_control_kernel_properties.hpp +// CHECK-NEXT: ext/intel/experimental/kernel_execution_properties.hpp +// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp +// CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp +// CHECK-NEXT: detail/sycl_local_mem_builtins.hpp // CHECK-NEXT: detail/kernel_name_based_cache.hpp // CHECK-NEXT: detail/kernel_name_str_t.hpp // CHECK-NEXT: detail/reduction_forward.hpp @@ -140,14 +147,7 @@ // CHECK-NEXT: ext/oneapi/experimental/cluster_group_prop.hpp // CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp // CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp -// CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp // CHECK-NEXT: kernel.hpp -// CHECK-NEXT: detail/kernel_launch_helper.hpp -// CHECK-NEXT: ext/intel/experimental/fp_control_kernel_properties.hpp -// CHECK-NEXT: ext/intel/experimental/kernel_execution_properties.hpp -// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp -// CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp -// CHECK-NEXT: detail/sycl_local_mem_builtins.hpp // CHECK-NEXT: sampler.hpp // CHECK-NEXT: feature_test.hpp // CHECK-EMPTY: