Add test case for long bitcast.

Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

3 files changed, 275 insertions, 0 deletions

diff --git a/kernels/compiler_long_bitcast.cl b/kernels/compiler_long_bitcast.cl
new file mode 100644
index 00000000..4b008ef1
--- /dev/null
+++ b/kernels/compiler_long_bitcast.cl
@@ -0,0 +1,47 @@
+__kernel void compiler_bitcast_char8_to_long(__global char8 *src, __global ulong *dst)
+{
+  int tid = get_global_id(0);
+  char8 v = src[tid];
+  ulong dl = as_ulong(v);
+  dst[tid] = dl;
+}
+
+__kernel void compiler_bitcast_long_to_char8(__global ulong *src,  __global uchar8 *dst)
+{
+  int tid = get_global_id(0);
+  ulong v = src[tid];
+  uchar8 dl = as_uchar8(v);
+  dst[tid] = dl;
+}
+
+__kernel void compiler_bitcast_int2_to_long(__global int2 *src, __global ulong *dst)
+{
+  int tid = get_global_id(0);
+  int2 v = src[tid];
+  ulong dl = as_ulong(v);
+  dst[tid] = dl;
+}
+
+__kernel void compiler_bitcast_long_to_int2(__global ulong *src,  __global uint2 *dst)
+{
+  int tid = get_global_id(0);
+  ulong v = src[tid];
+  uint2 dl = as_uint2(v);
+  dst[tid] = dl;
+}
+
+__kernel void compiler_bitcast_short4_to_long(__global short4 *src, __global ulong *dst)
+{
+  int tid = get_global_id(0);
+  short4 v = src[tid];
+  ulong dl = as_ulong(v);
+  dst[tid] = dl;
+}
+
+__kernel void compiler_bitcast_long_to_short4(__global ulong *src,  __global ushort4 *dst)
+{
+  int tid = get_global_id(0);
+  ulong v = src[tid];
+  ushort4 dl = as_ushort4(v);
+  dst[tid] = dl;
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 0f56c6c9..f8fb9c6b 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -182,6 +182,7 @@ set (utests_sources
   compiler_long_asr.cpp
   compiler_long_mult.cpp
   compiler_long_cmp.cpp
+  compiler_long_bitcast.cpp
   compiler_function_argument3.cpp
   compiler_function_qualifiers.cpp
   compiler_bool_cross_basic_block.cpp
diff --git a/utests/compiler_long_bitcast.cpp b/utests/compiler_long_bitcast.cpp
new file mode 100644
index 00000000..5bd962de
--- /dev/null
+++ b/utests/compiler_long_bitcast.cpp
@@ -0,0 +1,227 @@
+#include <cstdint>
+#include <cstring>
+#include <iostream>
+#include "utest_helper.hpp"
+
+void compiler_bitcast_char8_to_long(void)
+{
+  const size_t n = 64;
+  const int v = 8;
+  char src[n * v];
+  uint64_t *dst = (uint64_t *)src;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_long_bitcast", "compiler_bitcast_char8_to_long");
+  OCL_CREATE_BUFFER(buf[0], 0, sizeof(src), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, sizeof(src), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  globals[0] = n;
+  locals[0] = 16;
+
+  for (int32_t i = 0; i < (int32_t) n*v; ++i) {
+    src[i] = (char)rand();
+  }
+
+  OCL_MAP_BUFFER(0);
+  memcpy(buf_data[0], src, sizeof(src));
+  OCL_UNMAP_BUFFER(0);
+
+  // Run the kernel on GPU
+  OCL_NDRANGE(1);
+
+  // Compare
+  OCL_MAP_BUFFER(1);
+  for (int32_t i = 0; i < (int32_t) n; ++i) {
+    OCL_ASSERT(((uint64_t *)(buf_data[1]))[i] == dst[i]);
+    //printf("ref is 0x%lx, result is 0x%lx\n", dst[i], ((int64_t *)(buf_data[1]))[i]);
+  }
+  OCL_UNMAP_BUFFER(1);
+}
+
+void compiler_bitcast_long_to_char8(void)
+{
+  const size_t n = 64;
+  const int v = 8;
+  ulong src[n];
+  char *dst = (char *)src;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_long_bitcast", "compiler_bitcast_char8_to_long");
+  OCL_CREATE_BUFFER(buf[0], 0, sizeof(src), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, sizeof(src), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  globals[0] = n;
+  locals[0] = 16;
+
+  for (int32_t i = 0; i < (int32_t) n; ++i) {
+    src[i] = ((int64_t)rand() << 32) + rand();
+  }
+
+  OCL_MAP_BUFFER(0);
+  memcpy(buf_data[0], src, sizeof(src));
+  OCL_UNMAP_BUFFER(0);
+
+  // Run the kernel on GPU
+  OCL_NDRANGE(1);
+
+  // Compare
+  OCL_MAP_BUFFER(1);
+  for (int32_t i = 0; i < (int32_t) n*v; ++i) {
+    OCL_ASSERT(((char *)(buf_data[1]))[i] == dst[i]);
+//    printf("ref is 0x%2x, result is 0x%2x\n", dst[i], ((char *)(buf_data[1]))[i]);
+  }
+  OCL_UNMAP_BUFFER(1);
+}
+
+void compiler_bitcast_int2_to_long(void)
+{
+  const size_t n = 64;
+  const int v = 2;
+  int src[n * v];
+  uint64_t *dst = (uint64_t *)src;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_long_bitcast", "compiler_bitcast_int2_to_long");
+  OCL_CREATE_BUFFER(buf[0], 0, sizeof(src), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, sizeof(src), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  globals[0] = n;
+  locals[0] = 16;
+
+  for (int32_t i = 0; i < (int32_t) n*v; ++i) {
+    src[i] = (char)rand();
+  }
+
+  OCL_MAP_BUFFER(0);
+  memcpy(buf_data[0], src, sizeof(src));
+  OCL_UNMAP_BUFFER(0);
+
+  // Run the kernel on GPU
+  OCL_NDRANGE(1);
+
+  // Compare
+  OCL_MAP_BUFFER(1);
+  for (int32_t i = 0; i < (int32_t) n; ++i) {
+    OCL_ASSERT(((uint64_t *)(buf_data[1]))[i] == dst[i]);
+    //printf("ref is 0x%lx, result is 0x%lx\n", dst[i], ((int64_t *)(buf_data[1]))[i]);
+  }
+  OCL_UNMAP_BUFFER(1);
+}
+
+void compiler_bitcast_long_to_int2(void)
+{
+  const size_t n = 64;
+  const int v = 2;
+  ulong src[n];
+  uint32_t *dst = (uint32_t *)src;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_long_bitcast", "compiler_bitcast_long_to_int2");
+  OCL_CREATE_BUFFER(buf[0], 0, sizeof(src), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, sizeof(src), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  globals[0] = n;
+  locals[0] = 16;
+
+  for (int32_t i = 0; i < (int32_t) n; ++i) {
+    src[i] = ((int64_t)i << 32) + i;
+  }
+
+  OCL_MAP_BUFFER(0);
+  memcpy(buf_data[0], src, sizeof(src));
+  OCL_UNMAP_BUFFER(0);
+
+  // Run the kernel on GPU
+  OCL_NDRANGE(1);
+
+  // Compare
+  OCL_MAP_BUFFER(1);
+  for (int32_t i = 0; i < (int32_t) n*v; ++i) {
+    OCL_ASSERT(((uint32_t *)(buf_data[1]))[i] == dst[i]);
+    //printf("ref is 0x%2x, result is 0x%2x\n", dst[i], ((uint32_t *)(buf_data[1]))[i]);
+  }
+  OCL_UNMAP_BUFFER(1);
+}
+
+void compiler_bitcast_short4_to_long(void)
+{
+  const size_t n = 64;
+  const int v = 4;
+  short src[n * v];
+  uint64_t *dst = (uint64_t *)src;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_long_bitcast", "compiler_bitcast_short4_to_long");
+  OCL_CREATE_BUFFER(buf[0], 0, sizeof(src), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, sizeof(src), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  globals[0] = n;
+  locals[0] = 16;
+
+  for (int32_t i = 0; i < (int32_t) n*v; ++i) {
+    src[i] = (char)rand();
+  }
+
+  OCL_MAP_BUFFER(0);
+  memcpy(buf_data[0], src, sizeof(src));
+  OCL_UNMAP_BUFFER(0);
+
+  // Run the kernel on GPU
+  OCL_NDRANGE(1);
+
+  // Compare
+  OCL_MAP_BUFFER(1);
+  for (int32_t i = 0; i < (int32_t) n; ++i) {
+    OCL_ASSERT(((uint64_t *)(buf_data[1]))[i] == dst[i]);
+    //printf("ref is 0x%lx, result is 0x%lx\n", dst[i], ((int64_t *)(buf_data[1]))[i]);
+  }
+  OCL_UNMAP_BUFFER(1);
+}
+
+void compiler_bitcast_long_to_short4(void)
+{
+  const size_t n = 64;
+  const int v = 4;
+  ulong src[n];
+  uint16_t *dst = (uint16_t *)src;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_long_bitcast", "compiler_bitcast_long_to_short4");
+  OCL_CREATE_BUFFER(buf[0], 0, sizeof(src), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, sizeof(src), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  globals[0] = n;
+  locals[0] = 16;
+
+  for (int32_t i = 0; i < (int32_t) n; ++i) {
+    src[i] = ((int64_t)rand() << 32) + rand();
+  }
+
+  OCL_MAP_BUFFER(0);
+  memcpy(buf_data[0], src, sizeof(src));
+  OCL_UNMAP_BUFFER(0);
+
+  // Run the kernel on GPU
+  OCL_NDRANGE(1);
+
+  // Compare
+  OCL_MAP_BUFFER(1);
+  for (int32_t i = 0; i < (int32_t) n*v; ++i) {
+    OCL_ASSERT(((uint16_t *)(buf_data[1]))[i] == dst[i]);
+    //printf("ref is 0x%2x, result is 0x%2x\n", dst[i], ((uint16_t *)(buf_data[1]))[i]);
+  }
+  OCL_UNMAP_BUFFER(1);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_bitcast_char8_to_long);
+MAKE_UTEST_FROM_FUNCTION(compiler_bitcast_long_to_char8);
+MAKE_UTEST_FROM_FUNCTION(compiler_bitcast_int2_to_long);
+MAKE_UTEST_FROM_FUNCTION(compiler_bitcast_long_to_int2);
+MAKE_UTEST_FROM_FUNCTION(compiler_bitcast_short4_to_long);
+MAKE_UTEST_FROM_FUNCTION(compiler_bitcast_long_to_short4);