[SYCL] Generate warning about kernel argument size on all devices (#2509)

* [SYCL] Generate warning about kernel argument size on all devices

Warning about possible device limitation for kernel argument
size should be generated for all devices.

Signed-off-by: Elizabeth Andrews <[email protected]>

Author: elizabethandrews

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -11004,8 +11004,8 @@ def err_sycl_restrict : Error<
   "nor constant-initialized"
   "}0">;
 def warn_sycl_kernel_too_big_args : Warning<
-  "size of kernel arguments (%0 bytes) exceeds supported maximum of %1 bytes "
-  "on GPU">, InGroup<SyclStrict>;
+  "size of kernel arguments (%0 bytes) may exceed the supported maximum "
+  "of %1 bytes on some devices">, InGroup<SyclStrict>;
 def err_sycl_virtual_types : Error<
   "No class with a vtable can be used in a SYCL kernel or any code included in the kernel">;
 def note_sycl_recursive_function_declared_here: Note<"function implemented using recursion declared here">;

clang/lib/Sema/SemaSYCL.cpp

@@ -58,7 +58,7 @@ enum KernelInvocationKind {
 
 const static std::string InitMethodName = "__init";
 const static std::string FinalizeMethodName = "__finalize";
-constexpr unsigned GPUMaxKernelArgsSize = 2048;
+constexpr unsigned MaxKernelArgsSize = 2048;
 
 namespace {
 
@@ -1697,11 +1697,9 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
       : SyclKernelFieldHandler(S), KernelLoc(Loc) {}
 
   ~SyclKernelArgsSizeChecker() {
-    if (SemaRef.Context.getTargetInfo().getTriple().getSubArch() ==
-        llvm::Triple::SPIRSubArch_gen)
-      if (SizeOfParams > GPUMaxKernelArgsSize)
-        SemaRef.Diag(KernelLoc, diag::warn_sycl_kernel_too_big_args)
-            << SizeOfParams << GPUMaxKernelArgsSize;
+    if (SizeOfParams > MaxKernelArgsSize)
+      SemaRef.Diag(KernelLoc, diag::warn_sycl_kernel_too_big_args)
+          << SizeOfParams << MaxKernelArgsSize;
   }
 
   bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final {

clang/test/SemaSYCL/args-size-overflow.cpp

@@ -1,7 +1,5 @@
-// RUN: %clang_cc1 -fsycl -triple spir64_gen -DGPU -fsycl-is-device -fsyntax-only -verify %s
 // RUN: %clang_cc1 -fsycl -triple spir64 -fsycl-is-device -fsyntax-only -verify %s
-// RUN: %clang_cc1 -fsycl -triple spir64_gen -Wno-sycl-strict -fsycl-is-device -fsyntax-only -verify %s
-// RUN: %clang_cc1 -fsycl -triple spir64_gen -Werror=sycl-strict -DERROR -fsycl-is-device -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsycl -triple spir64 -Werror=sycl-strict -DERROR -fsycl-is-device -fsyntax-only -verify %s
 
 #include "Inputs/sycl.hpp"
 class Foo;
@@ -13,12 +11,10 @@ __attribute__((sycl_kernel)) void kernel(F KernelFunc) {
 
 template <typename Name, typename F>
 void parallel_for(F KernelFunc) {
-#ifdef GPU
-  // expected-warning@+6 {{size of kernel arguments (7994 bytes) exceeds supported maximum of 2048 bytes on GPU}}
-#elif ERROR
-  // expected-error@+4 {{size of kernel arguments (7994 bytes) exceeds supported maximum of 2048 bytes on GPU}}
+#ifdef ERROR
+  // expected-error@+4 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}}
 #else
-  // expected-no-diagnostics
+  // expected-warning@+2 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}}
 #endif
   kernel<Name>(KernelFunc);
 }
@@ -35,8 +31,6 @@ void use() {
     int Array[1991];
   } Args;
   auto L = [=]() { (void)Args; };
-#if defined(GPU) || defined(ERROR)
-  // expected-note@+2 {{in instantiation of function template specialization 'parallel_for<Foo}}
-#endif
+  // expected-note@+1 {{in instantiation of function template specialization 'parallel_for<Foo}}
   parallel_for<Foo>(L);
 }