[Beignet] [PATCH 4/5] GBE: handle ConstantExpr in program-scope variable handling.

[Ruiling Song]

although we have eliminate ConstantExpr in llvm instructions,
but in program scope variable, we still meet ConstantExpr.
So, we handle it here. also enhance the test case to hit it.

Signed-off-by: Ruiling Song <ruiling.song at intel.com>
---
 backend/src/llvm/llvm_gen_backend.cpp | 45 ++++++++++++++++++++++++++++++++++-
 kernels/compiler_program_global.cl    | 11 ++++++++-
 2 files changed, 54 insertions(+), 2 deletions(-)

diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
index 2b6875c..2942fa2 100644
--- a/backend/src/llvm/llvm_gen_backend.cpp
+++ b/backend/src/llvm/llvm_gen_backend.cpp
@@ -1439,6 +1439,36 @@ namespace gbe
     Type::TypeID id = type->getTypeID();
 
     GBE_ASSERT(c);
+    if (isa<ConstantExpr>(c)) {
+      const ConstantExpr *expr = dyn_cast<ConstantExpr>(c);
+      Value *pointer = expr->getOperand(0);
+      if (expr->getOpcode() == Instruction::GetElementPtr) {
+        uint32_t constantOffset = 0;
+        CompositeType* CompTy = cast<CompositeType>(pointer->getType());
+        for(uint32_t op=1; op<expr->getNumOperands(); ++op) {
+            int32_t TypeIndex;
+            ConstantInt* ConstOP = dyn_cast<ConstantInt>(expr->getOperand(op));
+            GBE_ASSERTM(ConstOP != NULL, "must be constant index");
+            TypeIndex = ConstOP->getZExtValue();
+            GBE_ASSERT(TypeIndex >= 0);
+            constantOffset += getGEPConstOffset(unit, CompTy, TypeIndex);
+            CompTy = dyn_cast<CompositeType>(CompTy->getTypeAtIndex(TypeIndex));
+        }
+
+        ir::Constant cc = unit.getConstantSet().getConstant(pointer->getName());
+        unsigned int defOffset = cc.getOffset();
+        relocs.push_back(ir::RelocEntry(offset, defOffset + constantOffset));
+
+        uint32_t size = getTypeByteSize(unit, type);
+        memset((char*)mem+offset, 0, size);
+        offset += size;
+      } else if (expr->isCast()) {
+        Constant *constPtr = cast<Constant>(pointer);
+        getConstantData(constPtr, mem, offset, relocs);
+        offset += getTypeByteSize(unit, type);
+      }
+      return;
+    }
     if (isa<GlobalVariable>(c)) {
       const GlobalVariable *GV = cast<GlobalVariable>(c);
 
@@ -1537,8 +1567,21 @@ namespace gbe
           offset += sizeof(double);
           break;
         }
-      default:
+      case Type::TypeID::HalfTyID:
+        {
+          const ConstantFP *cf = dyn_cast<ConstantFP>(c);
+          llvm::APFloat apf = cf->getValueAPF();
+          llvm::APInt api = apf.bitcastToAPInt();
+          uint64_t v64 = api.getZExtValue();
+          uint16_t v16 = static_cast<uint16_t>(v64);
+          *(unsigned short *)((char*)mem+offset) = v16;
+          offset += sizeof(short);
+          break;
+        }
+      default: {
+        c->dump();
         NOT_IMPLEMENTED;
+               }
     }
   }
 
diff --git a/kernels/compiler_program_global.cl b/kernels/compiler_program_global.cl
index 405c53f..fbe030f 100644
--- a/kernels/compiler_program_global.cl
+++ b/kernels/compiler_program_global.cl
@@ -31,6 +31,9 @@ global int a = 1;
 global int b = 2;
 global int * constant gArr[2]= {&a, &b};
 
+global int a_var[1] = {0};
+global int *p_var = a_var;
+
 __kernel void compiler_program_global0(const global int *src, int dynamic) {
   size_t gid = get_global_id(0);
   /* global read/write */
@@ -60,9 +63,15 @@ __kernel void compiler_program_global1(global int *dst, int dynamic) {
   dst[12] = *p;
   dst[13] = s;
   dst[14] = l;
-  dst[15] = *gArr[dynamic];
+  if (p_var == a_var)
+    dst[15] = *gArr[dynamic];
 
   if (gid < 11)
     dst[gid] = ba[gid];
 }
 
+__kernel void nouse(int dynamic) {
+  c[0].s1 = &s2;
+  p_var = a+dynamic;
+}
+
-- 
2.4.1