From 338b39014dd095eea043be021e7bf18c09236a51 Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Wed, 29 Nov 2023 20:21:44 +0100 Subject: [PATCH] [SYCL] Provide a mechanism to silence incorrect sub group size warning (#11991) As a follow up to https://github.com/intel/llvm/pull/11687, this PR adds a mechanism to silence the warning using dedicated switch `-Wno-incorrect-sub-group-size` that is wrapped in the `-Wno-attribute` group. --- clang/include/clang/Basic/DiagnosticGroups.td | 2 ++ .../clang/Basic/DiagnosticSemaKinds.td | 2 +- .../reqd-sub-group-size-silence-warnings.cpp | 24 +++++++++++++++++++ 3 files changed, 27 insertions(+), 1 deletion(-) create mode 100644 clang/test/SemaSYCL/reqd-sub-group-size-silence-warnings.cpp diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td index 35fe7ccf636ae..575c1b18bb8cd 100644 --- a/clang/include/clang/Basic/DiagnosticGroups.td +++ b/clang/include/clang/Basic/DiagnosticGroups.td @@ -780,9 +780,11 @@ def SyclIvdepAttribute : DiagGroup<"ivdep-compat">; def IndependentClassAttribute : DiagGroup<"IndependentClass-attribute">; def UnknownAttributes : DiagGroup<"unknown-attributes">; def IgnoredAttributes : DiagGroup<"ignored-attributes", [SyclIvdepAttribute]>; +def IncorrectSubGroupSize: DiagGroup<"incorrect-sub-group-size">; def AcceptedAttributes : DiagGroup<"accepted-attributes">; def Attributes : DiagGroup<"attributes", [UnknownAttributes, IgnoredAttributes, + IncorrectSubGroupSize, AcceptedAttributes]>; def UnknownSanitizers : DiagGroup<"unknown-sanitizers">; def UnnamedTypeTemplateArgs : DiagGroup<"unnamed-type-template-args", diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index d7d11f7bcf137..3e2f1a0998897 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -3539,7 +3539,7 @@ def warn_attribute_on_direct_kernel_callee_only : Warning<"%0 attribute allowed" def warn_reqd_sub_group_attribute_n : Warning<"attribute argument %0 is invalid and will be ignored; %1 " "requires sub_group size %2">, - InGroup; + InGroup; def warn_nothrow_attribute_ignored : Warning<"'nothrow' attribute conflicts with" " exception specification; attribute ignored">, InGroup; diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-silence-warnings.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-silence-warnings.cpp new file mode 100644 index 0000000000000..500f68bd18ca2 --- /dev/null +++ b/clang/test/SemaSYCL/reqd-sub-group-size-silence-warnings.cpp @@ -0,0 +1,24 @@ +// RUN: %clang_cc1 -fsycl-is-device -triple nvptx -internal-isystem %S/Inputs -std=c++2b -verify -Wno-incorrect-sub-group-size %s +// RUN: %clang_cc1 -fsycl-is-device -triple nvptx -internal-isystem %S/Inputs -std=c++2b -verify -Wno-attributes %s +// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx90a -internal-isystem %S/Inputs -std=c++2b -verify -Wno-incorrect-sub-group-size %s +// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx90a -internal-isystem %S/Inputs -std=c++2b -verify -Wno-attributes %s +// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx1010 -internal-isystem %S/Inputs -std=c++2b -verify -Wno-incorrect-sub-group-size %s +// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx1010 -internal-isystem %S/Inputs -std=c++2b -verify -Wno-attributes %s +// +// Sub group size of 8 is incompatible with both CUDA and HIP, expect it to be +// silenced. Check both the dedicated switch '-Wno-incorrect-sub-group-size' and +// the catch all '-Wno-attributes'. +#include "sycl.hpp" + + +// expected-no-diagnostics +int main() { + + sycl::queue Q; + + Q.submit([&](sycl::handler &h) { + h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); + }); + + return 0; +}