From 700b5f13a0406d7f57fe7e13bdeef92a9d77dc4a Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 5 Nov 2025 23:18:15 -0800 Subject: [PATCH 1/2] [SYCL][clang] Clarify warning message for potential property conflicts The compiler currently warns when properties are used in conjunction with SYCL 2020 properties, warning about potentially ignored conflicting properties. This commit makes the warning messages a little clearer to help users understand what particular properties may be of concern. Signed-off-by: Larsen, Steffen --- .../clang/Basic/DiagnosticSemaKinds.td | 5 +- clang/lib/Sema/SemaSYCLDeclAttr.cpp | 17 ++-- ...tr-add-ir-attributes-function-conflict.cpp | 80 +++++++++---------- 3 files changed, 54 insertions(+), 48 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 772200a3d5f0b..aa4fa900dcde2 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -13413,8 +13413,9 @@ def err_sycl_add_ir_attribute_invalid_filter : Error< "initializer list in the first argument of %0 must contain only string " "literals">; def warn_sycl_old_and_new_kernel_attributes : Warning< - "kernel has both attribute %0 and kernel properties; conflicting properties " - "are ignored">, InGroup; + "kernel has both attribute %0 and kernel properties; if the kernel " + "properties contains the property \"%1\" it will be ignored">, + InGroup; def err_attribute_argument_is_not_valid : Error< "%0 attribute requires integer constant value 0 or 1">; diff --git a/clang/lib/Sema/SemaSYCLDeclAttr.cpp b/clang/lib/Sema/SemaSYCLDeclAttr.cpp index 58ee6e81534be..c485d755a0def 100644 --- a/clang/lib/Sema/SemaSYCLDeclAttr.cpp +++ b/clang/lib/Sema/SemaSYCLDeclAttr.cpp @@ -3153,14 +3153,19 @@ void SemaSYCL::checkSYCLAddIRAttributesFunctionAttrConflicts(Decl *D) { } // If there are potentially conflicting attributes, we issue a warning. - for (const auto *Attr : std::vector{ - D->getAttr(), - D->getAttr(), - D->getAttr(), - D->getAttr()}) + for (const auto [Attr, PotentialConflictProp] : + std::vector>{ + {D->getAttr(), + "sycl::ext::oneapi::experimental::work_group_size"}, + {D->getAttr(), + "sycl::ext::oneapi::experimental::work_group_size_hint"}, + {D->getAttr(), + "sycl::ext::oneapi::experimental::sub_group_size"}, + {D->getAttr(), + "sycl::ext::oneapi::experimental::device_has"}}) if (Attr) Diag(Attr->getLoc(), diag::warn_sycl_old_and_new_kernel_attributes) - << Attr; + << Attr << PotentialConflictProp; } void SemaSYCL::handleSYCLRegisteredKernels(Decl *D, const ParsedAttr &A) { diff --git a/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp b/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp index 5fec0b4961618..a4c9dac9a999b 100644 --- a/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp +++ b/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp @@ -92,87 +92,87 @@ int main() { EmptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); EmptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::device_has()]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::device_has()]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::device_has()]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::device_has()]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); } From 0ba65091d9495c0c927aff5115a4234bce68a138 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 6 Nov 2025 00:35:03 -0800 Subject: [PATCH 2/2] Fix mistake and address SYCL test Signed-off-by: Larsen, Steffen --- clang/lib/Sema/SemaSYCLDeclAttr.cpp | 4 +-- ...tr-add-ir-attributes-function-conflict.cpp | 32 +++++++++---------- .../properties_kernel_negative_device.cpp | 8 ++--- 3 files changed, 22 insertions(+), 22 deletions(-) diff --git a/clang/lib/Sema/SemaSYCLDeclAttr.cpp b/clang/lib/Sema/SemaSYCLDeclAttr.cpp index c485d755a0def..8ecccf7585d22 100644 --- a/clang/lib/Sema/SemaSYCLDeclAttr.cpp +++ b/clang/lib/Sema/SemaSYCLDeclAttr.cpp @@ -3158,9 +3158,9 @@ void SemaSYCL::checkSYCLAddIRAttributesFunctionAttrConflicts(Decl *D) { {D->getAttr(), "sycl::ext::oneapi::experimental::work_group_size"}, {D->getAttr(), - "sycl::ext::oneapi::experimental::work_group_size_hint"}, - {D->getAttr(), "sycl::ext::oneapi::experimental::sub_group_size"}, + {D->getAttr(), + "sycl::ext::oneapi::experimental::work_group_size_hint"}, {D->getAttr(), "sycl::ext::oneapi::experimental::device_has"}}) if (Attr) diff --git a/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp b/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp index a4c9dac9a999b..14bba96021be9 100644 --- a/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp +++ b/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp @@ -98,13 +98,13 @@ int main() { NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::device_has()]] {}); @@ -119,13 +119,13 @@ int main() { NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapperWithImplicit.kernel_single_task([]() [[sycl::device_has()]] {}); @@ -140,13 +140,13 @@ int main() { NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapperWithFilter.kernel_single_task([]() [[sycl::device_has()]] {}); @@ -161,13 +161,13 @@ int main() { NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} NonemptyWrapperWithImplicitAndFilter.kernel_single_task([]() [[sycl::device_has()]] {}); diff --git a/sycl/test/extensions/properties/properties_kernel_negative_device.cpp b/sycl/test/extensions/properties/properties_kernel_negative_device.cpp index 44718aa5f3125..15034ae34188b 100644 --- a/sycl/test/extensions/properties/properties_kernel_negative_device.cpp +++ b/sycl/test/extensions/properties/properties_kernel_negative_device.cpp @@ -3,7 +3,7 @@ #include template struct KernelFunctorWithWGSizeWithAttr { - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size" it will be ignored}} void operator() [[sycl::reqd_work_group_size(32)]] () const {} auto get(sycl::ext::oneapi::experimental::properties_tag) const { return sycl::ext::oneapi::experimental::properties{ @@ -12,7 +12,7 @@ template struct KernelFunctorWithWGSizeWithAttr { }; template struct KernelFunctorWithWGSizeHintWithAttr { - // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::work_group_size_hint" it will be ignored}} void operator() [[sycl::work_group_size_hint(32)]] () const {} auto get(sycl::ext::oneapi::experimental::properties_tag) const { return sycl::ext::oneapi::experimental::properties{ @@ -21,7 +21,7 @@ template struct KernelFunctorWithWGSizeHintWithAttr { }; template struct KernelFunctorWithSGSizeWithAttr { - // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::sub_group_size" it will be ignored}} void operator() [[sycl::reqd_sub_group_size(32)]] () const {} auto get(sycl::ext::oneapi::experimental::properties_tag) const { return sycl::ext::oneapi::experimental::properties{ @@ -30,7 +30,7 @@ template struct KernelFunctorWithSGSizeWithAttr { }; template struct KernelFunctorWithDeviceHasWithAttr { - // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}} + // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; if the kernel properties contains the property "sycl::ext::oneapi::experimental::device_has" it will be ignored}} void operator() [[sycl::device_has(sycl::aspect::cpu)]] () const {} auto get(sycl::ext::oneapi::experimental::properties_tag) const { return sycl::ext::oneapi::experimental::properties{