[SYCL][NFC] remove SYCL 2020 incompatible tests (intel#14261)

This PR removes the SYCL 2017 deprecated tests as a continuation of
intel#14239 in preparation to land intel#13411

clang/test/CodeGenSYCL/integration_header_ppmacros.cpp

@@ -1,7 +1,5 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -sycl-std=2020 -fsycl-int-header=%t.h %s
-// RUN: FileCheck -input-file=%t.h %s --check-prefix=CHECK-SYCL2020
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -sycl-std=2017 -fsycl-int-header=%t.h %s
-// RUN: FileCheck -input-file=%t.h %s --check-prefix=CHECK-SYCL2017
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-int-header=%t.h %s
+// RUN: FileCheck -input-file=%t.h %s
 // RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-range-rounding=disable -fsycl-int-header=%t.h %s
 // RUN: FileCheck -input-file=%t.h %s --check-prefix=CHECK-RANGE
 // RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-range-rounding=force -fsycl-int-header=%t.h %s
@@ -18,19 +16,9 @@
 int main() {
   sycl::kernel_single_task<class first_kernel>([]() {});
 }
-// CHECK-SYCL2020: #ifndef SYCL_LANGUAGE_VERSION
-// CHECK-SYCL2020-NEXT: #define SYCL_LANGUAGE_VERSION 202001
-// CHECK-SYCL2020-NEXT: #endif //SYCL_LANGUAGE_VERSION
-// CHECK-SYCL2020-NOT: #define CL_SYCL_LANGUAGE_VERSION 121
-// CHECK-SYCL2020-NOT: #define SYCL_LANGUAGE_VERSION 201707
-
-// CHECK-SYCL2017: #ifndef CL_SYCL_LANGUAGE_VERSION
-// CHECK-SYCL2017-NEXT: #define CL_SYCL_LANGUAGE_VERSION 121
-// CHECK-SYCL2017-NEXT: #endif //CL_SYCL_LANGUAGE_VERSION
-// CHECK-SYCL2017: #ifndef SYCL_LANGUAGE_VERSION
-// CHECK-SYCL2017-NEXT: #define SYCL_LANGUAGE_VERSION 201707
-// CHECK-SYCL2017-NEXT: #endif //SYCL_LANGUAGE_VERSION
-// CHECK-SYCL2017-NOT: #define SYCL_LANGUAGE_VERSION 202001
+// CHECK: #ifndef SYCL_LANGUAGE_VERSION
+// CHECK-NEXT: #define SYCL_LANGUAGE_VERSION 202001
+// CHECK-NEXT: #endif //SYCL_LANGUAGE_VERSION
 
 // CHECK-RANGE: #ifndef __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__
 // CHECK-RANGE-NEXT: #define __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ 1

clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 
@@ -16,9 +16,6 @@ class Functor {
   [[intel::no_global_work_offset(SIZE)]] void operator()() const {}
 };
 
-template <int N>
-[[intel::no_global_work_offset(N)]] void func() {}
-
 int main() {
   q.submit([&](handler &h) {
     Foo boo;
@@ -32,10 +29,6 @@ int main() {
 
     Functor<1> f;
     h.single_task<class kernel_name4>(f);
-
-    h.single_task<class kernel_name5>([]() {
-      func<1>();
-    });
   });
   return 0;
 }
@@ -44,6 +37,5 @@ int main() {
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !no_global_work_offset ![[NUM5]]
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} ![[NUM4:[0-9]+]]
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !no_global_work_offset ![[NUM5]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !no_global_work_offset ![[NUM5]]
 // CHECK-NOT: ![[NUM4]]  = !{i32 0}
 // CHECK: ![[NUM5]] = !{}

clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 
@@ -16,9 +16,6 @@ class Functor {
   [[intel::max_global_work_dim(SIZE)]] void operator()() const {}
 };
 
-template <int N>
-[[intel::max_global_work_dim(N)]] void func() {}
-
 int main() {
   q.submit([&](handler &h) {
     Foo boo;
@@ -29,17 +26,12 @@ int main() {
 
     Functor<2> f;
     h.single_task<class kernel_name3>(f);
-
-    h.single_task<class kernel_name4>([]() {
-      func<2>();
-    });
   });
   return 0;
 }
 
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !max_global_work_dim ![[NUM1:[0-9]+]]
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !max_global_work_dim ![[NUM2:[0-9]+]]
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !max_global_work_dim ![[NUM2]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !max_global_work_dim ![[NUM2]]
 // CHECK: ![[NUM1]] = !{i32 1}
 // CHECK: ![[NUM2]] = !{i32 2}

clang/test/CodeGenSYCL/intel-max-work-group-size.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 
@@ -21,8 +21,6 @@ class Functor {
   [[intel::max_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {}
 };
 
-template <int N, int N1, int N2>
-[[intel::max_work_group_size(N, N1, N2)]] void func() {}
 
 int main() {
   q.submit([&](handler &h) {
@@ -38,9 +36,6 @@ int main() {
     Functor<2, 2, 2> f;
     h.single_task<class kernel_name4>(f);
 
-    h.single_task<class kernel_name5>([]() {
-      func<4, 4, 4>();
-    });
   });
   return 0;
 }
@@ -49,9 +44,7 @@ int main() {
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !max_work_group_size ![[NUM8:[0-9]+]]
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !max_work_group_size ![[NUM6:[0-9]+]]
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !max_work_group_size ![[NUM2:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !max_work_group_size ![[NUM4:[0-9]+]]
 // CHECK: ![[NUM1]] = !{i32 1, i32 1, i32 1}
 // CHECK: ![[NUM8]] = !{i32 8, i32 8, i32 8}
 // CHECK: ![[NUM6]] = !{i32 6, i32 3, i32 1}
 // CHECK: ![[NUM2]] = !{i32 2, i32 2, i32 2}
-// CHECK: ![[NUM4]] = !{i32 4, i32 4, i32 4}

clang/test/CodeGenSYCL/intel-restrict.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device %s -emit-llvm -sycl-std=2017 -triple spir64-unknown-unknown -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device %s -emit-llvm -triple spir64-unknown-unknown -o - | FileCheck %s
 
 template <typename name, typename Func>
 __attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {

clang/test/CodeGenSYCL/kernel-by-reference.cpp

@@ -1,37 +1,23 @@
-// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -sycl-std=2017 -DSYCL2017 %s
-// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -sycl-std=2020 -DSYCL2020 %s
-// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -Wno-sycl-strict -DNODIAG %s
-// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -sycl-std=2020 -Wno-sycl-strict -DNODIAG %s
+// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only %s
+// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -Wno-sycl-strict %s
 
 #include "sycl.hpp"
 
 using namespace sycl;
 
+// expected-no-diagnostics
+
 int simple_add(int i) {
   return i + 1;
 }
 
 // ensure both compile.
 int main() {
   queue q;
-#if defined(SYCL2020)
-  // expected-warning@#KernelSingleTask2017 {{passing kernel functions by value is deprecated in SYCL 2020}}
-  // expected-note@+3 {{in instantiation of function template specialization}}
-#endif
-  q.submit([&](handler &h) {
-    h.single_task_2017<class sycl2017>([]() { simple_add(10); });
-  });
 
-#if defined(SYCL2017)
-  // expected-warning@#KernelSingleTask {{passing of kernel functions by reference is a SYCL 2020 extension}}
-  // expected-note@+3 {{in instantiation of function template specialization}}
-#endif
   q.submit([&](handler &h) {
     h.single_task<class sycl2020>([]() { simple_add(11); });
   });
 
   return 0;
 }
-#if defined(NODIAG)
-// expected-no-diagnostics
-#endif

clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple nvptx-unknown-unknown -target-cpu sm_90 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx-unknown-unknown -target-cpu sm_90 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
 
 // Test correct handling of maximum work group size, minimum work groups per
 // compute unit and maximum work groups per multi-processor attributes, that
@@ -25,15 +25,6 @@ template <int N> class Functor {
   operator()() const {}
 };
 
-template <int N>
-[[intel::max_work_group_size(N, 4, 8), intel::min_work_groups_per_cu(N),
-  intel::max_work_groups_per_mp(N)]] void
-zoo() {}
-
-[[intel::max_work_group_size(2, 4, 8), intel::min_work_groups_per_cu(2),
-  intel::max_work_groups_per_mp(4)]] void
-bar() {}
-
 int main() {
   q.submit([&](handler &h) {
     // Test attribute argument size.
@@ -49,21 +40,13 @@ int main() {
     // Test class template argument.
     Functor<6> f;
     h.single_task<class kernel_name3>(f);
-
-    // Test attribute is propagated.
-    h.single_task<class kernel_name4>([]() { bar(); });
-
-    // Test function template argument.
-    h.single_task<class kernel_name5>([]() { zoo<16>(); });
   });
   return 0;
 }
 
 // CHECK: define dso_local void @{{.*}}kernel_name1() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC:[0-9]+]] !max_work_groups_per_mp ![[MWGPM:[0-9]+]] !max_work_group_size ![[MWGS:[0-9]+]]
 // CHECK: define dso_local void @{{.*}}kernel_name2() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC:[0-9]+]] !max_work_groups_per_mp ![[MWGPM:[0-9]+]] !max_work_group_size ![[MWGS:[0-9]+]]
 // CHECK: define dso_local void @{{.*}}kernel_name3() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC_MWGPM:[0-9]+]] !max_work_groups_per_mp ![[MWGPC_MWGPM]] !max_work_group_size ![[MWGS_2:[0-9]+]]
-// CHECK: define dso_local void @{{.*}}kernel_name4() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC:[0-9]+]] !max_work_groups_per_mp ![[MWGPM:[0-9]+]] !max_work_group_size ![[MWGS:[0-9]+]]
-// CHECK: define dso_local void @{{.*}}kernel_name5() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC_MWGPM_2:[0-9]+]] !max_work_groups_per_mp ![[MWGPC_MWGPM_2]] !max_work_group_size ![[MWGS_3:[0-9]+]]
 
 // CHECK: {{.*}}@{{.*}}kernel_name1, !"maxntidx", i32 8}
 // CHECK: {{.*}}@{{.*}}kernel_name1, !"maxntidy", i32 4}
@@ -95,31 +78,9 @@ int main() {
 // CHECK: {{.*}}@{{.*}}Functor{{.*}}, !"maxntidz", i32 6}
 // CHECK: {{.*}}@{{.*}}Functor{{.*}}, !"minctasm", i32 6}
 // CHECK: {{.*}}@{{.*}}Functor{{.*}}, !"maxclusterrank", i32 6}
-// CHECK: {{.*}}@{{.*}}kernel_name4, !"maxntidx", i32 8}
-// CHECK: {{.*}}@{{.*}}kernel_name4, !"maxntidy", i32 4}
-// CHECK: {{.*}}@{{.*}}kernel_name4, !"maxntidz", i32 2}
-// CHECK: {{.*}}@{{.*}}kernel_name4, !"minctasm", i32 2}
-// CHECK: {{.*}}@{{.*}}kernel_name4, !"maxclusterrank", i32 4}
-// CHECK: {{.*}}@{{.*}}bar{{.*}}, !"maxntidx", i32 8}
-// CHECK: {{.*}}@{{.*}}bar{{.*}}, !"maxntidy", i32 4}
-// CHECK: {{.*}}@{{.*}}bar{{.*}}, !"maxntidz", i32 2}
-// CHECK: {{.*}}@{{.*}}bar{{.*}}, !"minctasm", i32 2}
-// CHECK: {{.*}}@{{.*}}bar{{.*}}, !"maxclusterrank", i32 4}
-// CHECK: {{.*}}@{{.*}}kernel_name5, !"maxntidx", i32 8}
-// CHECK: {{.*}}@{{.*}}kernel_name5, !"maxntidy", i32 4}
-// CHECK: {{.*}}@{{.*}}kernel_name5, !"maxntidz", i32 16}
-// CHECK: {{.*}}@{{.*}}kernel_name5, !"minctasm", i32 16}
-// CHECK: {{.*}}@{{.*}}kernel_name5, !"maxclusterrank", i32 16}
-// CHECK: {{.*}}@{{.*}}zoo{{.*}}, !"maxntidx", i32 8}
-// CHECK: {{.*}}@{{.*}}zoo{{.*}}, !"maxntidy", i32 4}
-// CHECK: {{.*}}@{{.*}}zoo{{.*}}, !"maxntidz", i32 16}
-// CHECK: {{.*}}@{{.*}}zoo{{.*}}, !"minctasm", i32 16}
-// CHECK: {{.*}}@{{.*}}zoo{{.*}}, !"maxclusterrank", i32 16}
 
 // CHECK: ![[MWGPC]] = !{i32 2}
 // CHECK: ![[MWGPM]] = !{i32 4}
 // CHECK: ![[MWGS]] = !{i32 8, i32 4, i32 2}
 // CHECK: ![[MWGPC_MWGPM]] = !{i32 6}
 // CHECK: ![[MWGS_2]] = !{i32 8, i32 4, i32 6}
-// CHECK: ![[MWGPC_MWGPM_2]] = !{i32 16}
-// CHECK: ![[MWGS_3]] = !{i32 8, i32 4, i32 16}

clang/test/CodeGenSYCL/num-simd-work-items.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 
@@ -16,8 +16,6 @@ class Functor {
   [[intel::num_simd_work_items(SIZE)]] void operator()() const {}
 };
 
-template <int N>
-[[intel::num_simd_work_items(N)]] void func() {}
 
 int main() {
   q.submit([&](handler &h) {
@@ -30,18 +28,13 @@ int main() {
     Functor<2> f;
     h.single_task<class kernel_name3>(f);
 
-    h.single_task<class kernel_name4>([]() {
-      func<4>();
-    });
   });
   return 0;
 }
 
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !num_simd_work_items ![[NUM1:[0-9]+]]
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !num_simd_work_items ![[NUM42:[0-9]+]]
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !num_simd_work_items ![[NUM2:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !num_simd_work_items ![[NUM4:[0-9]+]]
 // CHECK: ![[NUM1]] = !{i32 1}
 // CHECK: ![[NUM42]] = !{i32 42}
 // CHECK: ![[NUM2]] = !{i32 2}
-// CHECK: ![[NUM4]] = !{i32 4}

clang/test/CodeGenSYCL/reqd-sub-group-size.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -sycl-std=2017 -triple spir64-unknown-unknown -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple spir64-unknown-unknown -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 
@@ -10,15 +10,6 @@ class Functor16 {
   [[intel::reqd_sub_group_size(16)]] void operator()() const {}
 };
 
-[[intel::reqd_sub_group_size(8)]] void foo() {}
-
-class Functor8 {
-public:
-  void operator()() const {
-    foo();
-  }
-};
-
 template <int SIZE>
 class Functor2 {
 public:
@@ -33,9 +24,6 @@ int main() {
     Functor16 f16;
     h.single_task<class kernel_name1>(f16);
 
-    Functor8 f8;
-    h.single_task<class kernel_name2>(f8);
-
     h.single_task<class kernel_name3>(
         []() [[intel::reqd_sub_group_size(4)]]{});
 
@@ -50,11 +38,8 @@ int main() {
 }
 
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE16:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE8:[0-9]+]]
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE4:[0-9]+]]
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE2:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE2]]
 // CHECK: ![[SGSIZE16]] = !{i32 16}
-// CHECK: ![[SGSIZE8]] = !{i32 8}
 // CHECK: ![[SGSIZE4]] = !{i32 4}
 // CHECK: ![[SGSIZE2]] = !{i32 2}