[Beignet] [PATCH] test 64bit-integer comparing
Xing, Homer
homer.xing at intel.com
Wed Aug 14 01:05:31 PDT 2013
Aha, I forgot to add a dependency. Thanks Zhigang.
-----Original Message-----
From: Zhigang Gong [mailto:zhigang.gong at linux.intel.com]
Sent: Wednesday, August 14, 2013 4:03 PM
To: Xing, Homer
Cc: beignet at lists.freedesktop.org
Subject: Re: [Beignet] [PATCH] test 64bit-integer comparing
This is a bug introduced by your last patch. You introduced a new cmp instruction but forgot to add corresponding dependency. I will fix it for you. Both patch LGTM, I will push them latter. Thanks.
On Wed, Aug 14, 2013 at 02:23:51PM +0800, Homer Hsing wrote:
> 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
>
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list