5 files changed, 75 insertions, 0 deletions
diff --git a/kernels/compiler_reqd_sub_group_size.cl b/kernels/compiler_reqd_sub_group_size.cl
new file mode 100644
index 00000000..0ce70e9c
--- /dev/null
+++ b/kernels/compiler_reqd_sub_group_size.cl
@@ -0,0 +1,5 @@
+__attribute__((intel_reqd_sub_group_size(SIMD_SIZE)))
+__kernel void compiler_reqd_sub_group_size(global int* src)
+{
+
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index afef07fc..ebbf0f56 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -289,6 +289,7 @@ set (utests_sources
   compiler_sub_group_shuffle_down.cpp
   compiler_sub_group_shuffle_up.cpp
   compiler_sub_group_shuffle_xor.cpp
+  compiler_reqd_sub_group_size.cpp
   builtin_global_linear_id.cpp
   builtin_local_linear_id.cpp
   multi_queue_events.cpp
diff --git a/utests/compiler_reqd_sub_group_size.cpp b/utests/compiler_reqd_sub_group_size.cpp
new file mode 100644
index 00000000..37d96fe5
--- /dev/null
+++ b/utests/compiler_reqd_sub_group_size.cpp
@@ -0,0 +1,46 @@
+#include "utest_helper.hpp"
+#include<string>
+#include<sstream>
+#include<iostream>
+
+using namespace std;
+
+void compiler_reqd_sub_group_size(void)
+{
+  if (!cl_check_reqd_subgroup())
+    return;
+
+  size_t param_value_size;
+  OCL_CALL(clGetDeviceInfo, device, CL_DEVICE_SUB_GROUP_SIZES_INTEL,
+           0, NULL, &param_value_size);
+
+  size_t* param_value = new size_t[param_value_size];
+  OCL_CALL(clGetDeviceInfo, device, CL_DEVICE_SUB_GROUP_SIZES_INTEL,
+           param_value_size, param_value, NULL);
+
+  const char* opt = "-D SIMD_SIZE=";
+  for( uint32_t i = 0; i < param_value_size / sizeof(size_t) ; ++i)
+  {
+    ostringstream ss;
+    uint32_t simd_size = param_value[i];
+    ss << opt << simd_size;
+    //cout << "options: " << ss.str() << endl;
+    OCL_CALL(cl_kernel_init, "compiler_reqd_sub_group_size.cl", "compiler_reqd_sub_group_size",
+                             SOURCE, ss.str().c_str());
+    size_t SIMD_SIZE = 0;
+    OCL_CALL(utestclGetKernelSubGroupInfoKHR,kernel,device, CL_KERNEL_COMPILE_SUB_GROUP_SIZE_INTEL,0, NULL,sizeof(size_t),&SIMD_SIZE,NULL);
+    //cout << SIMD_SIZE << " with " << simd_size << endl;
+    OCL_ASSERT(SIMD_SIZE == simd_size);
+
+    cl_ulong SPILL_SIZE = 0xFFFFFFFF;
+    OCL_CALL(clGetKernelWorkGroupInfo, kernel, device, CL_KERNEL_SPILL_MEM_SIZE_INTEL, sizeof(cl_ulong), &SPILL_SIZE, NULL);
+    //cout << "spill size: " << SPILL_SIZE << endl;
+    OCL_ASSERT(SPILL_SIZE == 0);
+
+    clReleaseProgram(program);
+    program = NULL;
+  }
+  delete[] param_value;
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_reqd_sub_group_size);
diff --git a/utests/utest_helper.cpp b/utests/utest_helper.cpp
index f4487c13..2e826bc6 100644
--- a/utests/utest_helper.cpp
+++ b/utests/utest_helper.cpp
@@ -1139,3 +1139,23 @@ float as_float(uint32_t i)
   _tmp._uint = i;
   return _tmp._float;
 }
+
+int cl_check_reqd_subgroup(void)
+{
+  if (!cl_check_subgroups())
+    return 0;
+  std::string extStr;
+  size_t param_value_size;
+  OCL_CALL(clGetDeviceInfo, device, CL_DEVICE_EXTENSIONS, 0, 0, &param_value_size);
+  std::vector<char> param_value(param_value_size);
+  OCL_CALL(clGetDeviceInfo, device, CL_DEVICE_EXTENSIONS, param_value_size,
+           param_value.empty() ? NULL : &param_value.front(), &param_value_size);
+  if (!param_value.empty())
+    extStr = std::string(&param_value.front(), param_value_size-1);
+
+  if (std::strstr(extStr.c_str(), "cl_intel_required_subgroup_size") == NULL) {
+    printf("No cl_intel_required_subgroup_size, Skip!");
+    return 0;
+  }
+  return 1;
+}
diff --git a/utests/utest_helper.hpp b/utests/utest_helper.hpp
index 5dc381e3..c3040087 100644
--- a/utests/utest_helper.hpp
+++ b/utests/utest_helper.hpp
@@ -328,4 +328,7 @@ extern float as_float(uint32_t i);
 extern uint32_t as_uint(float f);
 /* Check is intel subgroups short enabled. */
 extern int cl_check_subgroups_short(void);
+
+/* Check is intel_required_subgroup_size enabled. */
+extern int cl_check_reqd_subgroup(void);
 #endif /* __UTEST_HELPER_HPP__ */