[Beignet] [PATCH] test 64bit-integer comparing
Homer Hsing
homer.xing at intel.com
Tue Aug 13 23:23:51 PDT 2013
only work when OCL_POST_ALLOC_INSN_SCHEDULE=0
because the post alloc scheduler puts CMP after SEL, but in IR,
CMP is before SEL, like this
GT.int64 %34 %31 %33
LOADI.int64 %38 3
LOADI.int64 %39 4
SEL.int64 %35 %34 %38 %39
Signed-off-by: Homer Hsing <homer.xing at intel.com>
---
kernels/compiler_long_cmp.cl | 29 +++++++++++
utests/CMakeLists.txt | 1 +
utests/compiler_long_cmp.cpp | 117 +++++++++++++++++++++++++++++++++++++++++++
3 files changed, 147 insertions(+)
create mode 100644 kernels/compiler_long_cmp.cl
create mode 100644 utests/compiler_long_cmp.cpp
diff --git a/kernels/compiler_long_cmp.cl b/kernels/compiler_long_cmp.cl
new file mode 100644
index 0000000..90dfb60
--- /dev/null
+++ b/kernels/compiler_long_cmp.cl
@@ -0,0 +1,29 @@
+kernel void compiler_long_cmp_l(global long *src1, global long *src2, global long *dst) {
+ int i = get_global_id(0);
+ dst[i] = (src1[i] < src2[i]) ? 3 : 4;
+}
+
+kernel void compiler_long_cmp_le(global long *src1, global long *src2, global long *dst) {
+ int i = get_global_id(0);
+ dst[i] = (src1[i] <= src2[i]) ? 3 : 4;
+}
+
+kernel void compiler_long_cmp_g(global long *src1, global long *src2, global long *dst) {
+ int i = get_global_id(0);
+ dst[i] = (src1[i] > src2[i]) ? 3 : 4;
+}
+
+kernel void compiler_long_cmp_ge(global long *src1, global long *src2, global long *dst) {
+ int i = get_global_id(0);
+ dst[i] = (src1[i] >= src2[i]) ? 3 : 4;
+}
+
+kernel void compiler_long_cmp_eq(global long *src1, global long *src2, global long *dst) {
+ int i = get_global_id(0);
+ dst[i] = (src1[i] == src2[i]) ? 3 : 4;
+}
+
+kernel void compiler_long_cmp_neq(global long *src1, global long *src2, global long *dst) {
+ int i = get_global_id(0);
+ dst[i] = (src1[i] != src2[i]) ? 3 : 4;
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 0c98914..d3311ae 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -131,6 +131,7 @@ set (utests_sources
compiler_long_shr.cpp
compiler_long_asr.cpp
compiler_long_mult.cpp
+ compiler_long_cmp.cpp
utest_assert.cpp
utest.cpp
utest_file_map.cpp
diff --git a/utests/compiler_long_cmp.cpp b/utests/compiler_long_cmp.cpp
new file mode 100644
index 0000000..3775556
--- /dev/null
+++ b/utests/compiler_long_cmp.cpp
@@ -0,0 +1,117 @@
+#include <cstdint>
+#include <cstring>
+#include <iostream>
+#include "utest_helper.hpp"
+
+void compiler_long_cmp(void)
+{
+ const size_t n = 16;
+ int64_t src1[n], src2[n];
+
+ src1[0] = (int64_t)1 << 63, src2[0] = 0x7FFFFFFFFFFFFFFFll;
+ src1[1] = (int64_t)1 << 63, src2[1] = ((int64_t)1 << 63) | 1;
+ src1[2] = -1ll, src2[2] = 0;
+ src1[3] = ((int64_t)123 << 32) | 0x7FFFFFFF, src2[3] = ((int64_t)123 << 32) | 0x80000000;
+ src1[4] = 0x7FFFFFFFFFFFFFFFll, src2[4] = (int64_t)1 << 63;
+ src1[5] = ((int64_t)1 << 63) | 1, src2[5] = (int64_t)1 << 63;
+ src1[6] = 0, src2[6] = -1ll;
+ src1[7] = ((int64_t)123 << 32) | 0x80000000, src2[7] = ((int64_t)123 << 32) | 0x7FFFFFFF;
+ for(size_t i=8; i<n; i++) {
+ src1[i] = i;
+ src2[i] = i;
+ }
+
+ globals[0] = n;
+ locals[0] = 16;
+
+ OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(int64_t), NULL);
+ OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(int64_t), NULL);
+ OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(int64_t), NULL);
+ OCL_MAP_BUFFER(0);
+ OCL_MAP_BUFFER(1);
+ memcpy(buf_data[0], src1, sizeof(src1));
+ memcpy(buf_data[1], src2, sizeof(src2));
+ OCL_UNMAP_BUFFER(0);
+ OCL_UNMAP_BUFFER(1);
+
+
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_long_cmp", "compiler_long_cmp_l");
+ 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]);
+ OCL_NDRANGE(1);
+ OCL_MAP_BUFFER(2);
+ for (int32_t i = 0; i < (int32_t) n; ++i) {
+ int64_t *dest = (int64_t *)buf_data[2];
+ int64_t x = (src1[i] < src2[i]) ? 3 : 4;
+ OCL_ASSERT(x == dest[i]);
+ }
+ OCL_UNMAP_BUFFER(2);
+
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_long_cmp", "compiler_long_cmp_le");
+ 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]);
+ OCL_NDRANGE(1);
+ OCL_MAP_BUFFER(2);
+ for (int32_t i = 0; i < (int32_t) n; ++i) {
+ int64_t *dest = (int64_t *)buf_data[2];
+ int64_t x = (src1[i] <= src2[i]) ? 3 : 4;
+ OCL_ASSERT(x == dest[i]);
+ }
+ OCL_UNMAP_BUFFER(2);
+
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_long_cmp", "compiler_long_cmp_g");
+ 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]);
+ OCL_NDRANGE(1);
+ OCL_MAP_BUFFER(2);
+ for (int32_t i = 0; i < (int32_t) n; ++i) {
+ int64_t *dest = (int64_t *)buf_data[2];
+ int64_t x = (src1[i] > src2[i]) ? 3 : 4;
+ OCL_ASSERT(x == dest[i]);
+ }
+ OCL_UNMAP_BUFFER(2);
+
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_long_cmp", "compiler_long_cmp_ge");
+ 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]);
+ OCL_NDRANGE(1);
+ OCL_MAP_BUFFER(2);
+ for (int32_t i = 0; i < (int32_t) n; ++i) {
+ int64_t *dest = (int64_t *)buf_data[2];
+ int64_t x = (src1[i] >= src2[i]) ? 3 : 4;
+ OCL_ASSERT(x == dest[i]);
+ }
+ OCL_UNMAP_BUFFER(2);
+
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_long_cmp", "compiler_long_cmp_eq");
+ 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]);
+ OCL_NDRANGE(1);
+ OCL_MAP_BUFFER(2);
+ for (int32_t i = 0; i < (int32_t) n; ++i) {
+ int64_t *dest = (int64_t *)buf_data[2];
+ int64_t x = (src1[i] == src2[i]) ? 3 : 4;
+ OCL_ASSERT(x == dest[i]);
+ }
+ OCL_UNMAP_BUFFER(2);
+
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_long_cmp", "compiler_long_cmp_neq");
+ 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]);
+ OCL_NDRANGE(1);
+ OCL_MAP_BUFFER(2);
+ for (int32_t i = 0; i < (int32_t) n; ++i) {
+ int64_t *dest = (int64_t *)buf_data[2];
+ int64_t x = (src1[i] != src2[i]) ? 3 : 4;
+ OCL_ASSERT(x == dest[i]);
+ }
+ OCL_UNMAP_BUFFER(2);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_long_cmp);
--
1.8.1.2
More information about the Beignet
mailing list