One minor comments, On Fri, Jan 17, 2014 at 04:20:02PM +0800, Yongjia Zhang wrote: > utests: compiler_private_data_overflow is aimed at checking out whether > the private data stack size is smaller than 64. This test case is not to test stack size which is smaller than 64 but test with a stack larger than 64*4*4 = 1KB size per lane. Our previous implementation only allocates a 1KB stach per lane no matter how much stack is used actually, so this new case will fail.
Will push it latter, thanks. > > Signed-off-by: Yongjia Zhang<[email protected]> > --- > kernels/compiler_private_data_overflow.cl | 10 ++++++++++ > utests/CMakeLists.txt | 1 + > utests/compiler_private_data_overflow.cpp | 15 +++++++++++++++ > 3 files changed, 26 insertions(+) > create mode 100644 kernels/compiler_private_data_overflow.cl > create mode 100644 utests/compiler_private_data_overflow.cpp > > diff --git a/kernels/compiler_private_data_overflow.cl > b/kernels/compiler_private_data_overflow.cl > new file mode 100644 > 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 8c64844..47ea9ce 100644 > --- a/utests/CMakeLists.txt > +++ b/utests/CMakeLists.txt > @@ -165,6 +165,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 > 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 ); > -- > 1.8.3.2 > > > _______________________________________________ > Beignet mailing list > [email protected] > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
