diff --git a/SYCL/OptionalKernelFeatures/is_compatible.cpp b/SYCL/OptionalKernelFeatures/is_compatible.cpp
index 05204976be..787c346d79 100644
--- a/SYCL/OptionalKernelFeatures/is_compatible.cpp
+++ b/SYCL/OptionalKernelFeatures/is_compatible.cpp
@@ -1,11 +1,11 @@
 // requires: cpu, gpu, accelerator
-// UNSUPPORTED: hip
-// FIXME: enable the test back, see intel/llvm#8146
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -O0 %s -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
+#include <algorithm>
+
 #include <sycl/sycl.hpp>
 
 [[sycl::device_has(sycl::aspect::cpu)]] void foo(){};
@@ -17,59 +17,113 @@ class KernelGPU;
 class KernelACC;
 class GoodWGSize;
 class WrongReqWGSize;
+class GoodSubGroupSize;
+class WrongReqSubGroupSize;
 
 constexpr int SIZE = 2;
+constexpr int SIZE2 = 32;
 
 int main() {
   bool Compatible = true;
   bool Called = false;
-  sycl::device Dev;
-  sycl::queue Q(Dev);
+  try {
+    sycl::device Dev;
+    sycl::queue Q(Dev);
 
-  if (sycl::is_compatible<KernelCPU>(Dev)) {
-    Q.submit(
-        [&](sycl::handler &h) { h.single_task<KernelCPU>([=]() { foo(); }); });
-    Q.wait();
-    Compatible &= Dev.is_cpu();
-    Called = true;
-  }
-  if (sycl::is_compatible<KernelGPU>(Dev)) {
-    Q.submit(
-        [&](sycl::handler &h) { h.single_task<KernelGPU>([=]() { bar(); }); });
-    Q.wait();
-    Compatible &= Dev.is_gpu();
-    Called = true;
-  }
-  if (sycl::is_compatible<KernelACC>(Dev)) {
-    Q.submit(
-        [&](sycl::handler &h) { h.single_task<KernelACC>([=]() { baz(); }); });
-    Q.wait();
-    Compatible &= Dev.is_accelerator();
-    Called = true;
-  }
+    if (sycl::is_compatible<KernelCPU>(Dev)) {
+      Q.submit([&](sycl::handler &h) {
+        h.single_task<KernelCPU>([=]() { foo(); });
+      });
+      Q.wait();
+      Compatible &= Dev.is_cpu();
+      Called = true;
+    }
+    if (sycl::is_compatible<KernelGPU>(Dev)) {
+      Q.submit([&](sycl::handler &h) {
+        h.single_task<KernelGPU>([=]() { bar(); });
+      });
+      Q.wait();
+      Compatible &= Dev.is_gpu();
+      Called = true;
+    }
+    if (sycl::is_compatible<KernelACC>(Dev)) {
+      Q.submit([&](sycl::handler &h) {
+        h.single_task<KernelACC>([=]() { baz(); });
+      });
+      Q.wait();
+      Compatible &= Dev.is_accelerator();
+      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<GoodWGSize>(Dev)) {
+      Q.submit([&](sycl::handler &h) {
+        h.parallel_for<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 (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)]] {});
-    });
+    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<WrongReqWGSize>(
+            sycl::range<1>(2),
+            [=](sycl::item<1> it) [[sycl::reqd_work_group_size(INT_MAX)]] {});
+      });
+    }
+
+    const auto SupportedSubGroupSizes =
+        Dev.get_info<sycl::info::device::sub_group_sizes>();
+    if (sycl::is_compatible<GoodSubGroupSize>(Dev)) {
+      Q.submit([&](sycl::handler &h) {
+        h.parallel_for<GoodSubGroupSize>(
+            sycl::range<1>(4),
+            [=](sycl::item<1> it) [[sycl::reqd_sub_group_size(SIZE2)]] {});
+      });
+      Q.wait();
+
+      if (!std::any_of(SupportedSubGroupSizes.cbegin(),
+                       SupportedSubGroupSizes.cend(),
+                       [](int i) { return i == SIZE2; })) {
+        Compatible &= false;
+      }
+      Called = true;
+    } else {
+      Compatible &= false;
+      Called = false;
+    }
+    if (std::any_of(SupportedSubGroupSizes.cbegin(),
+                    SupportedSubGroupSizes.cend(),
+                    [](int i) { return i == INT_MAX; })) {
+      assert(false &&
+             "sycl::is_compatible<WrongReqSubGroupSize> must be false");
+    }
+    if (sycl::is_compatible<WrongReqSubGroupSize>(Dev)) {
+      assert(false &&
+             "sycl::is_compatible<WrongReqSubGroupSize> must be false");
+      Q.submit([&](sycl::handler &h) {
+        h.parallel_for<WrongReqSubGroupSize>(
+            sycl::range<1>(2),
+            [=](sycl::item<1> it) [[sycl::reqd_sub_group_size(INT_MAX)]] {});
+      });
+    }
+  } catch (sycl::exception const &E) {
+    assert(E.code() == sycl::errc::build && "unexpected exception code");
+    size_t pos1 = static_cast<std::string>(E.what()).find(
+        "Unsupported required sub group size");
+    size_t pos2 =
+        static_cast<std::string>(E.what()).find("WrongReqSubGroupSize");
+    if (pos1 == std::string::npos || pos2 == std::string::npos) {
+      assert(false && "unexpected exception message");
+    }
   }
 
   return (Compatible && Called) ? 0 : 1;