[Beignet] [PATCH 4/4] add [opencl-1.2] test case runtime_cl.

[xionghu.luo at intel.com]

From: Luo <xionghu.luo at intel.com>

---
 utests/CMakeLists.txt           |   1 +
 utests/runtime_compile_link.cpp | 127 ++++++++++++++++++++++++++++++++++++++++
 2 files changed, 128 insertions(+)
 create mode 100644 utests/runtime_compile_link.cpp

diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 5f0649f..c6d4098 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -153,6 +153,7 @@ set (utests_sources
   runtime_createcontext.cpp
   runtime_null_kernel_arg.cpp
   runtime_event.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..df55ab8
--- /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_cl(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_cl);
-- 
1.8.1.2