Added tests for shorts and bytes operations

diff --git a/kernels/compiler_sub_bytes.cl b/kernels/compiler_sub_bytes.cl
new file mode 100644 (file)
index 0000000..f058561
--- /dev/null
+++ b/kernels/compiler_sub_bytes.cl
@@ -0,0 +1,7 @@
+__kernel void
+compiler_sub_bytes(__global char *src0, __global char *src1, __global char *dst)
+{
+  int id = (int)get_global_id(0);
+  dst[id] = src0[id] - src1[id];
+}
+

diff --git a/kernels/compiler_sub_shorts.cl b/kernels/compiler_sub_shorts.cl
new file mode 100644 (file)
index 0000000..d26de7f
--- /dev/null
+++ b/kernels/compiler_sub_shorts.cl
@@ -0,0 +1,7 @@
+__kernel void
+compiler_sub_shorts(__global short *src0, __global short *src1, __global short *dst)
+{
+  int id = (int)get_global_id(0);
+  dst[id] = src0[id] - src1[id];
+}
+

diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index a6e9e5a..00818e1 100644 (file)
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -9,9 +9,11 @@ ADD_LIBRARY(utests SHARED
             utest.cpp
             app_mandelbrot.cpp
             compiler_write_only.cpp
+            compiler_sub_bytes.cpp
+            compiler_sub_shorts.cpp
             compiler_copy_buffer.cpp
             compiler_copy_buffer_row.cpp
-compiler_byte_scatter.cpp
+            compiler_byte_scatter.cpp
             compiler_short_scatter.cpp
             compiler_if_else.cpp
             compiler_unstructured_branch0.cpp

diff --git a/utests/compiler_sub_bytes.cpp b/utests/compiler_sub_bytes.cpp
new file mode 100644 (file)
index 0000000..f2262f1
--- /dev/null
+++ b/utests/compiler_sub_bytes.cpp
@@ -0,0 +1,54 @@
+/* 
+ * Copyright © 2012 Intel Corporation
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library. If not, see <http://www.gnu.org/licenses/>.
+ *
+ * Author: Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+
+#include "utest_helper.hpp"
+
+static void compiler_sub_bytes(void)
+{
+  const size_t n = 16;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_sub_bytes");
+  buf_data[0] = (int8_t*) malloc(sizeof(int8_t) * n);
+  buf_data[1] = (int8_t*) malloc(sizeof(int8_t) * n);
+  for (uint32_t i = 0; i < n; ++i) ((int8_t*)buf_data[0])[i] = (int8_t) rand();
+  for (uint32_t i = 0; i < n; ++i) ((int8_t*)buf_data[1])[i] = (int8_t) rand();
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(int8_t), buf_data[0]);
+  OCL_CREATE_BUFFER(buf[1], CL_MEM_COPY_HOST_PTR, n * sizeof(int8_t), buf_data[0]);
+  OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(int8_t), NULL);
+
+  // Run the kernel
+  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]);
+  globals[0] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  // Check result
+  OCL_MAP_BUFFER(2);
+  for (uint32_t i = 0; i < n; ++i)
+    OCL_ASSERT(((int8_t*)buf_data[2])[i] = ((int8_t*)buf_data[0])[i] - ((int8_t*)buf_data[1])[i]);
+  free(buf_data[0]);
+  free(buf_data[1]);
+  buf_data[0] = buf_data[1] = NULL;
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_sub_bytes);
+

diff --git a/utests/compiler_sub_shorts.cpp b/utests/compiler_sub_shorts.cpp
new file mode 100644 (file)
index 0000000..bf4271f
--- /dev/null
+++ b/utests/compiler_sub_shorts.cpp
@@ -0,0 +1,55 @@
+/* 
+ * Copyright © 2012 Intel Corporation
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library. If not, see <http://www.gnu.org/licenses/>.
+ *
+ * Author: Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+
+#include "utest_helper.hpp"
+
+static void compiler_sub_shorts(void)
+{
+  const size_t n = 16;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_sub_shorts");
+  buf_data[0] = (int16_t*) malloc(sizeof(int16_t) * n);
+  buf_data[1] = (int16_t*) malloc(sizeof(int16_t) * n);
+  for (uint32_t i = 0; i < n; ++i) ((int16_t*)buf_data[0])[i] = (int16_t) rand();
+  for (uint32_t i = 0; i < n; ++i) ((int16_t*)buf_data[1])[i] = (int16_t) rand();
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(int16_t), buf_data[0]);
+  OCL_CREATE_BUFFER(buf[1], CL_MEM_COPY_HOST_PTR, n * sizeof(int16_t), buf_data[0]);
+  OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(int16_t), NULL);
+
+  // Run the kernel
+  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]);
+  globals[0] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  // Check result
+  OCL_MAP_BUFFER(2);
+  for (uint32_t i = 0; i < n; ++i)
+    OCL_ASSERT(((int16_t*)buf_data[2])[i] = ((int16_t*)buf_data[0])[i] - ((int16_t*)buf_data[1])[i]);
+  free(buf_data[0]);
+  free(buf_data[1]);
+  buf_data[0] = buf_data[1] = NULL;
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_sub_shorts);
+
+