From 08aa9d21f135ce05c5e313ea5afe9bc98b3ce041 Mon Sep 17 00:00:00 2001
From: Dmitry Vodopyanov <dmitry.vodopyanov@intel.com>
Date: Thu, 19 Jan 2023 12:17:57 -0500
Subject: [PATCH 1/3] [SYCL] Add test for sycl::is_compatible supports
 reqd_work_group_size

Impl: https://github.com/intel/llvm/pull/8056
---
 SYCL/OptionalKernelFeatures/is_compatible.cpp | 30 +++++++++++++++++++
 1 file changed, 30 insertions(+)

diff --git a/SYCL/OptionalKernelFeatures/is_compatible.cpp b/SYCL/OptionalKernelFeatures/is_compatible.cpp
index 62e8aa6cb6..6624f6a1e2 100644
--- a/SYCL/OptionalKernelFeatures/is_compatible.cpp
+++ b/SYCL/OptionalKernelFeatures/is_compatible.cpp
@@ -13,6 +13,10 @@
 class KernelCPU;
 class KernelGPU;
 class KernelACC;
+class GoodWGSize;
+class WrongReqWGSize;
+
+constexpr int SIZE = 2;
 
 int main() {
   bool Compatible = true;
@@ -42,5 +46,31 @@ int main() {
     Called = true;
   }
 
+  if (sycl::is_compatible<GoodWGSize>(Dev)) {
+    Q.submit(
+        [&](sycl::handler &h) { 
+          h.parallel_for<class GoodWGSize>(sycl::range<2>(4, 2),
+                   [=](sycl::item<2> it) [[sycl::reqd_work_group_size(SIZE, SIZE)]] {});
+        });
+    Q.wait();
+    Compatible &= (Dev.get_info<sycl::info::device::max_work_group_size>() > (SIZE + SIZE));
+    Called = true;
+  }
+
+  if (sycl::is_compatible<WrongReqWGSize>(Dev)) {
+    Q.submit(
+        [&](sycl::handler &h) { 
+          h.parallel_for<class WrongReqWGSize>(sycl::range<1>(2),
+                   [=](sycl::item<1> it) [[sycl::reqd_work_group_size(INT_MAX)]] {});
+        });
+    Q.wait();
+    Compatible &= false;
+    Called = false;
+  }
+  else {
+    Compatible &= true;
+    Called = true;
+  }
+
   return (Compatible && Called) ? 0 : 1;
 }

From 06773b99620ba4196434a50e1ce342e5642f9838 Mon Sep 17 00:00:00 2001
From: Dmitry Vodopyanov <dmitry.vodopyanov@intel.com>
Date: Thu, 26 Jan 2023 07:53:51 -0500
Subject: [PATCH 2/3] Apply CR comments

---
 SYCL/OptionalKernelFeatures/is_compatible.cpp | 13 +++++--------
 1 file changed, 5 insertions(+), 8 deletions(-)

diff --git a/SYCL/OptionalKernelFeatures/is_compatible.cpp b/SYCL/OptionalKernelFeatures/is_compatible.cpp
index 6624f6a1e2..a2c1108ea8 100644
--- a/SYCL/OptionalKernelFeatures/is_compatible.cpp
+++ b/SYCL/OptionalKernelFeatures/is_compatible.cpp
@@ -53,23 +53,20 @@ int main() {
                    [=](sycl::item<2> it) [[sycl::reqd_work_group_size(SIZE, SIZE)]] {});
         });
     Q.wait();
-    Compatible &= (Dev.get_info<sycl::info::device::max_work_group_size>() > (SIZE + SIZE));
+    Compatible &= (Dev.get_info<sycl::info::device::max_work_group_size>() > (SIZE * SIZE));
     Called = true;
   }
 
+  if (Dev.get_info<sycl::info::device::max_work_group_size>() > INT_MAX) {
+    Compatible &= true;
+  }
   if (sycl::is_compatible<WrongReqWGSize>(Dev)) {
+    assert(false && "sycl::is_compatible<WrongReqWGSize> must be false");
     Q.submit(
         [&](sycl::handler &h) { 
           h.parallel_for<class WrongReqWGSize>(sycl::range<1>(2),
                    [=](sycl::item<1> it) [[sycl::reqd_work_group_size(INT_MAX)]] {});
         });
-    Q.wait();
-    Compatible &= false;
-    Called = false;
-  }
-  else {
-    Compatible &= true;
-    Called = true;
   }
 
   return (Compatible && Called) ? 0 : 1;

From b68c763437840c3216fc9a722414b23da00c3b14 Mon Sep 17 00:00:00 2001
From: Dmitry Vodopyanov <dmitry.vodopyanov@intel.com>
Date: Fri, 27 Jan 2023 07:26:21 -0500
Subject: [PATCH 3/3] Fix clang-format

---
 SYCL/OptionalKernelFeatures/is_compatible.cpp | 23 ++++++++++---------
 1 file changed, 12 insertions(+), 11 deletions(-)

diff --git a/SYCL/OptionalKernelFeatures/is_compatible.cpp b/SYCL/OptionalKernelFeatures/is_compatible.cpp
index a2c1108ea8..91e8d89afb 100644
--- a/SYCL/OptionalKernelFeatures/is_compatible.cpp
+++ b/SYCL/OptionalKernelFeatures/is_compatible.cpp
@@ -47,13 +47,14 @@ int main() {
   }
 
   if (sycl::is_compatible<GoodWGSize>(Dev)) {
-    Q.submit(
-        [&](sycl::handler &h) { 
-          h.parallel_for<class GoodWGSize>(sycl::range<2>(4, 2),
-                   [=](sycl::item<2> it) [[sycl::reqd_work_group_size(SIZE, SIZE)]] {});
-        });
+    Q.submit([&](sycl::handler &h) {
+      h.parallel_for<class GoodWGSize>(
+          sycl::range<2>(4, 2),
+          [=](sycl::item<2> it) [[sycl::reqd_work_group_size(SIZE, SIZE)]] {});
+    });
     Q.wait();
-    Compatible &= (Dev.get_info<sycl::info::device::max_work_group_size>() > (SIZE * SIZE));
+    Compatible &= (Dev.get_info<sycl::info::device::max_work_group_size>() >
+                   (SIZE * SIZE));
     Called = true;
   }
 
@@ -62,11 +63,11 @@ int main() {
   }
   if (sycl::is_compatible<WrongReqWGSize>(Dev)) {
     assert(false && "sycl::is_compatible<WrongReqWGSize> must be false");
-    Q.submit(
-        [&](sycl::handler &h) { 
-          h.parallel_for<class WrongReqWGSize>(sycl::range<1>(2),
-                   [=](sycl::item<1> it) [[sycl::reqd_work_group_size(INT_MAX)]] {});
-        });
+    Q.submit([&](sycl::handler &h) {
+      h.parallel_for<class WrongReqWGSize>(
+          sycl::range<1>(2),
+          [=](sycl::item<1> it) [[sycl::reqd_work_group_size(INT_MAX)]] {});
+    });
   }
 
   return (Compatible && Called) ? 0 : 1;