PR feedback 2.

Author: jchlanda

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12003,7 +12003,7 @@ def err_sycl_special_type_num_init_method : Error<
   "types with 'sycl_special_class' attribute must have one and only one '__init' "
   "method defined">;
 def warn_launch_bounds_is_cuda_specific : Warning<
-  "%0 attribute ignored, only applicable when targetting Nvidia devices">,
+  "%0 attribute ignored, only applicable when targeting Nvidia devices">,
   InGroup<IgnoredAttributes>;
 
 def warn_cuda_maxclusterrank_sm_90 : Warning<

clang/lib/Sema/SemaDeclAttr.cpp

@@ -22,7 +22,6 @@
 #include "clang/AST/Mangle.h"
 #include "clang/AST/RecursiveASTVisitor.h"
 #include "clang/AST/Type.h"
-#include "clang/Basic/AttributeCommonInfo.h"
 #include "clang/Basic/CharInfo.h"
 #include "clang/Basic/Cuda.h"
 #include "clang/Basic/DarwinSDKInfo.h"

clang/test/CodeGenSYCL/lb_sm_90.cpp

@@ -1,10 +1,7 @@
 // RUN: %clang_cc1 -internal-isystem %S/Inputs -triple nvptx-unknown-unknown -target-cpu sm_90 -fsycl-is-device -Wno-c++23-extensions -S -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR
-// RUN: %clang_cc1 -internal-isystem %S/Inputs %s -triple nvptx64-nvidia-cuda -target-cpu sm_90 -fsycl-is-device -fsyntax-only -Wno-c++23-extensions -verify
-// expected-no-diagnostics
 
 // Maximum work groups per multi-processor, mapped to maxclusterrank PTX
-// directive, is an SM_90 feature, make sure that correct metadata is generated
-// and no warnings/errors are issued.
+// directive, is an SM_90 feature, make sure that correct metadata is generated.
 
 #include "sycl.hpp"
 

clang/test/SemaSYCL/lb_sm_70.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -internal-isystem %S/Inputs -triple nvptx-unknown-unknown -target-cpu sm_70 -fsycl-is-device -Wno-c++23-extensions -S -emit-llvm %s -o -ferror-limit=100 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -triple nvptx-unknown-unknown -target-cpu sm_70 -fsycl-is-device -Wno-c++23-extensions -S -emit-llvm %s -o -fsyntax-only -verify %s
 
 // Maximum work groups per multi-processor, mapped to maxclusterrank PTX
 // directive, is an SM_90 feature, make sure that correct warning is issued on

clang/test/SemaSYCL/lb_sm_90.cpp

@@ -0,0 +1,31 @@
+// RUN: %clang_cc1 -internal-isystem %S/Inputs %s -triple nvptx64-nvidia-cuda -target-cpu sm_90 -fsycl-is-device -fsyntax-only -Wno-c++23-extensions -verify
+// expected-no-diagnostics
+
+// Maximum work groups per multi-processor, mapped to maxclusterrank PTX
+// directive, is an SM_90 feature, make sure that no warnings/errors are issued.
+
+#include "sycl.hpp"
+
+template <int N1, int N2, int N3> class Functor {
+public:
+  [[intel::max_work_group_size(1, 1, N1), intel::min_work_groups_per_cu(N2),
+    intel::max_work_groups_per_mp(N3)]] void
+  operator()() const {}
+};
+
+int main() {
+  sycl::queue Q{};
+
+  Q.submit([&](sycl::handler &cgh) {
+    cgh.single_task<class T1>( [=] [[intel::max_work_group_size(1, 1, 256),
+                                          intel::min_work_groups_per_cu(2),
+                                          intel::max_work_groups_per_mp(4)]] (
+) { volatile int A = 42; });
+  });
+
+  Q.submit([&](sycl::handler &cgh) {
+    cgh.single_task<class F>(Functor<512, 8, 16>{});
+  });
+
+  return 0;
+}