Add utest compiler_private_data_overflow

utests: compiler_private_data_overflow is aimed to hit a larger than
1KB stack. It will fail with the old beignet which allocate 1KB stack
size no matter the actual usage of stack in the kernel.

Signed-off-by: Yongjia Zhang<zhang_yong_jia@126.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

diff --git a/kernels/compiler_private_data_overflow.cl b/kernels/compiler_private_data_overflow.cl
new file mode 100644 (file)
index 0000000..d0f557d
--- /dev/null
+++ b/kernels/compiler_private_data_overflow.cl
@@ -0,0 +1,10 @@
+kernel void compiler_private_data_overflow( __global int4 *output )
+{
+       int4 data[65];
+       for( int i=0; i<65; ++i )
+       {
+               data[i] = (int4)i;
+       }
+       if( get_global_id(0) == 1 )
+               *output = data[0];
+}

diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 07945c1..0614ee6 100644 (file)
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -166,6 +166,7 @@ set (utests_sources
   compiler_long_cmp.cpp
   compiler_function_argument3.cpp
   compiler_bool_cross_basic_block.cpp
+  compiler_private_data_overflow.cpp
   load_program_from_bin.cpp
   enqueue_copy_buf.cpp
   utest_assert.cpp

diff --git a/utests/compiler_private_data_overflow.cpp b/utests/compiler_private_data_overflow.cpp
new file mode 100644 (file)
index 0000000..0fa30a0
--- /dev/null
+++ b/utests/compiler_private_data_overflow.cpp
@@ -0,0 +1,15 @@
+#include "utest_helper.hpp"
+
+void compiler_private_data_overflow(void)
+{
+       OCL_CREATE_KERNEL( "compiler_private_data_overflow" );
+       OCL_CREATE_BUFFER( buf[0], 0, sizeof(cl_int4), NULL );
+       OCL_SET_ARG( 0, sizeof(cl_mem), &buf[0] );
+       globals[0] = 64;
+       locals[0] = 32;
+       OCL_NDRANGE(1);
+       OCL_MAP_BUFFER(0);
+       OCL_ASSERT( ((uint32_t *)buf_data[0])[0] == 0 );
+       OCL_UNMAP_BUFFER(0);
+}
+MAKE_UTEST_FROM_FUNCTION( compiler_private_data_overflow );