Ping. -----Original Message----- From: Pan, Xiuli Sent: Wednesday, February 24, 2016 10:58 AM To: [email protected] Cc: Pan, Xiuli <[email protected]> Subject: [PATCH OCL2.0 4/4] Utest: Add pipe related test
From: Pan Xiuli <[email protected]> Add test case for builtin with user struct type and int type and runtime tset for creatPipe and pipe query. Signed-off-by: Pan Xiuli <[email protected]> --- kernels/compiler_pipe_builtin.cl | 117 +++++++++++++++++++++++++++++++++++++++ utests/CMakeLists.txt | 4 +- utests/compiler_pipe_builtin.cpp | 70 +++++++++++++++++++++++ utests/runtime_pipe_query.cpp | 16 ++++++ 4 files changed, 206 insertions(+), 1 deletion(-) create mode 100644 kernels/compiler_pipe_builtin.cl create mode 100644 utests/compiler_pipe_builtin.cpp create mode 100644 utests/runtime_pipe_query.cpp diff --git a/kernels/compiler_pipe_builtin.cl b/kernels/compiler_pipe_builtin.cl new file mode 100644 index 0000000..4e8dcc4 --- /dev/null +++ b/kernels/compiler_pipe_builtin.cl @@ -0,0 +1,117 @@ +typedef struct{ + int a; + int b; +}mystruct; + +__kernel void compiler_pipe_convenience_write_int(write_only pipe int +p, __global int *src) { + int gid = get_global_id(0); + write_pipe(p, &src[gid]); +} +__kernel void compiler_pipe_convenience_read_int(read_only pipe int p, +__global int *dst) { + int gid = get_global_id(0); + read_pipe(p, &dst[gid]); +} +__kernel void compiler_pipe_convenience_write_mystruct(write_only pipe +mystruct p, __global mystruct *src) { + int gid = get_global_id(0); + write_pipe(p, &src[gid]); +} +__kernel void compiler_pipe_convenience_read_mystruct(read_only pipe +mystruct p, __global mystruct *dst) { + int gid = get_global_id(0); + read_pipe(p, &dst[gid]); +} + +__kernel void compiler_pipe_reserve_write_int(write_only pipe int p, +__global int *src) { + int gid = get_global_id(0); + reserve_id_t res_id = reserve_write_pipe(p, 1); + if(is_valid_reserve_id(res_id)) + { + write_pipe(p, res_id, 0, &src[gid]); + commit_write_pipe(p, res_id); + } +} +__kernel void compiler_pipe_reserve_read_int(read_only pipe int p, +__global int *dst) { + int gid = get_global_id(0); + reserve_id_t res_id = reserve_read_pipe(p, 1); + if(is_valid_reserve_id(res_id)) + { + read_pipe(p, res_id, 0, &dst[gid]); + commit_read_pipe(p, res_id); + } +} +__kernel void compiler_pipe_reserve_write_mystruct(write_only pipe +mystruct p, __global mystruct *src) { + int gid = get_global_id(0); + reserve_id_t res_id = reserve_write_pipe(p, 1); + if(is_valid_reserve_id(res_id)) + { + write_pipe(p, res_id, 0, &src[gid]); + commit_write_pipe(p, res_id); + } +} +__kernel void compiler_pipe_reserve_read_mystruct(read_only pipe +mystruct p, __global mystruct *dst) { + int gid = get_global_id(0); + reserve_id_t res_id = reserve_read_pipe(p, 1); + if(is_valid_reserve_id(res_id)) + { + read_pipe(p, res_id, 0, &dst[gid]); + commit_read_pipe(p, res_id); + } +} + +__kernel void compiler_pipe_workgroup_write_int(write_only pipe int p, +__global int *src) { + int gid = get_global_id(0); + reserve_id_t res_id = work_group_reserve_write_pipe(p, get_local_size(0)); + if(is_valid_reserve_id(res_id)) + { + write_pipe(p, res_id, get_local_id(0), &src[gid]); + work_group_commit_write_pipe(p, res_id); + } +} +__kernel void compiler_pipe_workgroup_read_int(read_only pipe int p, +__global int *dst) { + int gid = get_global_id(0); + reserve_id_t res_id = work_group_reserve_read_pipe(p, get_local_size(0)); + if(is_valid_reserve_id(res_id)) + { + read_pipe(p, res_id, get_local_id(0), &dst[gid]); + work_group_commit_read_pipe(p, res_id); + } +} +__kernel void compiler_pipe_workgroup_write_mystruct(write_only pipe +mystruct p, __global mystruct *src) { + int gid = get_global_id(0); + reserve_id_t res_id = work_group_reserve_write_pipe(p, get_local_size(0)); + if(is_valid_reserve_id(res_id)) + { + write_pipe(p, res_id, get_local_id(0), &src[gid]); + work_group_commit_write_pipe(p, res_id); + } +} +__kernel void compiler_pipe_workgroup_read_mystruct(read_only pipe +mystruct p, __global mystruct *dst) { + int gid = get_global_id(0); + reserve_id_t res_id = work_group_reserve_read_pipe(p, get_local_size(0)); + if(is_valid_reserve_id(res_id)) + { + read_pipe(p, res_id, get_local_id(0), &dst[gid]); + work_group_commit_read_pipe(p, res_id); + } +} + +__kernel void compiler_pipe_query(write_only pipe int p, __global uint +*src) { + int gid = get_global_id(0); + write_pipe(p,&gid); + if(gid == 0) { + src[0] = get_pipe_num_packets(p); + src[1] = get_pipe_max_packets(p); + } +} diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index d1443aa..a6d6df0 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -222,7 +222,9 @@ set (utests_sources runtime_use_host_ptr_image.cpp compiler_get_sub_group_size.cpp compiler_get_sub_group_id.cpp - compiler_sub_group_shuffle.cpp) + compiler_sub_group_shuffle.cpp + runtime_pipe_query.cpp + compiler_pipe_builtin.cpp) if (LLVM_VERSION_NODOT VERSION_GREATER 34) SET(utests_sources diff --git a/utests/compiler_pipe_builtin.cpp b/utests/compiler_pipe_builtin.cpp new file mode 100644 index 0000000..3c6ae9b --- /dev/null +++ b/utests/compiler_pipe_builtin.cpp @@ -0,0 +1,70 @@ +#include <string.h> +#include "utest_helper.hpp" +typedef struct{ + int a; + uint b; +}mystruct; + +#define PIPE_BUILTIN(TYPE,GROUP) \ +static void compiler_pipe_##GROUP##_##TYPE(void) \ { \ + const size_t w = 16; \ + uint32_t ans_host = 0; \ + uint32_t ans_device = 0; \ + /* pipe write kernel*/ \ + OCL_CREATE_KERNEL_FROM_FILE("compiler_pipe_builtin", +"compiler_pipe_"#GROUP"_write_"#TYPE); \ + OCL_CALL2(clCreatePipe, buf[0], ctx, 0, sizeof(TYPE), w, NULL);\ + OCL_CREATE_BUFFER(buf[1], CL_MEM_READ_WRITE, w * sizeof(TYPE), +NULL);\ + OCL_MAP_BUFFER(1);\ + for (uint32_t i = 0; i < w; i++)\ + ((uint32_t*)buf_data[1])[i] = i;\ + OCL_UNMAP_BUFFER(1);\ + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);\ + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);\ + globals[0] = w;\ + locals[0] = 16;\ + OCL_NDRANGE(1);\ + OCL_CALL(clReleaseKernel, kernel);\ + /* pipe read kernel */\ + OCL_CREATE_KERNEL_FROM_FILE("compiler_pipe_builtin", +"compiler_pipe_"#GROUP"_read_"#TYPE);\ + OCL_CREATE_BUFFER(buf[2], CL_MEM_READ_WRITE, w * sizeof(TYPE), +NULL);\ + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);\ + OCL_SET_ARG(1, sizeof(cl_mem), &buf[2]);\ + OCL_NDRANGE(1);\ + /* Check result */\ + OCL_MAP_BUFFER(2);\ + for (uint32_t i = 0; i < w; i++) {\ + ans_device += ((uint32_t*)buf_data[2])[i];\ + ans_host += i;\ + }\ + OCL_UNMAP_BUFFER(2);\ + OCL_ASSERT(ans_host == ans_device);\ +}\ +MAKE_UTEST_FROM_FUNCTION(compiler_pipe_##GROUP##_##TYPE); + +PIPE_BUILTIN(int, convenience) +PIPE_BUILTIN(mystruct, convenience) +PIPE_BUILTIN(int, reserve) +PIPE_BUILTIN(mystruct, reserve) +PIPE_BUILTIN(int, workgroup) +PIPE_BUILTIN(mystruct, workgroup) + +static void compiler_pipe_query(void) { + const size_t w = 32; + const size_t sz = 16; + /* pipe write kernel */ + OCL_CREATE_KERNEL_FROM_FILE("compiler_pipe_builtin", +"compiler_pipe_query"); + OCL_CALL2(clCreatePipe, buf[0], ctx, 0, sizeof(uint32_t), w, NULL); + OCL_CREATE_BUFFER(buf[1], CL_MEM_READ_WRITE, sz * sizeof(uint32_t), +NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = sz; + locals[0] = 16; + OCL_NDRANGE(1); + /*Check result */ + OCL_MAP_BUFFER(1); + OCL_ASSERT(sz == ((uint32_t *)buf_data[1])[0] && w == ((uint32_t +*)buf_data[1])[1]); + OCL_UNMAP_BUFFER(2); +} +MAKE_UTEST_FROM_FUNCTION(compiler_pipe_query); + diff --git a/utests/runtime_pipe_query.cpp b/utests/runtime_pipe_query.cpp new file mode 100644 index 0000000..73330b3 --- /dev/null +++ b/utests/runtime_pipe_query.cpp @@ -0,0 +1,16 @@ +#include <string.h> +#include "utest_helper.hpp" +static void runtime_pipe_query(void) { + const size_t w = 16; + const size_t sz = 8; + cl_uint retnum, retsz; + /* pipe write kernel */ + OCL_CALL2(clCreatePipe, buf[0], ctx, 0, sz, w, NULL); + OCL_CALL(clGetPipeInfo, buf[0], CL_PIPE_MAX_PACKETS, sizeof(retnum), +&retnum, NULL); + OCL_CALL(clGetPipeInfo, buf[0], CL_PIPE_PACKET_SIZE, sizeof(retsz), +&retsz, NULL); + + /*Check result */ + OCL_ASSERT(sz == retsz && w == retnum); } +MAKE_UTEST_FROM_FUNCTION(runtime_pipe_query); + -- 2.5.0 _______________________________________________ Beignet mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/beignet
