fix scalarizing of llvm phi node

llvm phi node can have odd number of args.
this patch also contains a test case.

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

4 files changed, 60 insertions, 1 deletions

diff --git a/backend/src/llvm/llvm_scalarize.cpp b/backend/src/llvm/llvm_scalarize.cpp
index 41674b6b..7a406164 100644
--- a/backend/src/llvm/llvm_scalarize.cpp
+++ b/backend/src/llvm/llvm_scalarize.cpp
@@ -383,7 +383,6 @@ namespace gbe {
 
     if (PHINode* phi = dyn_cast<PHINode>(inst)) {
       PHINode* res = PHINode::Create(GetBasicType(inst), phi->getNumIncomingValues());
-      assert(args.size() % 2 == 0 && "Odd number of arguments for a PHI");
 
       // Loop over pairs of operands: [Value*, BasicBlock*]
       for (unsigned int i = 0; i < args.size(); i++) {
diff --git a/kernels/compiler_vector_inc.cl b/kernels/compiler_vector_inc.cl
new file mode 100644
index 00000000..548dcb4f
--- /dev/null
+++ b/kernels/compiler_vector_inc.cl
@@ -0,0 +1,13 @@
+kernel void compiler_vector_inc(global char *dst, global char *src) {
+    size_t i = get_global_id(0);
+    char2 dst2 = vload2(i, dst);
+    if (src[i] == 0)
+      dst2++;
+    else if(src[i] == 1)
+      ++dst2;
+    else if(src[i] == 2)
+      dst2--;
+    else
+      --dst2;
+    vstore2(dst2, i, dst);
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 37261168..f18bd46f 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -102,6 +102,7 @@ set (utests_sources
   compiler_get_image_info.cpp
   compiler_vect_compare.cpp
   compiler_vector_load_store.cpp
+  compiler_vector_inc.cpp
   compiler_cl_finish.cpp
   get_cl_info.cpp
   builtin_atan2.cpp
diff --git a/utests/compiler_vector_inc.cpp b/utests/compiler_vector_inc.cpp
new file mode 100644
index 00000000..abc5408a
--- /dev/null
+++ b/utests/compiler_vector_inc.cpp
@@ -0,0 +1,46 @@
+#include <cstdint>
+#include <cstring>
+#include <iostream>
+#include "utest_helper.hpp"
+
+void compiler_vector_inc(void)
+{
+  const int n = 64;
+  char dst[n];
+  char src[n];
+
+  OCL_CREATE_KERNEL("compiler_vector_inc");
+  OCL_CREATE_BUFFER(buf[0], 0, n, NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, n, NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  globals[0] = n / 2;
+  locals[0] = 16;
+
+  for (int i = 0; i < n; ++i) {
+    dst[i] = i;
+    src[i] = (i / 2) % 4;
+  }
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  memcpy(buf_data[0], dst, n);
+  memcpy(buf_data[1], src, n);
+  OCL_UNMAP_BUFFER(0);
+  OCL_UNMAP_BUFFER(1);
+
+  OCL_NDRANGE(1);
+
+  OCL_MAP_BUFFER(0);
+  char *dest = ((char *)buf_data[0]);
+  for (int i=0; i<n; ++i) {
+    char wish;
+    if (src[i/2] < 2)
+      wish = dst[i] + 1;
+    else
+      wish = dst[i] - 1;
+    OCL_ASSERT(dest[i] == wish);
+  }
+  OCL_UNMAP_BUFFER(0);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_vector_inc);