add test for cl buffer created with CL_MEM_USE_HOST_PTRuserptr

Signed-off-by: Guo Yejun <yejun.guo@intel.com>

3 files changed, 48 insertions, 0 deletions

diff --git a/kernels/runtime_use_host_ptr_buffer.cl b/kernels/runtime_use_host_ptr_buffer.cl
new file mode 100644
index 00000000..dbaadf87
--- /dev/null
+++ b/kernels/runtime_use_host_ptr_buffer.cl
@@ -0,0 +1,6 @@
+__kernel void
+runtime_use_host_ptr_buffer(__global int* buf)
+{
+  int id = (int)get_global_id(0);
+  buf[id] = buf[id] / 2;
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index b0225466..f609cc2d 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -195,6 +195,7 @@ set (utests_sources
   compiler_constant_expr.cpp
   compiler_assignment_operation_in_if.cpp
   vload_bench.cpp
+  runtime_use_host_ptr_buffer.cpp
   utest_assert.cpp
   utest.cpp
   utest_file_map.cpp
@@ -214,6 +215,11 @@ else(GEN_PCI_ID)
   DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/../backend/src/gbe_bin_generater ${kernel_bin}.cl)
 endif(GEN_PCI_ID)
 
+if (DRM_INTEL_USERPTR)
+SET(CMAKE_CXX_FLAGS "-DHAS_USERPTR ${CMAKE_CXX_FLAGS}")
+SET(CMAKE_C_FLAGS "-DHAS_USERPTR ${CMAKE_C_FLAGS}")
+endif (DRM_INTEL_USERPTR)
+
 ADD_CUSTOM_TARGET(kernel_bin.bin
     DEPENDS ${kernel_bin}.bin)
 
diff --git a/utests/runtime_use_host_ptr_buffer.cpp b/utests/runtime_use_host_ptr_buffer.cpp
new file mode 100644
index 00000000..ca06f4be
--- /dev/null
+++ b/utests/runtime_use_host_ptr_buffer.cpp
@@ -0,0 +1,36 @@
+#include "utest_helper.hpp"
+
+static void runtime_use_host_ptr_buffer(void)
+{
+  const size_t n = 4096*100;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("runtime_use_host_ptr_buffer");
+  buf_data[0] = (uint32_t*) aligned_alloc(4096, sizeof(uint32_t) * n);
+  for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = i;
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_USE_HOST_PTR, n * sizeof(uint32_t), buf_data[0]);
+
+  // Run the kernel
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  globals[0] = n;
+  locals[0] = 256;
+  OCL_NDRANGE(1);
+
+  // Check result
+
+#ifdef HAS_USERPTR
+  OCL_FINISH();
+#else
+  void* mapptr = (int*)clEnqueueMapBuffer(queue, buf[0], CL_TRUE, CL_MAP_READ, 0, n*sizeof(uint32_t), 0, NULL, NULL, NULL);
+  OCL_ASSERT(mapptr == buf_data[0]);
+  clEnqueueUnmapMemObject(queue, buf[0], mapptr, 0, NULL, NULL);
+#endif
+
+  for (uint32_t i = 0; i < n; ++i)
+    OCL_ASSERT(((uint32_t*)buf_data[0])[i] == i / 2);
+
+  free(buf_data[0]);
+  buf_data[0] = NULL;
+}
+
+MAKE_UTEST_FROM_FUNCTION(runtime_use_host_ptr_buffer);