Could you change the file/folder names a little meaningful? Thanks! -----Original Message----- From: Beignet [mailto:[email protected]] On Behalf Of [email protected] Sent: Friday, June 06, 2014 2:19 PM To: [email protected] Cc: Luo, Xionghu Subject: [Beignet] [PATCH v4 3/3] add [opencl-1.2] test case runtime_compile_link.
From: Luo <[email protected]> Signed-off-by: Luo <[email protected]> --- kernels/multi2.h | 1 + kernels/multi_A.cl | 13 ++++ kernels/multi_B.cl | 9 +++ kernels/mydir/multi3.h | 4 ++ utests/CMakeLists.txt | 1 + utests/runtime_compile_link.cpp | 127 ++++++++++++++++++++++++++++++++++++++++ 6 files changed, 155 insertions(+) create mode 100644 kernels/multi2.h create mode 100644 kernels/multi_A.cl create mode 100644 kernels/multi_B.cl create mode 100644 kernels/mydir/multi3.h create mode 100644 utests/runtime_compile_link.cpp diff --git a/kernels/multi2.h b/kernels/multi2.h new file mode 100644 index 0000000..ae2c56e --- /dev/null +++ b/kernels/multi2.h @@ -0,0 +1 @@ +int comp_long(long x, long y); diff --git a/kernels/multi_A.cl b/kernels/multi_A.cl new file mode 100644 index 0000000..9282b8d --- /dev/null +++ b/kernels/multi_A.cl @@ -0,0 +1,13 @@ +#include "multi2.h" +#include "mydir/multi3.h" + +int comp_long(long x, long y) +{ + return x < y ; +} + +kernel void multi_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/multi_B.cl b/kernels/multi_B.cl new file mode 100644 index 0000000..de147eb --- /dev/null +++ b/kernels/multi_B.cl @@ -0,0 +1,9 @@ +#include "multi2.h" +#include "mydir/multi3.h" + +kernel void multi_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/kernels/mydir/multi3.h b/kernels/mydir/multi3.h new file mode 100644 index 0000000..4011278 --- /dev/null +++ b/kernels/mydir/multi3.h @@ -0,0 +1,4 @@ +inline int greater(long x, long y) +{ + return x > y ; +} 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..17fe413 --- /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="multi2.h"; + cl_program foo_pg; + init_program(header_file_name, ctx, &foo_pg); + + const char* myinc_file_name="mydir/multi3.h"; + cl_program myinc_pg; + init_program(myinc_file_name, ctx, &myinc_pg); + + const char* file_name_A="multi_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] = { "multi2.h", "mydir/multi3.h"}; + + 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="multi_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, "multi_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
