[Beignet] [PATCH] Add a load bool imm test case.

Yang Rong rong.r.yang at intel.com
Tue Aug 13 02:10:07 PDT 2013


Signed-off-by: Yang Rong <rong.r.yang at intel.com>
---
 kernels/compiler_load_bool_imm.cl |   12 ++++++++++++
 utests/CMakeLists.txt             |    1 +
 utests/compiler_load_bool_imm.cpp |   28 ++++++++++++++++++++++++++++
 3 files changed, 41 insertions(+)
 create mode 100644 kernels/compiler_load_bool_imm.cl
 create mode 100644 utests/compiler_load_bool_imm.cpp

diff --git a/kernels/compiler_load_bool_imm.cl b/kernels/compiler_load_bool_imm.cl
new file mode 100644
index 0000000..fda49b9
--- /dev/null
+++ b/kernels/compiler_load_bool_imm.cl
@@ -0,0 +1,12 @@
+__kernel void
+compiler_load_bool_imm(__global int *dst, __local int *localBuffer, int copiesPerWorkItem )
+{
+  int i;
+  for(i=0; i<copiesPerWorkItem; i++)
+    localBuffer[get_local_id(0)*copiesPerWorkItem+i] = copiesPerWorkItem;
+  barrier(CLK_LOCAL_MEM_FENCE);
+
+  for(i=0; i<copiesPerWorkItem; i++)
+    dst[get_global_id(0)*copiesPerWorkItem + i] = localBuffer[get_local_id(0)*copiesPerWorkItem+i];
+  barrier(CLK_LOCAL_MEM_FENCE);
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 3922220..4093463 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -87,6 +87,7 @@ set (utests_sources
   compiler_insn_selection_min.cpp
   compiler_insn_selection_max.cpp
   compiler_insn_selection_masked_min_max.cpp
+	compiler_load_bool_imm.cpp
   compiler_global_memory_barrier.cpp
   compiler_local_memory_two_ptr.cpp
   compiler_local_memory_barrier.cpp
diff --git a/utests/compiler_load_bool_imm.cpp b/utests/compiler_load_bool_imm.cpp
new file mode 100644
index 0000000..47ae9ad
--- /dev/null
+++ b/utests/compiler_load_bool_imm.cpp
@@ -0,0 +1,28 @@
+#include "utest_helper.hpp"
+
+static void compiler_load_bool_imm(void)
+{
+  const size_t n = 1024;
+  const size_t local_size = 16;
+  const int copiesPerWorkItem = 5;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_load_bool_imm");
+  OCL_CREATE_BUFFER(buf[0], 0, n * copiesPerWorkItem * sizeof(uint32_t), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, local_size*copiesPerWorkItem*sizeof(int), NULL); // 16 x int
+  OCL_SET_ARG(2, sizeof(int), &copiesPerWorkItem); // 16 x int
+
+  // Run the kernel
+  globals[0] = n;
+  locals[0] = local_size;
+  OCL_NDRANGE(1);
+  OCL_MAP_BUFFER(0);
+
+  // Check results
+  int *dst = (int*)buf_data[0];
+  for (uint32_t i = 0; i < n * copiesPerWorkItem; i++)
+    OCL_ASSERT(dst[i] == copiesPerWorkItem);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_load_bool_imm);
-- 
1.7.10.4



More information about the Beignet mailing list