LGTM -----Original Message----- From: Beignet [mailto:[email protected]] On Behalf Of [email protected] Sent: Tuesday, June 10, 2014 5:07 AM To: [email protected] Cc: Luo, Xionghu Subject: [Beignet] [PATCH v5 3/3] add [opencl-1.2] test case runtime_compile_link.
From: Luo <[email protected]> Signed-off-by: Luo <[email protected]> --- kernels/include/runtime_compile_link_inc.h | 4 + kernels/runtime_compile_link.h | 1 + kernels/runtime_compile_link_a.cl | 13 +++ kernels/runtime_compile_link_b.cl | 9 ++ utests/CMakeLists.txt | 1 + utests/runtime_compile_link.cpp | 127 +++++++++++++++++++++++++++++ 6 files changed, 155 insertions(+) create mode 100644 kernels/include/runtime_compile_link_inc.h create mode 100644 kernels/runtime_compile_link.h create mode 100644 kernels/runtime_compile_link_a.cl create mode 100644 kernels/runtime_compile_link_b.cl create mode 100644 utests/runtime_compile_link.cpp diff --git a/kernels/include/runtime_compile_link_inc.h b/kernels/include/runtime_compile_link_inc.h new file mode 100644 index 0000000..4011278 --- /dev/null +++ b/kernels/include/runtime_compile_link_inc.h @@ -0,0 +1,4 @@ +inline int greater(long x, long y) +{ + return x > y ; +} diff --git a/kernels/runtime_compile_link.h b/kernels/runtime_compile_link.h new file mode 100644 index 0000000..ae2c56e --- /dev/null +++ b/kernels/runtime_compile_link.h @@ -0,0 +1 @@ +int comp_long(long x, long y); diff --git a/kernels/runtime_compile_link_a.cl b/kernels/runtime_compile_link_a.cl new file mode 100644 index 0000000..b17861f --- /dev/null +++ b/kernels/runtime_compile_link_a.cl @@ -0,0 +1,13 @@ +#include "runtime_compile_link.h" +#include "include/runtime_compile_link_inc.h" + +int comp_long(long x, long y) +{ + return x < y ; +} + +kernel void runtime_compile_link_a(global long *src1, global long +*src2, global long *dst) { + int i = get_global_id(0); + int j = comp_long(src1[i], src2[i]); + dst[i] = j ? 3 : 4; +} diff --git a/kernels/runtime_compile_link_b.cl b/kernels/runtime_compile_link_b.cl new file mode 100644 index 0000000..89b5a2d --- /dev/null +++ b/kernels/runtime_compile_link_b.cl @@ -0,0 +1,9 @@ +#include "runtime_compile_link.h" +#include "include/runtime_compile_link_inc.h" + +kernel void runtime_compile_link_b(global long *src1, global long +*src2, global long *dst) { + int i = get_global_id(0); + int j = comp_long(src1[i], src2[i]); + dst[i] = j ? 3 : 4; + int k = greater(src1[i], src2[i]); +} diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index 698c9ff..bee3e8f 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -157,6 +157,7 @@ set (utests_sources runtime_event.cpp runtime_barrier_list.cpp runtime_marker_list.cpp + runtime_compile_link.cpp compiler_double.cpp compiler_double_2.cpp compiler_double_3.cpp diff --git a/utests/runtime_compile_link.cpp b/utests/runtime_compile_link.cpp new file mode 100644 index 0000000..8aeea31 --- /dev/null +++ b/utests/runtime_compile_link.cpp @@ -0,0 +1,127 @@ +#include <cstdint> +#include <cstring> +#include <iostream> +#include "utest_helper.hpp" +#include "utest_file_map.hpp" + +#define BUFFERSIZE 32*1024 + +int init_program(const char* name, cl_context ctx, cl_program *pg ) { + cl_int err; + char* ker_path = cl_do_kiss_path(name, device); + + cl_file_map_t *fm = cl_file_map_new(); err = cl_file_map_open(fm, + ker_path); if(err != CL_FILE_MAP_SUCCESS) + OCL_ASSERT(0); + const char *src = cl_file_map_begin(fm); + + *pg = clCreateProgramWithSource(ctx, 1, &src, NULL, &err); + free(ker_path); cl_file_map_delete(fm); return 0; + +} + +void runtime_compile_link(void) +{ + + cl_int err; + + const char* header_file_name="runtime_compile_link.h"; + cl_program foo_pg; + init_program(header_file_name, ctx, &foo_pg); + + const char* myinc_file_name="include/runtime_compile_link_inc.h"; + cl_program myinc_pg; + init_program(myinc_file_name, ctx, &myinc_pg); + + const char* file_name_A="runtime_compile_link_a.cl"; + cl_program program_A; + init_program(file_name_A, ctx, &program_A); + + cl_program input_headers[2] = { foo_pg, myinc_pg}; + const char * input_header_names[2] = {header_file_name, myinc_file_name}; + + err = clCompileProgram(program_A, + 0, NULL, // num_devices & device_list + NULL, // compile_options + 2, // num_input_headers + input_headers, + input_header_names, + NULL, NULL); + + OCL_ASSERT(err==CL_SUCCESS); + const char* file_name_B="runtime_compile_link_b.cl"; + cl_program program_B; + init_program(file_name_B, ctx, &program_B); + + err = clCompileProgram(program_B, + 0, NULL, // num_devices & device_list + NULL, // compile_options + 2, // num_input_headers + input_headers, + input_header_names, + NULL, NULL); + + OCL_ASSERT(err==CL_SUCCESS); + cl_program input_programs[2] = { program_A, program_B}; + cl_program linked_program = clLinkProgram(ctx, 0, NULL, NULL, 2, input_programs, NULL, NULL, &err); + + + OCL_ASSERT(linked_program != NULL); + OCL_ASSERT(err == CL_SUCCESS); + + // link success, run this kernel. + + const size_t n = 16; + int64_t src1[n], src2[n]; + + src1[0] = (int64_t)1 << 63, src2[0] = 0x7FFFFFFFFFFFFFFFll; + src1[1] = (int64_t)1 << 63, src2[1] = ((int64_t)1 << 63) | 1; + src1[2] = -1ll, src2[2] = 0; + src1[3] = ((int64_t)123 << 32) | 0x7FFFFFFF, src2[3] = ((int64_t)123 << 32) | 0x80000000; + src1[4] = 0x7FFFFFFFFFFFFFFFll, src2[4] = (int64_t)1 << 63; + src1[5] = ((int64_t)1 << 63) | 1, src2[5] = (int64_t)1 << 63; + src1[6] = 0, src2[6] = -1ll; + src1[7] = ((int64_t)123 << 32) | 0x80000000, src2[7] = ((int64_t)123 << 32) | 0x7FFFFFFF; + for(size_t i=8; i<n; i++) { + src1[i] = i; + src2[i] = i; + } + + globals[0] = n; + locals[0] = 16; + + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(int64_t), NULL); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(int64_t), NULL); + OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(int64_t), NULL); + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + memcpy(buf_data[0], src1, sizeof(src1)); + memcpy(buf_data[1], src2, sizeof(src2)); + OCL_UNMAP_BUFFER(0); + OCL_UNMAP_BUFFER(1); + + kernel = clCreateKernel(linked_program, "runtime_compile_link_a", &err); + + OCL_ASSERT(err == CL_SUCCESS); + + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]); + + clEnqueueNDRangeKernel(queue, kernel, 1, NULL, globals, locals, 0, NULL, NULL); + + OCL_MAP_BUFFER(2); + for (int32_t i = 0; i < (int32_t) n; ++i) { + int64_t *dest = (int64_t *)buf_data[2]; + int64_t x = (src1[i] < src2[i]) ? 3 : 4; + OCL_ASSERT(x == dest[i]); + } + OCL_UNMAP_BUFFER(2); + OCL_DESTROY_KERNEL_KEEP_PROGRAM(true); +} + +MAKE_UTEST_FROM_FUNCTION(runtime_compile_link); -- 1.8.1.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
