diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 04b535619dcb4..6bba881466fa7 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -3288,6 +3288,10 @@ def err_attribute_argument_is_zero : Error< def warn_attribute_argument_n_negative : Warning< "%0 attribute parameter %1 is negative and will be ignored">, InGroup; +def warn_reqd_sub_group_attribute_cuda_n_32 + : Warning<"attribute argument %0 is invalid and will be ignored; CUDA " + "requires sub_group size 32">, + InGroup; def err_property_function_in_objc_container : Error< "use of Objective-C property in function nested in Objective-C " "container not supported, move function outside its container">; diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 321efe710af02..fc92779a5faa8 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3886,6 +3886,10 @@ void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI, << CI << /*positive*/ 0; return; } + if (Context.getTargetInfo().getTriple().isNVPTX() && ArgVal != 32) { + Diag(E->getExprLoc(), diag::warn_reqd_sub_group_attribute_cuda_n_32) + << ArgVal.getSExtValue(); + } // Check to see if there's a duplicate attribute with different values // already applied to the declaration. diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-cuda.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-cuda.cpp new file mode 100644 index 0000000000000..2d9087ddee31d --- /dev/null +++ b/clang/test/SemaSYCL/reqd-sub-group-size-cuda.cpp @@ -0,0 +1,20 @@ +// RUN: %clang_cc1 -fsycl-is-device -triple nvptx -internal-isystem %S/Inputs -std=c++2b -verify %s +// +// This tests that a warning is returned when a sub group size other than 32 is +// requested in the CUDA backend via the reqd_sub_group_size() kernel attribute. +#include "sycl.hpp" + +int main() { + + sycl::queue Q; + + Q.submit([&](sycl::handler &h) { + h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; CUDA requires sub_group size 32}} + }); + + Q.submit([&](sycl::handler &h) { + h.single_task([=] [[sycl::reqd_sub_group_size(32)]] {}); + }); + + return 0; +}