Move test to clang/test/SemaSYCL

Author: jchlanda

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -11991,7 +11991,7 @@ def warn_launch_bounds_is_cuda_specific : Warning<
   InGroup<IgnoredAttributes>;
 
 def warn_cuda_maxclusterrank_sm_90 : Warning<
-  "maxclusterrank requires sm_90 or higher, CUDA arch provided: %0, ignoring "
+  "'maxclusterrank' requires sm_90 or higher, CUDA arch provided: %0, ignoring "
   "%1 attribute">, InGroup<IgnoredAttributes>;
 
 def err_bit_int_bad_size : Error<"%select{signed|unsigned}0 _BitInt must "

clang/test/SemaCUDA/launch_bounds.cu

@@ -11,7 +11,7 @@ __launch_bounds__(0x10000000000000000) void TestWayTooBigArg(void); // expected-
 
 __launch_bounds__(-128, 7) void TestNegArg1(void); // expected-warning {{'launch_bounds' attribute parameter 0 is negative and will be ignored}}
 __launch_bounds__(128, -7) void TestNegArg2(void); // expected-warning {{'launch_bounds' attribute parameter 1 is negative and will be ignored}}
-__launch_bounds__(128, 2, -8) void TestNegArg2(void); // expected-warning {{maxclusterrank requires sm_90 or higher, CUDA arch provided: sm_75, ignoring 'launch_bounds' attribute}}
+__launch_bounds__(128, 2, -8) void TestNegArg2(void); // expected-warning {{'maxclusterrank' requires sm_90 or higher, CUDA arch provided: sm_75, ignoring 'launch_bounds' attribute}}
 
 __launch_bounds__(1, 2, 3, 4) void Test4Args(void); // expected-error {{'launch_bounds' attribute takes no more than 3 arguments}}
 __launch_bounds__() void TestNoArgs(void); // expected-error {{'launch_bounds' attribute takes at least 1 argument}}
@@ -49,4 +49,4 @@ __launch_bounds__(Args) void TestTemplateVariadicArgs(void) {} // expected-error
 template <int... Args>
 __launch_bounds__(1, Args) void TestTemplateVariadicArgs2(void) {} // expected-error {{expression contains unexpanded parameter pack 'Args'}}
 
-__launch_bounds__(1, 2, 3) void Test3Args(void); // expected-warning {{maxclusterrank requires sm_90 or higher, CUDA arch provided: sm_75, ignoring 'launch_bounds' attribute}}
+__launch_bounds__(1, 2, 3) void Test3Args(void); // expected-warning {{'maxclusterrank' requires sm_90 or higher, CUDA arch provided: sm_75, ignoring 'launch_bounds' attribute}}

sycl/test/launch_bounds/lb_sm_70.cpp renamed to clang/test/SemaSYCL/lb_sm_70.cpp

@@ -1,62 +1,53 @@
-//  REQUIRES: cuda
+// REQUIRES: cuda
 
-// RUN: not %clangxx -ferror-limit=100 -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_70 -fsyntax-only %s -o - 2>&1 | FileCheck %s
-
-// NOTE: we can not use the `-verify` run alongside
-// `expected-error`/`expected-warnings` as the diagnostics come from the device
-// compilation, which happen in temporary files, while `expected-...` are
-// placed in the main file, causing clang to complain at the file mismatch
+// RUN: %clangxx -ferror-limit=100 -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_70 -fsycl-device-only -fsyntax-only -Xclang -verify %s
 
 #include <sycl/sycl.hpp>
 
 template <int N1, int N2, int N3> class Functor {
 public:
+  // expected-warning@+2 {{'maxclusterrank' requires sm_90 or higher, CUDA arch provided: sm_70, ignoring 'max_work_groups_per_mp' attribute}}
   [[intel::max_work_group_size(1, 1, N1), intel::min_work_groups_per_cu(N2),
     intel::max_work_groups_per_mp(N3)]] void
-  // CHECK: maxclusterrank requires sm_90 or higher, CUDA arch provided: sm_70, ignoring 'max_work_groups_per_mp' attribute
   operator()() const {}
 };
 
 int main() {
   sycl::queue Q{};
 
-  sycl::range<1> Gws(32);
-  sycl::range<1> Lws(32);
-
   Q.submit([&](sycl::handler &cgh) {
+     // expected-warning@+5 {{'maxclusterrank' requires sm_90 or higher, CUDA arch provided: sm_70, ignoring 'max_work_groups_per_mp' attribute}}
      cgh.single_task<class T1>(
-         sycl::nd_range<1>(Gws, Lws),
          [=]()
              [[intel::max_work_group_size(1, 1, 256),
                intel::min_work_groups_per_cu(2),
-               // CHECK: maxclusterrank requires sm_90 or higher, CUDA arch provided: sm_70, ignoring 'max_work_groups_per_mp' attribute
                intel::max_work_groups_per_mp(4)]] { volatile int A = 42; });
 
      constexpr float A = 2.0;
+     // expected-error@+5 {{'min_work_groups_per_cu' attribute requires parameter 0 to be an integer constant}}
+     // expected-warning@+5 {{'maxclusterrank' requires sm_90 or higher, CUDA arch provided: sm_70, ignoring 'max_work_groups_per_mp' attribute}}
      cgh.single_task<class T2>(
          [=]()
              [[intel::max_work_group_size(1, 1, 256),
                intel::min_work_groups_per_cu(A),
-               // CHECK: 'min_work_groups_per_cu' attribute requires parameter 0 to be an integer constant
                intel::max_work_groups_per_mp(4)]] { volatile int A = 42; });
-     // CHECK: maxclusterrank requires sm_90 or higher, CUDA arch provided: sm_70, ignoring 'max_work_groups_per_mp' attribute
 
+     // expected-error@+3 {{'min_work_groups_per_cu' attribute requires parameter 0 to be an integer constant}}
      cgh.single_task<class T3>(
          [=]() [[intel::max_work_group_size(1, 1, 256),
                  intel::min_work_groups_per_cu(2147483647 + 1)]]
-         // CHECK: 'min_work_groups_per_cu' attribute requires parameter 0 to be an integer constant
          { volatile int A = 42; });
 
+     // expected-warning@+4 {{attribute 'min_work_groups_per_cu' is already applied with different arguments}}
+     // expected-note@+2 {{previous attribute is here}}
      cgh.single_task<class T4>([=]() [[intel::max_work_group_size(1, 1, 256),
                                        intel::min_work_groups_per_cu(4),
                                        intel::min_work_groups_per_cu(8)]] {
-       // CHECK: attribute 'min_work_groups_per_cu' is already applied with different arguments
-       // CHECK: note: previous attribute is here
        volatile int A = 42;
      });
 
+     // expected-error@+1 {{'min_work_groups_per_cu' attribute must be greater than 0}}
      cgh.single_task<class T5>([=]() [[intel::min_work_groups_per_cu(-8)]] {
-       // CHECK: 'min_work_groups_per_cu' attribute must be greater than 0
        volatile int A = 42;
      });
    }).wait_and_throw();