This patch LGTM, will push after the first 2 patches got merged. Thanks, Zhigang Gong.
On Mon, May 11, 2015 at 02:02:54PM +0800, Yang Rong wrote: > 1. Enable compiler_argument_structure_indirect. > 2. Add compiler_argument_structure_indirect, which has select address and > load argument instruction. > > Signed-off-by: Yang Rong <[email protected]> > --- > kernels/compiler_argument_structure_indirect.cl | 4 +-- > kernels/compiler_argument_structure_select.cl | 18 ++++++++++++ > utests/CMakeLists.txt | 2 ++ > utests/compiler_argument_structure_indirect.cpp | 7 +++-- > utests/compiler_argument_structure_select.cpp | 37 > +++++++++++++++++++++++++ > 5 files changed, 63 insertions(+), 5 deletions(-) > create mode 100644 kernels/compiler_argument_structure_select.cl > create mode 100644 utests/compiler_argument_structure_select.cpp > > diff --git a/kernels/compiler_argument_structure_indirect.cl > b/kernels/compiler_argument_structure_indirect.cl > index c4b062f..6fd7873 100644 > --- a/kernels/compiler_argument_structure_indirect.cl > +++ b/kernels/compiler_argument_structure_indirect.cl > @@ -1,7 +1,7 @@ > -struct hop { int x[16]; }; > +struct hop { int a, x[16]; }; > > __kernel void > -compiler_argument_structure(__global int *dst, struct hop h) > +compiler_argument_structure_indirect(__global int *dst, struct hop h) > { > int id = (int)get_global_id(0); > dst[id] = h.x[get_local_id(0)]; > diff --git a/kernels/compiler_argument_structure_select.cl > b/kernels/compiler_argument_structure_select.cl > new file mode 100644 > index 0000000..295acf4 > --- /dev/null > +++ b/kernels/compiler_argument_structure_select.cl > @@ -0,0 +1,18 @@ > +typedef struct { > + int offset; > + int threshold0; > + int threshold1; > +}hop; > + > +__kernel void compiler_argument_structure_select(__global int *dst, hop h) > +{ > + int i = get_global_id (0); > + int threshold=0; > + if (i == 0) { > + threshold = h.threshold0; > + } else { > + threshold = h.threshold1; > + } > + dst[i] = threshold; > +} > + > diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt > index dcb3385..c9d96db 100644 > --- a/utests/CMakeLists.txt > +++ b/utests/CMakeLists.txt > @@ -39,6 +39,8 @@ set (utests_sources > compiler_box_blur.cpp > compiler_insert_to_constant.cpp > compiler_argument_structure.cpp > + compiler_argument_structure_indirect.cpp > + compiler_argument_structure_select.cpp > compiler_arith_shift_right.cpp > compiler_mixed_pointer.cpp > compiler_array0.cpp > diff --git a/utests/compiler_argument_structure_indirect.cpp > b/utests/compiler_argument_structure_indirect.cpp > index a4584d5..b54432e 100644 > --- a/utests/compiler_argument_structure_indirect.cpp > +++ b/utests/compiler_argument_structure_indirect.cpp > @@ -1,6 +1,6 @@ > #include "utest_helper.hpp" > > -struct hop { int x[16]; }; > +struct hop { int a, x[16]; }; > > void compiler_argument_structure_indirect(void) > { > @@ -21,8 +21,9 @@ void compiler_argument_structure_indirect(void) > OCL_MAP_BUFFER(0); > > // Check results > - for (uint32_t i = 0; i < n; ++i) > - OCL_ASSERT(((uint32_t*)buf_data[0])[i] == 7); > + for (uint32_t i = 0; i < n; ++i ) { > + OCL_ASSERT(((uint32_t*)buf_data[0])[i] == (i%16)); > + } > } > > MAKE_UTEST_FROM_FUNCTION(compiler_argument_structure_indirect); > diff --git a/utests/compiler_argument_structure_select.cpp > b/utests/compiler_argument_structure_select.cpp > new file mode 100644 > index 0000000..b46e745 > --- /dev/null > +++ b/utests/compiler_argument_structure_select.cpp > @@ -0,0 +1,37 @@ > +#include "utest_helper.hpp" > + > +struct hop{ > + int offset; > + int threshold0; > + int threshold1; > +}; > + > +void compiler_argument_structure_select(void) > +{ > + const size_t n = 2048; > + hop h; > + > + // Setup kernel and buffers > + OCL_CREATE_KERNEL("compiler_argument_structure_select"); > + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); > + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); > + h.offset = 2; > + h.threshold0 = 5; > + h.threshold1 = 7; > + OCL_SET_ARG(1, sizeof(hop), &h); > + > + // Run the kernel > + globals[0] = n; > + locals[0] = 16; > + OCL_NDRANGE(1); > + OCL_MAP_BUFFER(0); > + > + // Check results > + OCL_ASSERT(((uint32_t*)buf_data[0])[0] == 5); > + for (uint32_t i = 1; i < n; ++i ) { > + OCL_ASSERT(((uint32_t*)buf_data[0])[i] == 7); > + } > +} > + > +MAKE_UTEST_FROM_FUNCTION(compiler_argument_structure_select); > + > -- > 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
