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..8ecccf7585d22 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::sub_group_size"}, + {D->getAttr(), + "sycl::ext::oneapi::experimental::work_group_size_hint"}, + {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..14bba96021be9 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::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; 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}} 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::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; 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}} 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::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; 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}} 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::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; 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}} 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::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; 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}} 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::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; 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}} 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::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; 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}} 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::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; 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}} 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)]] {}); } 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{