summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
-rw-r--r--kernels/compiler_reqd_sub_group_size.cl5
-rw-r--r--utests/CMakeLists.txt1
-rw-r--r--utests/compiler_reqd_sub_group_size.cpp46
-rw-r--r--utests/utest_helper.cpp20
-rw-r--r--utests/utest_helper.hpp3
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__ */