[Beignet] [PATCH v5 3/3] add [opencl-1.2] test case runtime_compile_link.
xionghu.luo at intel.com
xionghu.luo at intel.com
Mon Jun 9 14:07:18 PDT 2014
From: Luo <xionghu.luo at intel.com>
Signed-off-by: Luo <xionghu.luo at intel.com>
---
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
More information about the Beignet
mailing list