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
