[Beignet] OpenCL vector type compare behaviour

[Zhigang Gong]

I just wrote a patch for this bug.

Edward, could you have a try with it and let me know whether it works. Thanks.

>From 8f089e00a5b236883dd8c0f8e29ce595a2e861af Mon Sep 17 00:00:00 2001
From: Zhigang Gong <zhigang.gong at linux.intel.com>
Date: Wed, 26 Jun 2013 13:00:58 +0800
Subject: [PATCH] GBE: fixed the bug when sext a i1 to i8/i16/i32.

We need to extent it to -1 rather than 1.
Reported by Edward Ching <edward.k.ching at gmail.com>.

Signed-off-by: Zhigang Gong <zhigang.gong at linux.intel.com>
---
 backend/src/llvm/llvm_gen_backend.cpp | 7 ++++++-
 1 file changed, 6 insertions(+), 1 deletion(-)

diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
index 5b7754c..a6612f5 100644
--- a/backend/src/llvm/llvm_gen_backend.cpp
+++ b/backend/src/llvm/llvm_gen_backend.cpp
@@ -1480,7 +1480,12 @@ namespace gbe
         if (srcType == ir::TYPE_BOOL) {
           const ir::RegisterFamily family = getFamily(dstType);
           const ir::ImmediateIndex zero = ctx.newIntegerImmediate(0, dstType);
-          const ir::ImmediateIndex one = ctx.newIntegerImmediate(1, dstType);
+          ir::ImmediateIndex one;
+          if (I.getOpcode() == Instruction::SExt
+              && (dstType == ir::TYPE_S8 || dstType == ir::TYPE_S16 || dstType == ir::TYPE_S32))
+            one = ctx.newIntegerImmediate(-1, dstType);
+          else
+            one = ctx.newIntegerImmediate(1, dstType);
           const ir::Register zeroReg = ctx.reg(family);
           const ir::Register oneReg = ctx.reg(family);
           ctx.LOADI(dstType, zeroReg, zero);
-- 
1.7.11.7



On Wed, Jun 26, 2013 at 01:11:27AM +0000, Zou, Nanhai wrote:
> Thanks for reporting this.
> Intel SDK is right.
> I have check the spec, OpenCL spec said if compare relation is true it should return -1 with all bit set.
> We will fix that.
> 
> Thanks
> Zou Nanhai
> 
> From: beignet-bounces+nanhai.zou=intel.com at lists.freedesktop.org [mailto:beignet-bounces+nanhai.zou=intel.com at lists.freedesktop.org] On Behalf Of Edward Ching
> Sent: Wednesday, June 26, 2013 8:58 AM
> To: beignet at lists.freedesktop.org
> Subject: [Beignet] OpenCL vector type compare behaviour
> 
> What is the correct behaviour when comparing vector type values in OpenCL?
> e.g.
> 
> __kernel void
> compiler_int4_comp(__global int4 *src, __global int4 *dst)
> {
>   int4 v1=(int4)(0,1,1,0);
>   int4 v2=(int4)(1,0,0,0);
>   int4 mask;
>  ...
>  ...
>   mask = v1 < v2;
>   ...
>   ...
> }
> 
> Using Beignet OpenCL, the value in mask after executing the mask=v1<v2 statement,  is (0x1, 0x0, 0x0, 0x1).
> However, using Intel OpenCL SDK 2012 and 2013, the result is (0xffffffff,0x0, 0x0, 0xffffffff).
> Which one is the correct behavior? There are algorithms (such as the bitonic sort in Intel SDK) that assumes the Intel OpenCL SDK behavior, and this kind of subtlety is a pain to debug.
> /Ed
> 

> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet