[Beignet] [PATCH 31/31 V2] Add test case for long bitcast.
junyan.he at inbox.com
junyan.he at inbox.com
Sun Jan 18 23:35:04 PST 2015
From: Junyan He <junyan.he at linux.intel.com>
Signed-off-by: Junyan He <junyan.he at linux.intel.com>
---
kernels/compiler_long_bitcast.cl | 47 ++++++++
utests/CMakeLists.txt | 1 +
utests/compiler_long_bitcast.cpp | 227 ++++++++++++++++++++++++++++++++++++++
3 files changed, 275 insertions(+)
create mode 100644 kernels/compiler_long_bitcast.cl
create mode 100644 utests/compiler_long_bitcast.cpp
diff --git a/kernels/compiler_long_bitcast.cl b/kernels/compiler_long_bitcast.cl
new file mode 100644
index 0000000..4b008ef
--- /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 0f56c6c..f8fb9c6 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 0000000..5bd962d
--- /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);
--
1.7.9.5
More information about the Beignet
mailing list