From: Pan Xiuli <xiuli....@intel.com> Check the device supported subgroup sizes, and use intel_reqd_sub_group_size to build kernels in these size. Then check if there is spill for each kernel. V2: Fix memory leak
Signed-off-by: Pan Xiuli <xiuli....@intel.com> --- kernels/compiler_reqd_sub_group_size.cl | 5 ++++ utests/CMakeLists.txt | 1 + utests/compiler_reqd_sub_group_size.cpp | 46 +++++++++++++++++++++++++++++++++ utests/utest_helper.cpp | 20 ++++++++++++++ utests/utest_helper.hpp | 3 +++ 5 files changed, 75 insertions(+) create mode 100644 kernels/compiler_reqd_sub_group_size.cl create mode 100644 utests/compiler_reqd_sub_group_size.cpp diff --git a/kernels/compiler_reqd_sub_group_size.cl b/kernels/compiler_reqd_sub_group_size.cl new file mode 100644 index 0000000..0ce70e9 --- /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 afef07f..ebbf0f5 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 0000000..37d96fe --- /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, ¶m_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 f4487c1..2e826bc 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, ¶m_value_size); + std::vector<char> param_value(param_value_size); + OCL_CALL(clGetDeviceInfo, device, CL_DEVICE_EXTENSIONS, param_value_size, + param_value.empty() ? NULL : ¶m_value.front(), ¶m_value_size); + if (!param_value.empty()) + extStr = std::string(¶m_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 5dc381e..c304008 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__ */ -- 2.7.4 _______________________________________________ Beignet mailing list Beignet@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/beignet