[Beignet] [PATCH 3/3] Utset: Add test case for cl_intel_required_subgroup_size extension
Yang, Rong R
rong.r.yang at intel.com
Tue Jun 13 08:39:57 UTC 2017
> -----Original Message-----
> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
> Xiuli Pan
> Sent: Monday, June 5, 2017 16:28
> To: beignet at lists.freedesktop.org
> Cc: Pan, Xiuli <xiuli.pan at intel.com>
> Subject: [Beignet] [PATCH 3/3] Utset: Add test case for
> cl_intel_required_subgroup_size extension
>
> From: Pan Xiuli <xiuli.pan at 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.
>
> Signed-off-by: Pan Xiuli <xiuli.pan at intel.com>
> ---
> kernels/compiler_reqd_sub_group_size.cl | 5 ++++
> utests/CMakeLists.txt | 1 +
> utests/compiler_reqd_sub_group_size.cpp | 45
> +++++++++++++++++++++++++++++++++
> utests/utest_helper.cpp | 20 +++++++++++++++
> utests/utest_helper.hpp | 3 +++
> 5 files changed, 74 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
> b7ef742..07df957 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -286,6 +286,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..5e2545b
> --- /dev/null
> +++ b/utests/compiler_reqd_sub_group_size.cpp
> @@ -0,0 +1,45 @@
> +#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];
Forgot to free it?
More information about the Beignet
mailing list