Fix some math function error in simd16.

INT DIV splite to simd8 but forget to set quarter_control.
Will fail when predication enable.
Change the atomic test case to trigger this bug.

Signed-off-by: Yang Rong <rong.r.yang@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

diff --git a/backend/src/backend/gen_encoder.cpp b/backend/src/backend/gen_encoder.cpp
index 43658e8..e10a04b 100644 (file)
--- a/backend/src/backend/gen_encoder.cpp
+++ b/backend/src/backend/gen_encoder.cpp
@@ -1048,6 +1048,7 @@ namespace gbe
      if (function == GEN_MATH_FUNCTION_INT_DIV_QUOTIENT ||
          function == GEN_MATH_FUNCTION_INT_DIV_REMAINDER) {
         insn->header.execution_size = GEN_WIDTH_8;
+        insn->header.quarter_control = GEN_COMPRESSION_Q1;
 
         if(this->curr.execWidth == 16) {
           GenInstruction *insn2 = this->next(GEN_OPCODE_MATH);
@@ -1058,6 +1059,7 @@ namespace gbe
           insn2->header.destreg_or_condmod = function;
           this->setHeader(insn2);
           insn2->header.execution_size = GEN_WIDTH_8;
+          insn2->header.quarter_control = GEN_COMPRESSION_Q2;
           this->setDst(insn2, new_dest);
           this->setSrc0(insn2, new_src0);
           this->setSrc1(insn2, new_src1);

diff --git a/kernels/compiler_atomic_functions.cl b/kernels/compiler_atomic_functions.cl
index 24f17c2..61ce2f4 100644 (file)
--- a/kernels/compiler_atomic_functions.cl
+++ b/kernels/compiler_atomic_functions.cl
@@ -7,8 +7,8 @@ __kernel void compiler_atomic_functions(__global int *dst, __local int *tmp, __g
     case 1: atomic_dec(&tmp[i]); break;
     case 2: atomic_add(&tmp[i], src[lid]); break;
     case 3: atomic_sub(&tmp[i], src[lid]); break;
-    case 4: atomic_and(&tmp[i], ~(src[lid]<<(lid>>2))); break;
-    case 5: atomic_or (&tmp[i], src[lid]<<(lid>>2)); break;
+    case 4: atomic_and(&tmp[i], ~(src[lid]<<(lid / 4))); break;
+    case 5: atomic_or (&tmp[i], src[lid]<<(lid / 4)); break;
     case 6: atomic_xor(&tmp[i], src[lid]); break;
     case 7: atomic_min(&tmp[i], -src[lid]); break;
     case 8: atomic_max(&tmp[i], src[lid]); break;
@@ -23,8 +23,8 @@ __kernel void compiler_atomic_functions(__global int *dst, __local int *tmp, __g
     case 1: atomic_dec(&dst[i]); break;
     case 2: atomic_add(&dst[i], src[lid]); break;
     case 3: atomic_sub(&dst[i], src[lid]); break;
-    case 4: atomic_and(&dst[i], ~(src[lid]<<(lid>>2))); break;
-    case 5: atomic_or (&dst[i], src[lid]<<(lid>>2)); break;
+    case 4: atomic_and(&dst[i], ~(src[lid]<<(lid / 4))); break;
+    case 5: atomic_or (&dst[i], src[lid]<<(lid / 4)); break;
     case 6: atomic_xor(&dst[i], src[lid]); break;
     case 7: atomic_min(&dst[i], -src[lid]); break;
     case 8: atomic_max(&dst[i], src[lid]); break;